/*
    Copyright 2018 - The OPRECOMP Project Consortium, Alma Mater Studiorum
    Università di Bologna. All rights reserved.

    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.
    You may obtain a copy of the License at

        http://www.apache.org/licenses/LICENSE-2.0

    Unless required by applicable law or agreed to in writing, software
    distributed under the License is distributed on an "AS IS" BASIS,
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
*/

#include <stdlib.h>
#include <stdio.h>
#include <math.h>

#ifndef FLOAT
#define FLOAT double
#endif

#define     KERNEL_RADIUS 8
#define      KERNEL_W (2 * KERNEL_RADIUS + 1)
#define      DATA_W  4096
#define      DATA_H  4096
const int   DATA_SIZE = DATA_W * DATA_H * sizeof(FLOAT);
const int KERNEL_SIZE = KERNEL_W * sizeof(FLOAT);

__device__ __constant__ FLOAT d_Kernel[KERNEL_W];

__global__
void convolutionRowGPU(
    FLOAT *d_Result,
    FLOAT *d_Data,
    int dataW,
    int dataH,
    int kernelR
){
    int index=blockIdx.x * blockDim.x +threadIdx.x; // global thread id
    if(index >= dataW*dataH) return;
    int y = index/dataW;
    int x = index - y*dataW;
    int k, d;
    FLOAT sum;


    sum = FLOAT(0.0f);
    for(k = -kernelR; k <= kernelR; k++){
        d = x + k;
        if(d >= 0 && d < dataW)
            sum += d_Data[y * dataW + d] * d_Kernel[kernelR - k];
    }
    d_Result[y * dataW + x] = sum;
        
}


__global__
void convolutionColumnGPU(
    FLOAT *d_Result,
    FLOAT *d_Data,
    int dataW,
    int dataH,
    int kernelR
){
    int index=blockIdx.x * blockDim.x +threadIdx.x; // global thread id
    if(index >= dataW*dataH) return;
    int y = index/dataW;
    int x = index - y*dataW;
    int k, d;
    FLOAT sum;


    sum = FLOAT(0.0f);
    for(k = -kernelR; k <= kernelR; k++){
        d = y + k;
        if(d >= 0 && d < dataH)
            sum += d_Data[d * dataW + x] * d_Kernel[kernelR - k];
    }
    d_Result[y * dataW + x] = sum;
        
}



int main(int argc, char **argv){
    int i;
    
    FLOAT
        *h_Kernel,
        *h_DataA;

    FLOAT
        *d_DataA,
        *d_DataB;
       

    h_Kernel    = (FLOAT *)malloc(KERNEL_SIZE);
    h_DataA     = (FLOAT *)malloc(DATA_SIZE);  

    cudaMalloc( (void **)&d_DataA, DATA_SIZE);
    cudaMalloc( (void **)&d_DataB, DATA_SIZE);   

    FLOAT kernelSum = 0;
    for(i = 0; i < KERNEL_W; i++){
        FLOAT dist = (FLOAT)(i - KERNEL_RADIUS) / (FLOAT)KERNEL_RADIUS;
        h_Kernel[i] = expf(- dist * dist / 2);
        kernelSum += h_Kernel[i];
    }
    for(i = 0; i < KERNEL_W; i++)
        h_Kernel[i] /= kernelSum;

    srand(5497);
    for(i = 0; i < DATA_W * DATA_H; i++)
        h_DataA[i] = (FLOAT)rand() / (FLOAT)RAND_MAX;

    cudaMemcpyToSymbol(d_Kernel, h_Kernel, KERNEL_SIZE);
    cudaMemcpy(d_DataA, h_DataA, DATA_SIZE, cudaMemcpyHostToDevice);  

  int blockSize=256;
  int numBlocks = ((DATA_W * DATA_H)+blockSize-1)/blockSize;    
      //for(i = 0; i < DATA_W * DATA_H; i++)
       // printf("%.15f,", h_DataA[i]);

    convolutionRowGPU<<<numBlocks, blockSize>>>(
        d_DataB,
        d_DataA,
        DATA_W,
        DATA_H,
        KERNEL_RADIUS
    );
    //cudaMemcpy(h_DataA, d_DataB, DATA_SIZE, cudaMemcpyDeviceToHost);

    //for(i = 0; i < DATA_W * DATA_H; i++)
    //    printf("%.15f,", h_DataA[i]);
    convolutionColumnGPU<<<numBlocks, blockSize>>>(
        d_DataA,
        d_DataB,
        DATA_W,
        DATA_H,
        KERNEL_RADIUS
    );

    cudaMemcpy(h_DataA, d_DataA, DATA_SIZE, cudaMemcpyDeviceToHost);

    for(i = 0; i < DATA_W * DATA_H; i++)
        printf("%.15f,", h_DataA[i]);
    return 0;
}

	.file	"tmpxft_002dbe32_00000000-6_conv.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii
	.type	_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii, @function
_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%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	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%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	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z17convolutionRowGPUPdS_iii(%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
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii, .-_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii
	.globl	_Z17convolutionRowGPUPdS_iii
	.type	_Z17convolutionRowGPUPdS_iii, @function
_Z17convolutionRowGPUPdS_iii:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii
	.cfi_endproc
.LFE2053:
	.size	_Z17convolutionRowGPUPdS_iii, .-_Z17convolutionRowGPUPdS_iii
	.globl	_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii
	.type	_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii, @function
_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movl	%r8d, 4(%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	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%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	.L8
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z20convolutionColumnGPUPdS_iii(%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
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii, .-_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii
	.globl	_Z20convolutionColumnGPUPdS_iii
	.type	_Z20convolutionColumnGPUPdS_iii, @function
_Z20convolutionColumnGPUPdS_iii:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii
	.cfi_endproc
.LFE2055:
	.size	_Z20convolutionColumnGPUPdS_iii, .-_Z20convolutionColumnGPUPdS_iii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC5:
	.string	"%.15f,"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$136, %edi
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$64, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movl	$134217728, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	leaq	16(%rsp), %rdi
	movl	$134217728, %esi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	leaq	24(%rsp), %rdi
	movl	$134217728, %esi
	call	cudaMalloc@PLT
	xorps	%xmm2, %xmm2
.L14:
	leal	-8(%r12), %eax
	movsd	%xmm2, 8(%rsp)
	cvtsi2sdl	%eax, %xmm1
	mulsd	.LC1(%rip), %xmm1
	movaps	%xmm1, %xmm0
	xorps	.LC2(%rip), %xmm0
	mulsd	%xmm1, %xmm0
	mulsd	.LC3(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	call	expf@PLT
	movsd	8(%rsp), %xmm2
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, 0(%rbp,%r12,8)
	incq	%r12
	addsd	%xmm0, %xmm2
	cmpq	$17, %r12
	jne	.L14
	xorl	%eax, %eax
.L15:
	movsd	0(%rbp,%rax,8), %xmm0
	divsd	%xmm2, %xmm0
	movsd	%xmm0, 0(%rbp,%rax,8)
	incq	%rax
	cmpq	$17, %rax
	jne	.L15
	movl	$5497, %edi
	xorl	%r12d, %r12d
	call	srand@PLT
.L16:
	call	rand@PLT
	cvtsi2sdl	%eax, %xmm0
	divsd	.LC4(%rip), %xmm0
	movsd	%xmm0, (%rbx,%r12,8)
	incq	%r12
	cmpq	$16777216, %r12
	jne	.L16
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$136, %edx
	movq	%rbp, %rsi
	leaq	_ZL8d_Kernel(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$134217728, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$65537, %edi
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	salq	$16, %rdi
	movq	%rdx, 44(%rsp)
	movl	$1, 52(%rsp)
	movq	%rdi, 32(%rsp)
	movl	$1, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L17
	movq	16(%rsp), %rsi
	movq	24(%rsp), %rdi
	movl	$8, %r8d
	movl	$4096, %ecx
	movl	$4096, %edx
	call	_Z42__device_stub__Z17convolutionRowGPUPdS_iiiPdS_iii
.L17:
	movl	$16777217, %edx
	movl	$65537, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	salq	$16, %rdi
	movl	$1, %ecx
	movl	$1, %esi
	movq	%rdx, 44(%rsp)
	movl	$1, 52(%rsp)
	movq	%rdi, 32(%rsp)
	movl	$1, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$8, %r8d
	movl	$4096, %ecx
	movl	$4096, %edx
	call	_Z45__device_stub__Z20convolutionColumnGPUPdS_iiiPdS_iii
.L18:
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	xorl	%ebp, %ebp
	movl	$134217728, %edx
	leaq	.LC5(%rip), %r12
	call	cudaMemcpy@PLT
.L19:
	movsd	(%rbx,%rbp,8), %xmm0
	movq	%r12, %rsi
	movl	$2, %edi
	movb	$1, %al
	incq	%rbp
	call	__printf_chk@PLT
	cmpq	$16777216, %rbp
	jne	.L19
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$64, %rsp
	.cfi_def_cfa_offset 32
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC6:
	.string	"_Z20convolutionColumnGPUPdS_iii"
.LC7:
	.string	"_Z17convolutionRowGPUPdS_iii"
.LC8:
	.string	"d_Kernel"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -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	.LC6(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z20convolutionColumnGPUPdS_iii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z17convolutionRowGPUPdS_iii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC8(%rip), %rdx
	movl	$136, %r9d
	leaq	_ZL8d_Kernel(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.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
	.local	_ZL8d_Kernel
	.comm	_ZL8d_Kernel,136,32
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC1:
	.long	0
	.long	1069547520
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC2:
	.long	0
	.long	-2147483648
	.long	0
	.long	0
	.section	.rodata.cst8
	.align 8
.LC3:
	.long	0
	.long	1071644672
	.align 8
.LC4:
	.long	-4194304
	.long	1105199103
	.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:

/*
    Copyright 2018 - The OPRECOMP Project Consortium, Alma Mater Studiorum
    Università di Bologna. All rights reserved.

    Licensed under the Apache License, Version 2.0 (the "License");
    you may not use this file except in compliance with the License.
    You may obtain a copy of the License at

        http://www.apache.org/licenses/LICENSE-2.0

    Unless required by applicable law or agreed to in writing, software
    distributed under the License is distributed on an "AS IS" BASIS,
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
*/


#include <hip/hip_runtime.h>
#include <stdlib.h>
#include <stdio.h>
#include <math.h>

#ifndef FLOAT
#define FLOAT double
#endif

#define     KERNEL_RADIUS 8
#define      KERNEL_W (2 * KERNEL_RADIUS + 1)
#define      DATA_W  4096
#define      DATA_H  4096
const int   DATA_SIZE = DATA_W * DATA_H * sizeof(FLOAT);
const int KERNEL_SIZE = KERNEL_W * sizeof(FLOAT);

__device__ __constant__ FLOAT d_Kernel[KERNEL_W];

__global__
void convolutionRowGPU(
    FLOAT *d_Result,
    FLOAT *d_Data,
    int dataW,
    int dataH,
    int kernelR
){
    int index=blockIdx.x * blockDim.x +threadIdx.x; // global thread id
    if(index >= dataW*dataH) return;
    int y = index/dataW;
    int x = index - y*dataW;
    int k, d;
    FLOAT sum;


    sum = FLOAT(0.0f);
    for(k = -kernelR; k <= kernelR; k++){
        d = x + k;
        if(d >= 0 && d < dataW)
            sum += d_Data[y * dataW + d] * d_Kernel[kernelR - k];
    }
    d_Result[y * dataW + x] = sum;
        
}


__global__
void convolutionColumnGPU(
    FLOAT *d_Result,
    FLOAT *d_Data,
    int dataW,
    int dataH,
    int kernelR
){
    int index=blockIdx.x * blockDim.x +threadIdx.x; // global thread id
    if(index >= dataW*dataH) return;
    int y = index/dataW;
    int x = index - y*dataW;
    int k, d;
    FLOAT sum;


    sum = FLOAT(0.0f);
    for(k = -kernelR; k <= kernelR; k++){
        d = y + k;
        if(d >= 0 && d < dataH)
            sum += d_Data[d * dataW + x] * d_Kernel[kernelR - k];
    }
    d_Result[y * dataW + x] = sum;
        
}



int main(int argc, char **argv){
    int i;
    
    FLOAT
        *h_Kernel,
        *h_DataA;

    FLOAT
        *d_DataA,
        *d_DataB;
       

    h_Kernel    = (FLOAT *)malloc(KERNEL_SIZE);
    h_DataA     = (FLOAT *)malloc(DATA_SIZE);  

    hipMalloc( (void **)&d_DataA, DATA_SIZE);
    hipMalloc( (void **)&d_DataB, DATA_SIZE);   

    FLOAT kernelSum = 0;
    for(i = 0; i < KERNEL_W; i++){
        FLOAT dist = (FLOAT)(i - KERNEL_RADIUS) / (FLOAT)KERNEL_RADIUS;
        h_Kernel[i] = expf(- dist * dist / 2);
        kernelSum += h_Kernel[i];
    }
    for(i = 0; i < KERNEL_W; i++)
        h_Kernel[i] /= kernelSum;

    srand(5497);
    for(i = 0; i < DATA_W * DATA_H; i++)
        h_DataA[i] = (FLOAT)rand() / (FLOAT)RAND_MAX;

    hipMemcpyToSymbol(HIP_SYMBOL(d_Kernel), h_Kernel, KERNEL_SIZE);
    hipMemcpy(d_DataA, h_DataA, DATA_SIZE, hipMemcpyHostToDevice);  

  int blockSize=256;
  int numBlocks = ((DATA_W * DATA_H)+blockSize-1)/blockSize;    
      //for(i = 0; i < DATA_W * DATA_H; i++)
       // printf("%.15f,", h_DataA[i]);

    convolutionRowGPU<<<numBlocks, blockSize>>>(
        d_DataB,
        d_DataA,
        DATA_W,
        DATA_H,
        KERNEL_RADIUS
    );
    //cudaMemcpy(h_DataA, d_DataB, DATA_SIZE, cudaMemcpyDeviceToHost);

    //for(i = 0; i < DATA_W * DATA_H; i++)
    //    printf("%.15f,", h_DataA[i]);
    convolutionColumnGPU<<<numBlocks, blockSize>>>(
        d_DataA,
        d_DataB,
        DATA_W,
        DATA_H,
        KERNEL_RADIUS
    );

    hipMemcpy(h_DataA, d_DataA, DATA_SIZE, hipMemcpyDeviceToHost);

    for(i = 0; i < DATA_W * DATA_H; i++)
        printf("%.15f,", h_DataA[i]);
    return 0;
}

	.text
	.file	"conv.hip"
	.globl	_Z32__device_stub__convolutionRowGPUPdS_iii # -- Begin function _Z32__device_stub__convolutionRowGPUPdS_iii
	.type	_Z32__device_stub__convolutionRowGPUPdS_iii,@function
_Z32__device_stub__convolutionRowGPUPdS_iii: # @_Z32__device_stub__convolutionRowGPUPdS_iii
	.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	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 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	$_Z17convolutionRowGPUPdS_iii, %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	_Z32__device_stub__convolutionRowGPUPdS_iii, .Lfunc_end0-_Z32__device_stub__convolutionRowGPUPdS_iii
	.cfi_endproc
                                        # -- End function
	.globl	_Z35__device_stub__convolutionColumnGPUPdS_iii # -- Begin function _Z35__device_stub__convolutionColumnGPUPdS_iii
	.type	_Z35__device_stub__convolutionColumnGPUPdS_iii,@function
_Z35__device_stub__convolutionColumnGPUPdS_iii: # @_Z35__device_stub__convolutionColumnGPUPdS_iii
	.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	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	4(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 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	$_Z20convolutionColumnGPUPdS_iii, %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_end1:
	.size	_Z35__device_stub__convolutionColumnGPUPdS_iii, .Lfunc_end1-_Z35__device_stub__convolutionColumnGPUPdS_iii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI2_0:
	.quad	0x3fc0000000000000              # double 0.125
.LCPI2_1:
	.quad	0xbfe0000000000000              # double -0.5
.LCPI2_2:
	.quad	0x41dfffffffc00000              # double 2147483647
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$136, %edi
	callq	malloc
	movq	%rax, %r14
	movl	$134217728, %edi                # imm = 0x8000000
	callq	malloc
	movq	%rax, %rbx
	leaq	8(%rsp), %rdi
	movl	$134217728, %esi                # imm = 0x8000000
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$134217728, %esi                # imm = 0x8000000
	callq	hipMalloc
	xorpd	%xmm1, %xmm1
	xorl	%r15d, %r15d
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movsd	%xmm1, 24(%rsp)                 # 8-byte Spill
	leal	-8(%r15), %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	mulsd	.LCPI2_0(%rip), %xmm0
	mulsd	%xmm0, %xmm0
	mulsd	.LCPI2_1(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	callq	expf
	movsd	24(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, (%r14,%r15,8)
	addsd	%xmm0, %xmm1
	incq	%r15
	cmpq	$17, %r15
	jne	.LBB2_1
# %bb.2:                                # %.preheader.preheader
	xorl	%eax, %eax
.LBB2_3:                                # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movsd	(%r14,%rax,8), %xmm0            # xmm0 = mem[0],zero
	divsd	%xmm1, %xmm0
	movsd	%xmm0, (%r14,%rax,8)
	incq	%rax
	cmpq	$17, %rax
	jne	.LBB2_3
# %bb.4:
	movl	$5497, %edi                     # imm = 0x1579
	callq	srand
	xorl	%r15d, %r15d
.LBB2_5:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	divsd	.LCPI2_2(%rip), %xmm0
	movsd	%xmm0, (%rbx,%r15,8)
	incq	%r15
	cmpq	$16777216, %r15                 # imm = 0x1000000
	jne	.LBB2_5
# %bb.6:
	movabsq	$4294967552, %r15               # imm = 0x100000100
	movl	$d_Kernel, %edi
	movl	$136, %edx
	movq	%r14, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movq	8(%rsp), %rdi
	movl	$134217728, %edx                # imm = 0x8000000
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	65280(%r15), %r14
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_8
# %bb.7:
	movq	16(%rsp), %rdi
	movq	8(%rsp), %rsi
	movl	$4096, %edx                     # imm = 0x1000
	movl	$4096, %ecx                     # imm = 0x1000
	movl	$8, %r8d
	callq	_Z32__device_stub__convolutionRowGPUPdS_iii
.LBB2_8:
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_10
# %bb.9:
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	$4096, %edx                     # imm = 0x1000
	movl	$4096, %ecx                     # imm = 0x1000
	movl	$8, %r8d
	callq	_Z35__device_stub__convolutionColumnGPUPdS_iii
.LBB2_10:
	movq	8(%rsp), %rsi
	movl	$134217728, %edx                # imm = 0x8000000
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%r14d, %r14d
.LBB2_11:                               # =>This Inner Loop Header: Depth=1
	movsd	(%rbx,%r14,8), %xmm0            # xmm0 = mem[0],zero
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	incq	%r14
	cmpq	$16777216, %r14                 # imm = 0x1000000
	jne	.LBB2_11
# %bb.12:
	xorl	%eax, %eax
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17convolutionRowGPUPdS_iii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z20convolutionColumnGPUPdS_iii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$d_Kernel, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$136, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.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	d_Kernel,@object                # @d_Kernel
	.local	d_Kernel
	.comm	d_Kernel,136,16
	.type	_Z17convolutionRowGPUPdS_iii,@object # @_Z17convolutionRowGPUPdS_iii
	.section	.rodata,"a",@progbits
	.globl	_Z17convolutionRowGPUPdS_iii
	.p2align	3, 0x0
_Z17convolutionRowGPUPdS_iii:
	.quad	_Z32__device_stub__convolutionRowGPUPdS_iii
	.size	_Z17convolutionRowGPUPdS_iii, 8

	.type	_Z20convolutionColumnGPUPdS_iii,@object # @_Z20convolutionColumnGPUPdS_iii
	.globl	_Z20convolutionColumnGPUPdS_iii
	.p2align	3, 0x0
_Z20convolutionColumnGPUPdS_iii:
	.quad	_Z35__device_stub__convolutionColumnGPUPdS_iii
	.size	_Z20convolutionColumnGPUPdS_iii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%.15f,"
	.size	.L.str, 7

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z17convolutionRowGPUPdS_iii"
	.size	.L__unnamed_1, 29

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z20convolutionColumnGPUPdS_iii"
	.size	.L__unnamed_2, 32

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"d_Kernel"
	.size	.L__unnamed_3, 9

	.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 _Z32__device_stub__convolutionRowGPUPdS_iii
	.addrsig_sym _Z35__device_stub__convolutionColumnGPUPdS_iii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym d_Kernel
	.addrsig_sym _Z17convolutionRowGPUPdS_iii
	.addrsig_sym _Z20convolutionColumnGPUPdS_iii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17convolutionRowGPUPdS_iii ; -- Begin function _Z17convolutionRowGPUPdS_iii
	.globl	_Z17convolutionRowGPUPdS_iii
	.p2align	8
	.type	_Z17convolutionRowGPUPdS_iii,@function
_Z17convolutionRowGPUPdS_iii:           ; @_Z17convolutionRowGPUPdS_iii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b128 s[4:7], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mul_i32 s2, s5, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_7
; %bb.1:
	s_load_b128 s[8:11], s[0:1], 0x0
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	v_ashrrev_i32_e32 v2, 31, v1
	s_cmp_lt_i32 s6, 0
	s_cbranch_scc1 .LBB0_6
; %bb.2:                                ; %.lr.ph.preheader
	s_ashr_i32 s0, s4, 31
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v4, v1, v2
	s_add_i32 s1, s4, s0
	v_subrev_nc_u32_e32 v5, s6, v1
	s_xor_b32 s0, s1, s0
	s_lshl_b32 s2, s6, 1
	v_cvt_f32_u32_e32 v0, s0
	s_sub_i32 s1, 0, s0
	v_xor_b32_e32 v4, v4, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v0, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, s1, v0
	v_mul_hi_u32 v3, v0, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v3
	v_mul_hi_u32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s0
	v_sub_nc_u32_e32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v3, s0, v0
	v_cmp_le_u32_e32 vcc_lo, s0, v0
	v_cndmask_b32_e32 v0, v0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v3, s0, v0
	v_cmp_le_u32_e32 vcc_lo, s0, v0
	v_dual_cndmask_b32 v0, v0, v3 :: v_dual_mov_b32 v3, 0
	v_mov_b32_e32 v4, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v0, v2
	v_sub_nc_u32_e32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_subrev_nc_u32_e32 v0, s6, v0
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_lt_i32_e32 vcc_lo, -1, v0
	v_cmp_gt_i32_e64 s0, s4, v0
	s_and_b32 s1, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_5
; %bb.4:                                ;   in Loop: Header=BB0_3 Depth=1
	v_ashrrev_i32_e32 v6, 31, v5
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, d_Kernel@rel32@lo+4
	s_addc_u32 s7, s7, d_Kernel@rel32@hi+12
	s_ashr_i32 s3, s2, 31
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[12:13], s[2:3], 3
	v_lshlrev_b64 v[6:7], 3, v[5:6]
	s_add_u32 s6, s12, s6
	s_addc_u32 s7, s13, s7
	s_load_b64 s[6:7], s[6:7], 0x0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_u32 v6, vcc_lo, s10, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s11, v7, vcc_lo
	global_load_b64 v[6:7], v[6:7], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[3:4], v[6:7], s[6:7], v[3:4]
.LBB0_5:                                ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v0, 1, v0
	v_add_nc_u32_e32 v5, 1, v5
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s2, -1
	s_cbranch_scc1 .LBB0_3
.LBB0_6:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_store_b64 v[0:1], v[3:4], off
.LBB0_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17convolutionRowGPUPdS_iii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 8
		.amdhsa_next_free_sgpr 16
		.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	_Z17convolutionRowGPUPdS_iii, .Lfunc_end0-_Z17convolutionRowGPUPdS_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 480
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z20convolutionColumnGPUPdS_iii ; -- Begin function _Z20convolutionColumnGPUPdS_iii
	.globl	_Z20convolutionColumnGPUPdS_iii
	.p2align	8
	.type	_Z20convolutionColumnGPUPdS_iii,@function
_Z20convolutionColumnGPUPdS_iii:        ; @_Z20convolutionColumnGPUPdS_iii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x2c
	s_load_b128 s[4:7], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mul_i32 s2, s5, s4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB1_7
; %bb.1:
	s_load_b128 s[8:11], s[0:1], 0x0
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	v_ashrrev_i32_e32 v2, 31, v1
	s_cmp_lt_i32 s6, 0
	s_cbranch_scc1 .LBB1_6
; %bb.2:                                ; %.lr.ph.preheader
	s_ashr_i32 s0, s4, 31
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_nc_u32_e32 v4, v1, v2
	s_add_i32 s1, s4, s0
	s_xor_b32 s1, s1, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v4, v4, v2
	v_cvt_f32_u32_e32 v0, s1
	s_sub_i32 s2, 0, s1
	v_rcp_iflag_f32_e32 v0, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v0, v0
	v_mul_lo_u32 v3, s2, v0
	s_lshl_b32 s2, s6, 1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v0, v3
	v_add_nc_u32_e32 v0, v0, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v0, v4, v0
	v_mul_lo_u32 v3, v0, s1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v3, v4, v3
	v_add_nc_u32_e32 v4, 1, v0
	v_subrev_nc_u32_e32 v5, s1, v3
	v_cmp_le_u32_e32 vcc_lo, s1, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v0, v0, v4 :: v_dual_cndmask_b32 v3, v3, v5
	v_xor_b32_e32 v5, s0, v2
	s_mul_i32 s0, s6, s4
	v_add_nc_u32_e32 v4, 1, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s1, v3
	v_dual_cndmask_b32 v0, v0, v4 :: v_dual_mov_b32 v3, 0
	v_mov_b32_e32 v4, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v0, v5
	v_sub_nc_u32_e32 v0, v0, v5
	v_subrev_nc_u32_e32 v5, s0, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_subrev_nc_u32_e32 v0, s6, v0
.LBB1_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_lt_i32_e32 vcc_lo, -1, v0
	v_cmp_gt_i32_e64 s0, s5, v0
	s_and_b32 s1, vcc_lo, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s1
	s_cbranch_execz .LBB1_5
; %bb.4:                                ;   in Loop: Header=BB1_3 Depth=1
	v_ashrrev_i32_e32 v6, 31, v5
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, d_Kernel@rel32@lo+4
	s_addc_u32 s7, s7, d_Kernel@rel32@hi+12
	s_ashr_i32 s3, s2, 31
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[12:13], s[2:3], 3
	v_lshlrev_b64 v[6:7], 3, v[5:6]
	s_add_u32 s6, s12, s6
	s_addc_u32 s7, s13, s7
	s_load_b64 s[6:7], s[6:7], 0x0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_u32 v6, vcc_lo, s10, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s11, v7, vcc_lo
	global_load_b64 v[6:7], v[6:7], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[3:4], v[6:7], s[6:7], v[3:4]
.LBB1_5:                                ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v5, s4, v5
	v_add_nc_u32_e32 v0, 1, v0
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s2, -1
	s_cbranch_scc1 .LBB1_3
.LBB1_6:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_store_b64 v[0:1], v[3:4], off
.LBB1_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z20convolutionColumnGPUPdS_iii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 8
		.amdhsa_next_free_sgpr 16
		.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_end1:
	.size	_Z20convolutionColumnGPUPdS_iii, .Lfunc_end1-_Z20convolutionColumnGPUPdS_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 500
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	d_Kernel                ; @d_Kernel
	.type	d_Kernel,@object
	.section	.bss,"aw",@nobits
	.globl	d_Kernel
	.p2align	4, 0x0
d_Kernel:
	.zero	136
	.size	d_Kernel, 136

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 d_Kernel
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .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:           _Z17convolutionRowGPUPdS_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17convolutionRowGPUPdS_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .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:           _Z20convolutionColumnGPUPdS_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z20convolutionColumnGPUPdS_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
