#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void transformKernel(float *outputData, int width, int height, float theta, cudaTextureObject_t tex){
    // calculate normalized texture coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    float u = (float)x - (float)width/2; 
    float v = (float)y - (float)height/2; 
    float tu = u*cosf(theta) - v*sinf(theta); 
    float tv = v*cosf(theta) + u*sinf(theta); 

    tu /= (float)width; 
    tv /= (float)height; 

    // read from texture and write to global memory
    outputData[y*width + x] = tex2D<float>(tex, tu+0.5f, tv+0.5f);
}

void loadPGM(float* _imageData, const char* _imagePath, int* width, int* height){
    for(int i=0;i<256;i++){
        _imageData[i] = i+0.0; 
    }
    *width = 16;
    *height = 16;
    return;
}

int main(){
    /**
        1. load image and reference image
        2. allocate device memory for result
        3. [most important] allocate array and copy image data, then bind it to texture memory (as well as copy host data to device)
        4. issure kernel 
        5. copy result of device to host
        6. do post process
    */

    int width = 0;
    int height = 0;
    int size = 256*sizeof(float);
    const float angle = 0.5f; 
    // load image
    float* h_imageData = (float*) malloc(256*sizeof(float));
    const char* imagePath = "./test.pgm";
    loadPGM(h_imageData, imagePath, &width, &height);

    

    // load reference image
    //float* h_refData;
    //const char* refPath = "./testRef.pgm";
    //loadPGM(h_refData, refPath, );


    float* h_resData = (float*) malloc(256*sizeof(float));

    // allocate device memory for res, inorder to copy result from gpu to cpu.
    float* d_resData;
    cudaMalloc((void **) &d_resData, size);
    // allocate arry and copy image data, bind to texture memory 

    //first cudaChannelDesc
    //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); //low level, C style
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();                                         //high level, C++ style, using template
    
    //second, create cudaArray used for texture memory and fill it with h_imageData
    cudaArray* cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);
    cudaMemcpyToArray(cuArray, 0, 0, h_imageData, size, cudaMemcpyHostToDevice);

    //third, create cudaResourceDesc (with cudaArray) and cudaTextureDesc
    cudaResourceDesc texRes;
    memset(&texRes, 0, sizeof(cudaResourceDesc));

    texRes.resType = cudaResourceTypeArray;
    texRes.res.array.array = cuArray;

    cudaTextureDesc texDescr;
    memset(&texDescr, 0, sizeof(cudaTextureDesc));

    texDescr.normalizedCoords = true;
    texDescr.filterMode = cudaFilterModeLinear;
    texDescr.addressMode[0] = cudaAddressModeWrap;
    texDescr.addressMode[1] = cudaAddressModeWrap;
    texDescr.readMode = cudaReadModeElementType;

    // fourth, create cudaTextureObject
    cudaTextureObject_t tex;            //aka cudaTextureObject*
    cudaCreateTextureObject(&tex, &texRes, &texDescr, NULL);                //Bind cudaArray(imageData) to texture memory HERE!!!!

    //Done! Issue kernel       
    //Do not need warm up, just for testing
    dim3 threadsPerBlock(8, 8, 1);                                                 //
    dim3 numBlocks(width / threadsPerBlock.x, height / threadsPerBlock.y, 1);
    transformKernel<<<numBlocks,threadsPerBlock>>>(d_resData, width, height, angle, tex);

    //copy result from gpu to host
    cudaMemcpy(h_resData, d_resData, size, cudaMemcpyDeviceToHost);


    //Do post processing
    //sdkSavePGM(outputFilename, hOutputData, width, height);

    //check result
    for(int i=0;i<256;i++){
        printf("%lf ", h_resData[i]);
    }
}
	.file	"tmpxft_0037971d_00000000-6_simpleTexture.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z7loadPGMPfPKcPiS2_
	.type	_Z7loadPGMPfPKcPiS2_, @function
_Z7loadPGMPfPKcPiS2_:
.LFB2027:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L3:
	cvtsi2ssl	%eax, %xmm0
	movss	%xmm0, (%rdi,%rax,4)
	incq	%rax
	cmpq	$256, %rax
	jne	.L3
	movl	$16, (%rdx)
	movl	$16, (%rcx)
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z7loadPGMPfPKcPiS2_, .-_Z7loadPGMPfPKcPiS2_
	.globl	_Z39__device_stub__Z15transformKernelPfiifyPfiify
	.type	_Z39__device_stub__Z15transformKernelPfiifyPfiify, @function
_Z39__device_stub__Z15transformKernelPfiifyPfiify:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 20(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L6
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z15transformKernelPfiify(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L6:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L8
	call	__stack_chk_fail@PLT
.L8:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z39__device_stub__Z15transformKernelPfiifyPfiify, .-_Z39__device_stub__Z15transformKernelPfiifyPfiify
	.globl	_Z15transformKernelPfiify
	.type	_Z15transformKernelPfiify, @function
_Z15transformKernelPfiify:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z15transformKernelPfiifyPfiify
	.cfi_endproc
.LFE2054:
	.size	_Z15transformKernelPfiify, .-_Z15transformKernelPfiify
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"./test.pgm"
.LC2:
	.string	"%lf "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$1024, %edi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$224, %rsp
	.cfi_def_cfa_offset 272
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	60(%rsp), %r14
	call	malloc@PLT
	leaq	4(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	.LC0(%rip), %rsi
	movq	%rax, %rdi
	movq	%rax, %rbx
	call	_Z7loadPGMPfPKcPiS2_
	movl	$1024, %edi
	call	malloc@PLT
	leaq	8(%rsp), %rdi
	movl	$1024, %esi
	movq	%rax, %rbp
	call	cudaMalloc@PLT
	movl	$2, %r9d
	xorl	%r8d, %r8d
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	movl	$32, %esi
	movq	%r14, %rdi
	call	cudaCreateChannelDesc@PLT
	movslq	4(%rsp), %rcx
	xorl	%r8d, %r8d
	movq	%r14, %rsi
	movslq	(%rsp), %rdx
	leaq	16(%rsp), %rdi
	movq	%rcx, %r13
	movq	%rdx, %r12
	call	cudaMallocArray@PLT
	movq	16(%rsp), %rdi
	xorl	%edx, %edx
	xorl	%esi, %esi
	movl	$1, %r9d
	movl	$1024, %r8d
	movq	%rbx, %rcx
	call	cudaMemcpyToArray@PLT
	movq	16(%rsp), %rdx
	leaq	80(%rsp), %rsi
	xorl	%eax, %eax
	movq	%rsi, %rdi
	movl	$16, %ecx
	rep stosl
	movq	%rdx, 88(%rsp)
	leaq	144(%rsp), %rdx
	movl	$18, %ecx
	movq	%rdx, %rdi
	rep stosl
	leaq	24(%rsp), %rdi
	movl	$1, 184(%rsp)
	movl	$1, 156(%rsp)
	call	cudaCreateTextureObject@PLT
	movl	%r12d, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	shrl	$3, %eax
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$34359738376, %rdx
	movl	%eax, 48(%rsp)
	movl	%r13d, %eax
	shrl	$3, %eax
	movl	%eax, 52(%rsp)
	movq	48(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	movq	24(%rsp), %rcx
	movq	8(%rsp), %rdi
	movl	%r13d, %edx
	movl	%r12d, %esi
	movss	.LC1(%rip), %xmm0
	call	_Z39__device_stub__Z15transformKernelPfiifyPfiify
.L12:
	movq	8(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbp, %rdi
	xorl	%ebx, %ebx
	movl	$1024, %edx
	leaq	.LC2(%rip), %r12
	call	cudaMemcpy@PLT
.L13:
	cvtss2sd	0(%rbp,%rbx,4), %xmm0
	movq	%r12, %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	incq	%rbx
	cmpq	$256, %rbx
	jne	.L13
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$224, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC3:
	.string	"_Z15transformKernelPfiify"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC3(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z15transformKernelPfiify(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2056:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	1056964608
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#include <hip/hip_runtime.h>

#include <stdio.h>
__global__ void transformKernel(float *outputData, int width, int height, float theta, hipTextureObject_t tex){
    // calculate normalized texture coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

    float u = (float)x - (float)width/2; 
    float v = (float)y - (float)height/2; 
    float tu = u*cosf(theta) - v*sinf(theta); 
    float tv = v*cosf(theta) + u*sinf(theta); 

    tu /= (float)width; 
    tv /= (float)height; 

    // read from texture and write to global memory
    outputData[y*width + x] = tex2D<float>(tex, tu+0.5f, tv+0.5f);
}

void loadPGM(float* _imageData, const char* _imagePath, int* width, int* height){
    for(int i=0;i<256;i++){
        _imageData[i] = i+0.0; 
    }
    *width = 16;
    *height = 16;
    return;
}

int main(){
    /**
        1. load image and reference image
        2. allocate device memory for result
        3. [most important] allocate array and copy image data, then bind it to texture memory (as well as copy host data to device)
        4. issure kernel 
        5. copy result of device to host
        6. do post process
    */

    int width = 0;
    int height = 0;
    int size = 256*sizeof(float);
    const float angle = 0.5f; 
    // load image
    float* h_imageData = (float*) malloc(256*sizeof(float));
    const char* imagePath = "./test.pgm";
    loadPGM(h_imageData, imagePath, &width, &height);

    

    // load reference image
    //float* h_refData;
    //const char* refPath = "./testRef.pgm";
    //loadPGM(h_refData, refPath, );


    float* h_resData = (float*) malloc(256*sizeof(float));

    // allocate device memory for res, inorder to copy result from gpu to cpu.
    float* d_resData;
    hipMalloc((void **) &d_resData, size);
    // allocate arry and copy image data, bind to texture memory 

    //first cudaChannelDesc
    //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); //low level, C style
    hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();                                         //high level, C++ style, using template
    
    //second, create cudaArray used for texture memory and fill it with h_imageData
    hipArray* cuArray;
    hipMallocArray(&cuArray, &channelDesc, width, height);
    hipMemcpyToArray(cuArray, 0, 0, h_imageData, size, hipMemcpyHostToDevice);

    //third, create cudaResourceDesc (with cudaArray) and cudaTextureDesc
    hipResourceDesc texRes;
    memset(&texRes, 0, sizeof(hipResourceDesc));

    texRes.resType = hipResourceTypeArray;
    texRes.res.array.array = cuArray;

    hipTextureDesc texDescr;
    memset(&texDescr, 0, sizeof(hipTextureDesc));

    texDescr.normalizedCoords = true;
    texDescr.filterMode = hipFilterModeLinear;
    texDescr.addressMode[0] = hipAddressModeWrap;
    texDescr.addressMode[1] = hipAddressModeWrap;
    texDescr.readMode = hipReadModeElementType;

    // fourth, create cudaTextureObject
    hipTextureObject_t tex;            //aka cudaTextureObject*
    hipCreateTextureObject(&tex, &texRes, &texDescr, NULL);                //Bind cudaArray(imageData) to texture memory HERE!!!!

    //Done! Issue kernel       
    //Do not need warm up, just for testing
    dim3 threadsPerBlock(8, 8, 1);                                                 //
    dim3 numBlocks(width / threadsPerBlock.x, height / threadsPerBlock.y, 1);
    transformKernel<<<numBlocks,threadsPerBlock>>>(d_resData, width, height, angle, tex);

    //copy result from gpu to host
    hipMemcpy(h_resData, d_resData, size, hipMemcpyDeviceToHost);


    //Do post processing
    //sdkSavePGM(outputFilename, hOutputData, width, height);

    //check result
    for(int i=0;i<256;i++){
        printf("%lf ", h_resData[i]);
    }
}
	.text
	.file	"simpleTexture.hip"
	.globl	_Z30__device_stub__transformKernelPfiifP13__hip_texture # -- Begin function _Z30__device_stub__transformKernelPfiifP13__hip_texture
	.type	_Z30__device_stub__transformKernelPfiifP13__hip_texture,@function
_Z30__device_stub__transformKernelPfiifP13__hip_texture: # @_Z30__device_stub__transformKernelPfiifP13__hip_texture
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	12(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	8(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	4(%rsp), %rdx
	movss	%xmm0, (%rdx)
	leaq	32(%rsp), %r8
	movq	%rcx, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%r8, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z15transformKernelPfiifP13__hip_texture, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z30__device_stub__transformKernelPfiifP13__hip_texture, .Lfunc_end0-_Z30__device_stub__transformKernelPfiifP13__hip_texture
	.cfi_endproc
                                        # -- End function
	.globl	_Z7loadPGMPfPKcPiS2_            # -- Begin function _Z7loadPGMPfPKcPiS2_
	.type	_Z7loadPGMPfPKcPiS2_,@function
_Z7loadPGMPfPKcPiS2_:                   # @_Z7loadPGMPfPKcPiS2_
	.cfi_startproc
# %bb.0:
	xorl	%eax, %eax
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, (%rdi,%rax,4)
	incq	%rax
	cmpq	$256, %rax                      # imm = 0x100
	jne	.LBB1_1
# %bb.2:
	movl	$16, %eax
	movl	%eax, (%rdx)
	movl	%eax, (%rcx)
	retq
.Lfunc_end1:
	.size	_Z7loadPGMPfPKcPiS2_, .Lfunc_end1-_Z7loadPGMPfPKcPiS2_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI2_0:
	.long	0x3f000000                      # float 0.5
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$1024, %edi                     # imm = 0x400
	callq	malloc
	movq	%rax, %r14
	xorl	%eax, %eax
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, (%r14,%rax,4)
	incq	%rax
	cmpq	$256, %rax                      # imm = 0x100
	jne	.LBB2_1
# %bb.2:                                # %_Z7loadPGMPfPKcPiS2_.exit
	movl	$1024, %edi                     # imm = 0x400
	callq	malloc
	movq	%rax, %rbx
	movq	%rsp, %rdi
	movl	$1024, %esi                     # imm = 0x400
	callq	hipMalloc
	leaq	28(%rsp), %r12
	xorl	%r13d, %r13d
	movq	%r12, %rdi
	movl	$32, %esi
	xorl	%edx, %edx
	xorl	%ecx, %ecx
	xorl	%r8d, %r8d
	movl	$2, %r9d
	callq	hipCreateChannelDesc
	leaq	16(%rsp), %r15
	movl	$16, %edx
	movl	$16, %ecx
	movq	%r15, %rdi
	movq	%r12, %rsi
	xorl	%r8d, %r8d
	callq	hipMallocArray
	movq	(%r15), %rdi
	movl	$1, %ebp
	movl	$1024, %r8d                     # imm = 0x400
	xorl	%esi, %esi
	xorl	%edx, %edx
	movq	%r14, %rcx
	movl	$1, %r9d
	callq	hipMemcpyToArray
	xorps	%xmm0, %xmm0
	leaq	48(%rsp), %rsi
	movaps	%xmm0, (%rsi)
	movaps	%xmm0, 48(%rsi)
	movaps	%xmm0, 32(%rsi)
	movaps	%xmm0, 16(%rsi)
	movq	(%r15), %rax
	movq	%rax, 8(%rsi)
	leaq	120(%rsp), %rdx
	movups	%xmm0, 40(%rdx)
	movups	%xmm0, 8(%rdx)
	movq	%r13, 56(%rdx)
	movups	%xmm0, 24(%rdx)
	movl	%ebp, 40(%rdx)
	movl	%ebp, 12(%rdx)
	movq	%r13, (%rdx)
	leaq	8(%rsp), %rdi
	xorl	%ecx, %ecx
	callq	hipCreateTextureObject
	movabsq	$8589934594, %rdi               # imm = 0x200000002
	movabsq	$34359738376, %rdx              # imm = 0x800000008
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_4
# %bb.3:
	movq	(%rsp), %rdi
	movq	8(%rsp), %rcx
	movss	.LCPI2_0(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	movl	$16, %esi
	movl	$16, %edx
	callq	_Z30__device_stub__transformKernelPfiifP13__hip_texture
.LBB2_4:
	movq	(%rsp), %rsi
	movl	$1024, %edx                     # imm = 0x400
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%r14d, %r14d
.LBB2_5:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r14,4), %xmm0
	movl	$.L.str.1, %edi
	movb	$1, %al
	callq	printf
	incq	%r14
	cmpq	$256, %r14                      # imm = 0x100
	jne	.LBB2_5
# %bb.6:
	xorl	%eax, %eax
	addq	$184, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15transformKernelPfiifP13__hip_texture, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB4_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z15transformKernelPfiifP13__hip_texture,@object # @_Z15transformKernelPfiifP13__hip_texture
	.section	.rodata,"a",@progbits
	.globl	_Z15transformKernelPfiifP13__hip_texture
	.p2align	3, 0x0
_Z15transformKernelPfiifP13__hip_texture:
	.quad	_Z30__device_stub__transformKernelPfiifP13__hip_texture
	.size	_Z15transformKernelPfiifP13__hip_texture, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"%lf "
	.size	.L.str.1, 5

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z15transformKernelPfiifP13__hip_texture"
	.size	.L__unnamed_1, 41

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z30__device_stub__transformKernelPfiifP13__hip_texture
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z15transformKernelPfiifP13__hip_texture
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z15transformKernelPfiifP13__hip_texture ; -- Begin function _Z15transformKernelPfiifP13__hip_texture
	.globl	_Z15transformKernelPfiifP13__hip_texture
	.p2align	8
	.type	_Z15transformKernelPfiifP13__hip_texture,@function
_Z15transformKernelPfiifP13__hip_texture: ; @_Z15transformKernelPfiifP13__hip_texture
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	v_cmp_ngt_f32_e64 s2, 0x48000000, |s6|
	s_and_b32 s3, s6, 0x7fffffff
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, exec_lo, s2
	s_mov_b32 vcc_lo, s2
	s_cbranch_vccz .LBB0_2
; %bb.1:
	s_and_b32 s7, s3, 0x7fffff
	s_lshr_b32 s8, s3, 23
	s_bitset1_b32 s7, 23
	s_addk_i32 s8, 0xff88
	s_mul_hi_u32 s9, s7, 0xfe5163ab
	s_mul_i32 s10, s7, 0x3c439041
	s_mul_hi_u32 s11, s7, 0x3c439041
	s_add_u32 s9, s9, s10
	s_addc_u32 s10, 0, s11
	s_mul_i32 s11, s7, 0xdb629599
	s_mul_hi_u32 s12, s7, 0xdb629599
	s_add_u32 s10, s10, s11
	s_addc_u32 s11, 0, s12
	s_mul_i32 s12, s7, 0xf534ddc0
	s_mul_hi_u32 s13, s7, 0xf534ddc0
	s_add_u32 s11, s11, s12
	s_addc_u32 s12, 0, s13
	s_mul_i32 s13, s7, 0xfc2757d1
	s_mul_hi_u32 s16, s7, 0xfc2757d1
	s_add_u32 s12, s12, s13
	s_addc_u32 s13, 0, s16
	s_mul_i32 s16, s7, 0x4e441529
	s_mul_hi_u32 s17, s7, 0x4e441529
	s_add_u32 s13, s13, s16
	s_addc_u32 s16, 0, s17
	s_cmp_gt_u32 s8, 63
	s_mul_i32 s17, s7, 0xfe5163ab
	s_mul_hi_u32 s18, s7, 0xa2f9836e
	s_mul_i32 s7, s7, 0xa2f9836e
	s_cselect_b32 s19, s10, s12
	s_cselect_b32 s9, s9, s11
	s_cselect_b32 s10, s17, s10
	s_add_u32 s7, s16, s7
	s_addc_u32 s16, 0, s18
	s_cmp_gt_u32 s8, 63
	s_cselect_b32 s17, 0xffffffc0, 0
	s_cselect_b32 s11, s11, s13
	s_cselect_b32 s7, s12, s7
	s_cselect_b32 s12, s13, s16
	s_add_i32 s17, s17, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s17, 31
	s_cselect_b32 s8, 0xffffffe0, 0
	s_cselect_b32 s13, s11, s7
	s_cselect_b32 s7, s7, s12
	s_cselect_b32 s11, s19, s11
	s_cselect_b32 s12, s9, s19
	s_cselect_b32 s9, s10, s9
	s_add_i32 s8, s8, s17
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s8, 31
	s_cselect_b32 s10, 0xffffffe0, 0
	s_cselect_b32 s7, s13, s7
	s_cselect_b32 s13, s11, s13
	s_cselect_b32 s11, s12, s11
	s_cselect_b32 s9, s9, s12
	s_add_i32 s10, s10, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_sub_i32 s8, 32, s10
	s_cmp_eq_u32 s10, 0
	v_mov_b32_e32 v1, s8
	s_cselect_b32 s10, -1, 0
	v_alignbit_b32 v2, s7, s13, v1
	v_alignbit_b32 v3, s13, s11, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s8, v2
	v_cndmask_b32_e64 v2, v3, s13, s10
	s_delay_alu instid0(VALU_DEP_2)
	s_cselect_b32 s7, s7, s8
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_alignbit_b32 v3, s7, v2, 30
	s_bfe_u32 s12, s7, 0x1001d
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s8, 0, s12
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_xor_b32_e32 v3, s8, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v4, v3
	v_min_u32_e32 v4, 32, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v6, 23, v4
	v_alignbit_b32 v1, s11, s9, v1
	v_sub_nc_u32_e32 v5, 31, v4
	v_cndmask_b32_e64 v1, v1, s11, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v2, v2, v1, 30
	v_alignbit_b32 v1, v1, s9, 30
	v_xor_b32_e32 v2, s8, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_xor_b32_e32 v1, s8, v1
	s_lshr_b32 s8, s7, 29
	s_lshr_b32 s7, s7, 30
	s_lshl_b32 s8, s8, 31
	v_alignbit_b32 v3, v3, v2, v5
	v_alignbit_b32 v1, v2, v1, v5
	s_or_b32 s9, s8, 0.5
	s_add_i32 s7, s12, s7
	v_sub_nc_u32_e32 v6, s9, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v2, v3, v1, 9
	v_clz_i32_u32_e32 v5, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v5, 32, v5
	v_sub_nc_u32_e32 v7, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v1, v2, v1, v7
	v_lshrrev_b32_e32 v2, 9, v3
	v_lshrrev_b32_e32 v1, 9, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_or_b32_e32 v2, v2, v6
	v_add_nc_u32_e32 v4, v5, v4
	v_lshlrev_b32_e32 v3, 23, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v1, v1, v3
	v_mul_f32_e32 v3, 0x3fc90fda, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v1, 0x33000000, v1
	v_fma_f32 v4, 0x3fc90fda, v2, -v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v1, s8, v1
	v_fmamk_f32 v2, v2, 0x33a22168, v4
	s_mov_b32 s8, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v2, 0x3fc90fda, v1
	v_add_f32_e32 v1, v3, v2
	s_branch .LBB0_3
.LBB0_2:
	s_mov_b32 s8, -1
                                        ; implicit-def: $sgpr7
                                        ; implicit-def: $vgpr1
.LBB0_3:                                ; %Flow71
	v_mul_f32_e64 v3, 0x3f22f983, |s6|
	s_and_not1_b32 vcc_lo, exec_lo, s8
	s_cbranch_vccnz .LBB0_5
; %bb.4:
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f32_e32 v2, v3
	v_fma_f32 v1, 0xbfc90fda, v2, |s6|
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmamk_f32 v1, v2, 0xb3a22168, v1
	v_fmamk_f32 v1, v2, 0xa7c234c4, v1
	v_cvt_i32_f32_e32 v2, v2
	s_branch .LBB0_6
.LBB0_5:
	v_mov_b32_e32 v2, s7
.LBB0_6:                                ; %_ZL4cosff.exit
	s_mov_b32 vcc_lo, s2
	s_cbranch_vccz .LBB0_8
; %bb.7:
	s_and_b32 s7, s3, 0x7fffff
	s_lshr_b32 s8, s3, 23
	s_bitset1_b32 s7, 23
	s_addk_i32 s8, 0xff88
	s_mul_hi_u32 s9, s7, 0xfe5163ab
	s_mul_i32 s10, s7, 0x3c439041
	s_mul_hi_u32 s11, s7, 0x3c439041
	s_add_u32 s9, s9, s10
	s_addc_u32 s10, 0, s11
	s_mul_i32 s11, s7, 0xdb629599
	s_mul_hi_u32 s12, s7, 0xdb629599
	s_add_u32 s10, s10, s11
	s_addc_u32 s11, 0, s12
	s_mul_i32 s12, s7, 0xf534ddc0
	s_mul_hi_u32 s13, s7, 0xf534ddc0
	s_add_u32 s11, s11, s12
	s_addc_u32 s12, 0, s13
	s_mul_i32 s13, s7, 0xfc2757d1
	s_mul_hi_u32 s16, s7, 0xfc2757d1
	s_add_u32 s12, s12, s13
	s_addc_u32 s13, 0, s16
	s_mul_i32 s16, s7, 0x4e441529
	s_mul_hi_u32 s17, s7, 0x4e441529
	s_add_u32 s13, s13, s16
	s_addc_u32 s16, 0, s17
	s_cmp_gt_u32 s8, 63
	s_mul_i32 s17, s7, 0xfe5163ab
	s_mul_hi_u32 s18, s7, 0xa2f9836e
	s_mul_i32 s7, s7, 0xa2f9836e
	s_cselect_b32 s19, s10, s12
	s_cselect_b32 s9, s9, s11
	s_cselect_b32 s10, s17, s10
	s_add_u32 s7, s16, s7
	s_addc_u32 s16, 0, s18
	s_cmp_gt_u32 s8, 63
	s_cselect_b32 s17, 0xffffffc0, 0
	s_cselect_b32 s11, s11, s13
	s_cselect_b32 s7, s12, s7
	s_cselect_b32 s12, s13, s16
	s_add_i32 s17, s17, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s17, 31
	s_cselect_b32 s8, 0xffffffe0, 0
	s_cselect_b32 s13, s11, s7
	s_cselect_b32 s7, s7, s12
	s_cselect_b32 s11, s19, s11
	s_cselect_b32 s12, s9, s19
	s_cselect_b32 s9, s10, s9
	s_add_i32 s8, s8, s17
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s8, 31
	s_cselect_b32 s10, 0xffffffe0, 0
	s_cselect_b32 s7, s13, s7
	s_cselect_b32 s13, s11, s13
	s_cselect_b32 s11, s12, s11
	s_cselect_b32 s9, s9, s12
	s_add_i32 s10, s10, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_sub_i32 s8, 32, s10
	s_cmp_eq_u32 s10, 0
	v_mov_b32_e32 v4, s8
	s_cselect_b32 s10, -1, 0
	v_alignbit_b32 v5, s7, s13, v4
	v_alignbit_b32 v6, s13, s11, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s8, v5
	v_cndmask_b32_e64 v5, v6, s13, s10
	s_delay_alu instid0(VALU_DEP_2)
	s_cselect_b32 s7, s7, s8
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_alignbit_b32 v6, s7, v5, 30
	s_bfe_u32 s12, s7, 0x1001d
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s8, 0, s12
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_xor_b32_e32 v6, s8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v7, v6
	v_min_u32_e32 v7, 32, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v9, 23, v7
	v_alignbit_b32 v4, s11, s9, v4
	v_sub_nc_u32_e32 v8, 31, v7
	v_cndmask_b32_e64 v4, v4, s11, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v5, v5, v4, 30
	v_alignbit_b32 v4, v4, s9, 30
	v_xor_b32_e32 v5, s8, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_xor_b32_e32 v4, s8, v4
	s_lshr_b32 s8, s7, 29
	s_lshr_b32 s7, s7, 30
	s_lshl_b32 s8, s8, 31
	v_alignbit_b32 v6, v6, v5, v8
	v_alignbit_b32 v4, v5, v4, v8
	s_or_b32 s9, s8, 0.5
	s_add_i32 s7, s12, s7
	v_sub_nc_u32_e32 v9, s9, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v5, v6, v4, 9
	v_clz_i32_u32_e32 v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v8, 32, v8
	v_sub_nc_u32_e32 v10, 31, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v4, v5, v4, v10
	v_lshrrev_b32_e32 v5, 9, v6
	v_lshrrev_b32_e32 v4, 9, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_or_b32_e32 v5, v5, v9
	v_add_nc_u32_e32 v7, v8, v7
	v_lshlrev_b32_e32 v6, 23, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v4, v4, v6
	v_mul_f32_e32 v6, 0x3fc90fda, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, 0x33000000, v4
	v_fma_f32 v7, 0x3fc90fda, v5, -v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v4, s8, v4
	v_fmamk_f32 v5, v5, 0x33a22168, v7
	s_mov_b32 s8, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, 0x3fc90fda, v4
	v_add_f32_e32 v4, v6, v5
	s_branch .LBB0_9
.LBB0_8:
	s_mov_b32 s8, -1
                                        ; implicit-def: $sgpr7
                                        ; implicit-def: $vgpr4
.LBB0_9:                                ; %Flow70
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s8
	s_cbranch_vccnz .LBB0_11
; %bb.10:
	v_rndne_f32_e32 v5, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v4, 0xbfc90fda, v5, |s6|
	v_fmamk_f32 v4, v5, 0xb3a22168, v4
	s_delay_alu instid0(VALU_DEP_1)
	v_fmamk_f32 v4, v5, 0xa7c234c4, v4
	v_cvt_i32_f32_e32 v5, v5
	s_branch .LBB0_12
.LBB0_11:
	v_mov_b32_e32 v5, s7
.LBB0_12:                               ; %_ZL4sinff.exit
	s_mov_b32 vcc_lo, s2
	s_cbranch_vccz .LBB0_14
; %bb.13:
	s_and_b32 s7, s3, 0x7fffff
	s_lshr_b32 s8, s3, 23
	s_bitset1_b32 s7, 23
	s_addk_i32 s8, 0xff88
	s_mul_hi_u32 s9, s7, 0xfe5163ab
	s_mul_i32 s10, s7, 0x3c439041
	s_mul_hi_u32 s11, s7, 0x3c439041
	s_add_u32 s9, s9, s10
	s_addc_u32 s10, 0, s11
	s_mul_i32 s11, s7, 0xdb629599
	s_mul_hi_u32 s12, s7, 0xdb629599
	s_add_u32 s10, s10, s11
	s_addc_u32 s11, 0, s12
	s_mul_i32 s12, s7, 0xf534ddc0
	s_mul_hi_u32 s13, s7, 0xf534ddc0
	s_add_u32 s11, s11, s12
	s_addc_u32 s12, 0, s13
	s_mul_i32 s13, s7, 0xfc2757d1
	s_mul_hi_u32 s16, s7, 0xfc2757d1
	s_add_u32 s12, s12, s13
	s_addc_u32 s13, 0, s16
	s_mul_i32 s16, s7, 0x4e441529
	s_mul_hi_u32 s17, s7, 0x4e441529
	s_add_u32 s13, s13, s16
	s_addc_u32 s16, 0, s17
	s_cmp_gt_u32 s8, 63
	s_mul_i32 s17, s7, 0xfe5163ab
	s_mul_hi_u32 s18, s7, 0xa2f9836e
	s_mul_i32 s7, s7, 0xa2f9836e
	s_cselect_b32 s19, s10, s12
	s_cselect_b32 s9, s9, s11
	s_cselect_b32 s10, s17, s10
	s_add_u32 s7, s16, s7
	s_addc_u32 s16, 0, s18
	s_cmp_gt_u32 s8, 63
	s_cselect_b32 s17, 0xffffffc0, 0
	s_cselect_b32 s11, s11, s13
	s_cselect_b32 s7, s12, s7
	s_cselect_b32 s12, s13, s16
	s_add_i32 s17, s17, s8
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s17, 31
	s_cselect_b32 s8, 0xffffffe0, 0
	s_cselect_b32 s13, s11, s7
	s_cselect_b32 s7, s7, s12
	s_cselect_b32 s11, s19, s11
	s_cselect_b32 s12, s9, s19
	s_cselect_b32 s9, s10, s9
	s_add_i32 s8, s8, s17
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s8, 31
	s_cselect_b32 s10, 0xffffffe0, 0
	s_cselect_b32 s7, s13, s7
	s_cselect_b32 s13, s11, s13
	s_cselect_b32 s11, s12, s11
	s_cselect_b32 s9, s9, s12
	s_add_i32 s10, s10, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_sub_i32 s8, 32, s10
	s_cmp_eq_u32 s10, 0
	v_mov_b32_e32 v6, s8
	s_cselect_b32 s10, -1, 0
	v_alignbit_b32 v7, s7, s13, v6
	v_alignbit_b32 v8, s13, s11, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s8, v7
	v_cndmask_b32_e64 v7, v8, s13, s10
	s_delay_alu instid0(VALU_DEP_2)
	s_cselect_b32 s7, s7, s8
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_alignbit_b32 v8, s7, v7, 30
	s_bfe_u32 s12, s7, 0x1001d
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s8, 0, s12
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_xor_b32_e32 v8, s8, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v9, v8
	v_min_u32_e32 v9, 32, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v11, 23, v9
	v_alignbit_b32 v6, s11, s9, v6
	v_sub_nc_u32_e32 v10, 31, v9
	v_cndmask_b32_e64 v6, v6, s11, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v7, v7, v6, 30
	v_alignbit_b32 v6, v6, s9, 30
	v_xor_b32_e32 v7, s8, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_xor_b32_e32 v6, s8, v6
	s_lshr_b32 s8, s7, 29
	s_lshr_b32 s7, s7, 30
	s_lshl_b32 s8, s8, 31
	v_alignbit_b32 v8, v8, v7, v10
	v_alignbit_b32 v6, v7, v6, v10
	s_or_b32 s9, s8, 0.5
	s_add_i32 s7, s12, s7
	v_sub_nc_u32_e32 v11, s9, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v7, v8, v6, 9
	v_clz_i32_u32_e32 v10, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v10, 32, v10
	v_sub_nc_u32_e32 v12, 31, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v6, v7, v6, v12
	v_lshrrev_b32_e32 v7, 9, v8
	v_lshrrev_b32_e32 v6, 9, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_or_b32_e32 v7, v7, v11
	v_add_nc_u32_e32 v9, v10, v9
	v_lshlrev_b32_e32 v8, 23, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v6, v6, v8
	v_mul_f32_e32 v8, 0x3fc90fda, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, 0x33000000, v6
	v_fma_f32 v9, 0x3fc90fda, v7, -v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v6, s8, v6
	v_fmamk_f32 v7, v7, 0x33a22168, v9
	s_mov_b32 s8, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v7, 0x3fc90fda, v6
	v_add_f32_e32 v6, v8, v7
	s_branch .LBB0_15
.LBB0_14:
	s_mov_b32 s8, -1
                                        ; implicit-def: $sgpr7
                                        ; implicit-def: $vgpr6
.LBB0_15:                               ; %Flow69
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s8
	s_cbranch_vccnz .LBB0_17
; %bb.16:
	v_rndne_f32_e32 v7, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v6, 0xbfc90fda, v7, |s6|
	v_fmamk_f32 v6, v7, 0xb3a22168, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_fmamk_f32 v6, v7, 0xa7c234c4, v6
	v_cvt_i32_f32_e32 v7, v7
	s_branch .LBB0_18
.LBB0_17:
	v_mov_b32_e32 v7, s7
.LBB0_18:                               ; %_ZL4cosff.exit35
	s_mov_b32 vcc_lo, s2
	s_cbranch_vccz .LBB0_20
; %bb.19:
	s_and_b32 s2, s3, 0x7fffff
	s_lshr_b32 s7, s3, 23
	s_bitset1_b32 s2, 23
	s_addk_i32 s7, 0xff88
	s_mul_hi_u32 s8, s2, 0xfe5163ab
	s_mul_i32 s9, s2, 0x3c439041
	s_mul_hi_u32 s10, s2, 0x3c439041
	s_add_u32 s8, s8, s9
	s_addc_u32 s9, 0, s10
	s_mul_i32 s10, s2, 0xdb629599
	s_mul_hi_u32 s11, s2, 0xdb629599
	s_add_u32 s9, s9, s10
	s_addc_u32 s10, 0, s11
	s_mul_i32 s11, s2, 0xf534ddc0
	s_mul_hi_u32 s12, s2, 0xf534ddc0
	s_add_u32 s10, s10, s11
	s_addc_u32 s11, 0, s12
	s_mul_i32 s12, s2, 0xfc2757d1
	s_mul_hi_u32 s13, s2, 0xfc2757d1
	s_add_u32 s11, s11, s12
	s_addc_u32 s12, 0, s13
	s_mul_i32 s13, s2, 0x4e441529
	s_mul_hi_u32 s16, s2, 0x4e441529
	s_add_u32 s12, s12, s13
	s_addc_u32 s13, 0, s16
	s_cmp_gt_u32 s7, 63
	s_mul_i32 s16, s2, 0xfe5163ab
	s_mul_hi_u32 s17, s2, 0xa2f9836e
	s_mul_i32 s2, s2, 0xa2f9836e
	s_cselect_b32 s18, s9, s11
	s_cselect_b32 s8, s8, s10
	s_cselect_b32 s9, s16, s9
	s_add_u32 s2, s13, s2
	s_addc_u32 s13, 0, s17
	s_cmp_gt_u32 s7, 63
	s_cselect_b32 s16, 0xffffffc0, 0
	s_cselect_b32 s10, s10, s12
	s_cselect_b32 s2, s11, s2
	s_cselect_b32 s11, s12, s13
	s_add_i32 s16, s16, s7
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s16, 31
	s_cselect_b32 s7, 0xffffffe0, 0
	s_cselect_b32 s12, s10, s2
	s_cselect_b32 s2, s2, s11
	s_cselect_b32 s10, s18, s10
	s_cselect_b32 s11, s8, s18
	s_cselect_b32 s8, s9, s8
	s_add_i32 s7, s7, s16
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s7, 31
	s_cselect_b32 s9, 0xffffffe0, 0
	s_cselect_b32 s2, s12, s2
	s_cselect_b32 s12, s10, s12
	s_cselect_b32 s10, s11, s10
	s_cselect_b32 s8, s8, s11
	s_add_i32 s9, s9, s7
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_sub_i32 s7, 32, s9
	s_cmp_eq_u32 s9, 0
	v_mov_b32_e32 v8, s7
	s_cselect_b32 s9, -1, 0
	v_alignbit_b32 v9, s2, s12, v8
	v_alignbit_b32 v10, s12, s10, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_readfirstlane_b32 s7, v9
	v_cndmask_b32_e64 v9, v10, s12, s9
	s_delay_alu instid0(VALU_DEP_2)
	s_cselect_b32 s2, s2, s7
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_alignbit_b32 v10, s2, v9, 30
	s_bfe_u32 s7, s2, 0x1001d
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s11, 0, s7
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_xor_b32_e32 v10, s11, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_clz_i32_u32_e32 v11, v10
	v_min_u32_e32 v11, 32, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v13, 23, v11
	v_alignbit_b32 v8, s10, s8, v8
	v_sub_nc_u32_e32 v12, 31, v11
	v_cndmask_b32_e64 v8, v8, s10, s9
	s_delay_alu instid0(VALU_DEP_1)
	v_alignbit_b32 v9, v9, v8, 30
	v_alignbit_b32 v8, v8, s8, 30
	s_lshr_b32 s8, s2, 29
	s_lshr_b32 s2, s2, 30
	s_lshl_b32 s8, s8, 31
	v_xor_b32_e32 v9, s11, v9
	v_xor_b32_e32 v8, s11, v8
	s_or_b32 s9, s8, 0.5
	s_add_i32 s7, s7, s2
	v_sub_nc_u32_e32 v13, s9, v13
	v_alignbit_b32 v10, v10, v9, v12
	v_alignbit_b32 v8, v9, v8, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v9, v10, v8, 9
	v_clz_i32_u32_e32 v12, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_min_u32_e32 v12, 32, v12
	v_sub_nc_u32_e32 v14, 31, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_alignbit_b32 v8, v9, v8, v14
	v_lshrrev_b32_e32 v9, 9, v10
	v_lshrrev_b32_e32 v8, 9, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_or_b32_e32 v9, v9, v13
	v_add_nc_u32_e32 v11, v12, v11
	v_lshlrev_b32_e32 v10, 23, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v8, v8, v10
	v_mul_f32_e32 v10, 0x3fc90fda, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v8, 0x33000000, v8
	v_fma_f32 v11, 0x3fc90fda, v9, -v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_or_b32_e32 v8, s8, v8
	v_fmamk_f32 v9, v9, 0x33a22168, v11
	s_mov_b32 s8, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v9, 0x3fc90fda, v8
	v_add_f32_e32 v8, v10, v9
	s_branch .LBB0_21
.LBB0_20:
	s_mov_b32 s8, -1
                                        ; implicit-def: $sgpr7
                                        ; implicit-def: $vgpr8
.LBB0_21:                               ; %Flow
	s_load_b32 s2, s[0:1], 0x2c
	s_and_not1_b32 vcc_lo, exec_lo, s8
	s_cbranch_vccnz .LBB0_23
; %bb.22:
	v_rndne_f32_e32 v3, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v8, 0xbfc90fda, v3, |s6|
	v_fmamk_f32 v8, v3, 0xb3a22168, v8
	s_delay_alu instid0(VALU_DEP_1)
	v_fmamk_f32 v8, v3, 0xa7c234c4, v8
	v_cvt_i32_f32_e32 v3, v3
	s_branch .LBB0_24
.LBB0_23:
	v_mov_b32_e32 v3, s7
.LBB0_24:                               ; %_ZL4sinff.exit43
	v_mul_f32_e32 v12, v6, v6
	s_mov_b32 s8, 0x37d75334
	s_mov_b32 s9, 0xb94c1982
	v_cvt_f32_i32_e32 v13, s5
	v_lshlrev_b32_e32 v14, 30, v7
	v_dual_fmaak_f32 v10, s8, v12, 0xbab64f3b :: v_dual_and_b32 v7, 1, v7
	v_mul_f32_e32 v17, v1, v1
	v_fmaak_f32 v11, s9, v12, 0x3c0881c4
	v_bfe_u32 v9, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s7, s2, 16
	v_fmaak_f32 v10, v12, v10, 0x3d2aabf7
	v_cmp_eq_u32_e32 vcc_lo, 0, v7
	v_fmaak_f32 v11, v12, v11, 0xbe2aaa9d
	v_and_b32_e32 v0, 0x3ff, v0
	s_and_b32 s2, s2, 0xffff
	v_dual_fmaak_f32 v15, v12, v10, 0xbf000004 :: v_dual_lshlrev_b32 v18, 30, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_f32_e32 v16, v12, v11
	v_mad_u64_u32 v[10:11], null, s15, s7, v[9:10]
	v_lshlrev_b32_e32 v11, 30, v2
	v_fma_f32 v9, v12, v15, 1.0
	s_delay_alu instid0(VALU_DEP_4)
	v_fmac_f32_e32 v6, v6, v16
	v_and_b32_e32 v12, 0x80000000, v14
	v_fmaak_f32 v15, s9, v17, 0x3c0881c4
	v_and_b32_e32 v5, 1, v5
	v_cvt_f32_u32_e32 v16, v10
	v_cndmask_b32_e64 v9, -v6, v9, vcc_lo
	v_mad_u64_u32 v[6:7], null, s14, s2, v[0:1]
	v_cmp_class_f32_e64 vcc_lo, s6, 0x1f8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v16, -0.5, v13
	v_xor_b32_e32 v0, v12, v9
	v_mul_f32_e32 v12, v4, v4
	v_fmaak_f32 v14, s8, v17, 0xbab64f3b
	v_dual_fmaak_f32 v9, v17, v15, 0xbe2aaa9d :: v_dual_and_b32 v2, 1, v2
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v0, 0x7fc00000, v0, vcc_lo
	v_fmaak_f32 v19, s8, v12, 0xbab64f3b
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_dual_fmaak_f32 v7, v17, v14, 0x3d2aabf7 :: v_dual_mul_f32 v14, v8, v8
	v_fmaak_f32 v15, s9, v12, 0x3c0881c4
	v_mul_f32_e32 v9, v17, v9
	v_fmaak_f32 v19, v12, v19, 0x3d2aabf7
	s_delay_alu instid0(VALU_DEP_4)
	v_fmaak_f32 v7, v17, v7, 0xbf000004
	v_fmaak_f32 v20, s9, v14, 0x3c0881c4
	v_fmaak_f32 v15, v12, v15, 0xbe2aaa9d
	v_cmp_eq_u32_e64 s2, 0, v5
	v_fmaak_f32 v19, v12, v19, 0xbf000004
	v_fma_f32 v7, v17, v7, 1.0
	v_fmaak_f32 v20, v14, v20, 0xbe2aaa9d
	v_fmaak_f32 v21, s8, v14, 0xbab64f3b
	v_mul_f32_e32 v15, v12, v15
	v_fma_f32 v12, v12, v19, 1.0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_fmac_f32 v1, v1, v9 :: v_dual_mul_f32 v20, v14, v20
	v_dual_fmaak_f32 v21, v14, v21, 0x3d2aabf7 :: v_dual_fmac_f32 v4, v4, v15
	v_lshlrev_b32_e32 v15, 30, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_fmac_f32 v8, v8, v20 :: v_dual_and_b32 v3, 1, v3
	v_fmaak_f32 v21, v14, v21, 0xbf000004
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v4, v12, v4, s2
	v_and_b32_e32 v15, 0x80000000, v15
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_eq_u32_e64 s2, 0, v3
	v_fma_f32 v14, v14, v21, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v3, v14, v8, s2
	v_xor_b32_e32 v8, s3, v15
	v_cmp_eq_u32_e64 s2, 0, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v8, v3
	v_cndmask_b32_e64 v1, -v1, v7, s2
	v_cvt_f32_i32_e32 v7, s4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v3, s6, v3
	v_cndmask_b32_e32 v3, 0x7fc00000, v3, vcc_lo
	v_and_b32_e32 v18, 0x80000000, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v5, s3, v18
	v_xor_b32_e32 v4, v5, v4
	v_and_b32_e32 v5, 0x80000000, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v2, s6, v4
	v_cvt_f32_u32_e32 v4, v6
	v_xor_b32_e32 v1, v5, v1
	s_clause 0x1
	s_load_b64 s[6:7], s[0:1], 0x18
	s_load_b64 s[0:1], s[0:1], 0x0
	v_cndmask_b32_e32 v2, 0x7fc00000, v2, vcc_lo
	v_dual_fmac_f32 v4, -0.5, v7 :: v_dual_cndmask_b32 v1, 0x7fc00000, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v2, v16, v2
	v_mul_f32_e32 v3, v4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v1, v4, v1, -v2
	v_fmac_f32_e32 v3, v16, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_scale_f32 v0, null, v7, v7, v1
	v_div_scale_f32 v2, null, v13, v13, v3
	s_waitcnt lgkmcnt(0)
	s_clause 0x1
	s_load_b128 s[16:19], s[6:7], 0x30
	s_load_b32 s3, s[6:7], 0x28
	v_rcp_f32_e32 v4, v0
	v_rcp_f32_e32 v5, v2
	s_load_b256 s[8:15], s[6:7], 0x0
	s_waitcnt_depctr 0xfff
	v_fma_f32 v8, -v0, v4, 1.0
	v_fma_f32 v9, -v2, v5, 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v8, v4
	v_div_scale_f32 v8, vcc_lo, v1, v7, v1
	s_waitcnt lgkmcnt(0)
	s_bitcmp0_b32 s18, 20
	v_mul_f32_e32 v11, v8, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v14, -v0, v11, v8
	v_fmac_f32_e32 v11, v14, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v0, -v0, v11, v8
	v_div_fmas_f32 v0, v0, v4, v11
	v_cvt_f32_u32_e32 v4, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v0, v0, v7, v1
	v_dual_add_f32 v0, 0.5, v0 :: v_dual_fmac_f32 v5, v9, v5
	v_div_scale_f32 v9, s2, v3, v13, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_mov_b32 vcc_lo, s2
	v_mul_f32_e32 v12, v9, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v15, -v2, v12, v9
	v_fmac_f32_e32 v12, v15, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v2, -v2, v12, v9
	v_div_fmas_f32 v2, v2, v5, v12
	s_cselect_b32 vcc_lo, -1, 0
	s_bitcmp0_b32 s16, 15
	s_cselect_b32 s2, -1, 0
	s_bfe_u32 s5, s10, 0xe000e
	v_div_fixup_f32 v2, v2, v13, v3
	s_add_i32 s5, s5, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f32_u32_e32 v3, s5
	v_add_f32_e32 v2, 0.5, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v1, 1.0, v3, s2
	v_cndmask_b32_e64 v3, 1.0, v4, s2
	v_rcp_f32_e32 v4, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v1, v1, v2
	v_rcp_f32_e32 v5, v3
	v_mul_f32_e32 v3, v0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_floor_f32_e32 v1, v1
	v_floor_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, v4, v1
	v_mul_f32_e32 v3, v5, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v1, v2, v1 :: v_dual_cndmask_b32 v0, v0, v3
	image_sample_lz v2, v[0:1], s[8:15], s[16:19] dmask:0x1 dim:SQ_RSRC_IMG_2D
	v_mad_u64_u32 v[0:1], null, v10, s4, v[6:7]
	v_mov_b32_e32 v1, 0
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[0:1], v2, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15transformKernelPfiifP13__hip_texture
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 22
		.amdhsa_next_free_sgpr 20
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z15transformKernelPfiifP13__hip_texture, .Lfunc_end0-_Z15transformKernelPfiifP13__hip_texture
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4008
; NumSgprs: 22
; NumVgprs: 22
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 22
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z15transformKernelPfiifP13__hip_texture
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z15transformKernelPfiifP13__hip_texture.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     22
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
