/*
Two Dimensional (2D) Image Convolution in CUDA by Shared & Constant Memory: An Optimized way
After learning the concept of two dimension (2D) Convolution and its implementation in C language;
the next step is to learn to optimize it. As Convolution is one of the most Compute Intensive task in 
Image Processing, it is always better to save time required for it. So, today I am going to share a technique 
to optimize the Convolution process by using CUDA. Here we will use Shared Memory and Constant Memory resources
available in CUDA to get fastest implementation of Convolution.

Here is the strategy, 
We will store convolution mask in Constant Memory as all threads will need the same Mask. Constant Memory is 
the off-chip memory present on GPU card which is read-only memory and faster than Global memory. The Mask is 
passed to kernel with the keyword "__restrict__". Check for it in the kernel definition.
Instead of accessing complete image and performing the convolution operation on it; we will fix TILE size. 
The image pixels are copied to this TILE first and then convolution is performed accordingly. 
This TILE will be stored in the Shared Memory of GPU card which is on-chip memory and faster than 
the Global Memory.
Following is the sample code of the implementation. You can also get it from my GitHub repository.

*/
/*
 * Title: 2D Image Convolution on GPU by using Shared Memory and Constant Memory.
 *
 * Image Size: 2048 X 2048
 * Mask Size: 64 X 64
 * TILE_WIDTH 32
 *
 *
 * */
#include<stdio.h>
#include<cuda_runtime_api.h>
#include<cuda.h>
#include<stdlib.h>
#include<assert.h>

#define wbCheck(stmt)  do {                                                    \
        cudaError_t err = stmt;                                               \
        if (err != cudaSuccess) {                                             \
            printf( "Failed to run stmt %d ", __LINE__);                       \
            printf( "Got CUDA error ...  %s ", cudaGetErrorString(err));    \
            return -1;                                                        \
        }                                                                     \
    } while(0)


#define Mask_width  64
#define Mask_height 64
#define Mask_radius_x Mask_width/2
#define Mask_radius_y Mask_height/2
#define TILE_WIDTH 32  //16 X 16 TILE
#define w_x (TILE_WIDTH + Mask_width - 1) //Shared Memory Elements needed to load as per Mask Size
#define w_y (TILE_WIDTH + Mask_height - 1)
#define clamp(x) (min(max((x), 0.0), 1.0))

//@@ INSERT CODE HERE
__global__ void convolution(float *I, const float* __restrict__ M, float *P,int channels, int width, int height)
{
   __shared__ float N_ds[w_y][w_x];
   int k;
   for (k = 0; k < channels; k++)
   {
      //1. Phase to Load Data into Shared Memory. Each Thread loads multiple elements indexed by each Batch loading
    //1.dest: RMO ID 2. destY & destX: Row and Column of Shared Memory
    //3. srcY & srcX: Indexes to fetch data from input Image
    //4. src: RMO index of Input Image

    // First batch loading
      int dest = threadIdx.y * TILE_WIDTH + threadIdx.x,
         destY = dest / w_x, destX = dest % w_x,
         srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius_x,
         srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius_y,
         src = (srcY * width + srcX) * channels + k;
      if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
         N_ds[destY][destX] = I[src];
      else
         N_ds[destY][destX] = 0.0;

        for (int iter=1; iter <= (w_x * w_y) / (TILE_WIDTH*TILE_WIDTH); iter++)
        {
          // Second batch loading
          dest = threadIdx.y * TILE_WIDTH + threadIdx.x + iter*(TILE_WIDTH * TILE_WIDTH);
            destY = dest / w_x, destX = dest % w_x;
            srcY  = blockIdx.y * TILE_WIDTH + destY - Mask_radius_x;
            srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius_y;
            src = (srcY * width + srcX) * channels + k;
            if (destY < w_y && destX < w_x)
            {
                if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
                    N_ds[destY][destX] = I[src];
                else
                    N_ds[destY][destX] = 0.0;
            }
        }
      __syncthreads();

      float accum = 0;
      int y, x;
      for (y = 0; y < Mask_width; y++)
         for (x = 0; x < Mask_width; x++)
            accum += N_ds[threadIdx.y + y][threadIdx.x + x] * M[y * Mask_width + x];
      y = blockIdx.y * TILE_WIDTH + threadIdx.y;
      x = blockIdx.x * TILE_WIDTH + threadIdx.x;
      if (y < height && x < width)
         //P[(y * width + x) * channels + k] = clamp(accum);
       P[(y * width + x) * channels + k] = accum;
      __syncthreads();
   }
}

float convolution_2D_OnHost(float * N,float * M,int width, int height,int i,int j);

int main() {

    int maskRows=Mask_height; // Set it as per requirement of 64 X 32
    int maskColumns=Mask_width;

    int imageChannels=1;
    int imageWidth=2048;
    int imageHeight=2048;

    float * hostInputImageData;
    float * hostOutputImageData;
    float * hostMaskData;
    float * deviceInputImageData;
    float * deviceOutputImageData;
    float * deviceMaskData;
    float * outputImageOnHost;

    hostMaskData = (float *) malloc(sizeof(float)*maskRows*maskColumns);
    for(int i=0;i<maskRows*maskColumns;i++)//To set Mask of size 5*5 which has all values as 1
    {
     hostMaskData[i]=1.0;
    }

//Comment this assert code for 64 X 32
    //assert(maskRows == 5); /* mask height is fixed to 5 in this mp */
    //assert(maskColumns == 5); /* mask width is fixed to 5 in this mp */

    //To store Memory
    hostInputImageData = (float *) malloc(sizeof(float)*imageWidth*imageHeight);
    hostOutputImageData = (float *) malloc(sizeof(float)*imageWidth*imageHeight);
    outputImageOnHost = (float *) malloc(sizeof(float)*imageWidth*imageHeight);
    for(int i=0;i<imageWidth*imageHeight;i++)//To set Image data as 1.0
    {
     hostInputImageData[i]=1.0;
    }

    wbCheck(cudaMalloc((void **) &deviceInputImageData, imageWidth * imageHeight * imageChannels * sizeof(float)));
    wbCheck(cudaMalloc((void **) &deviceOutputImageData, imageWidth * imageHeight * imageChannels * sizeof(float)));
    wbCheck(cudaMalloc((void **) &deviceMaskData, maskRows * maskColumns * sizeof(float)));

    wbCheck(cudaMemcpy(deviceInputImageData,
               hostInputImageData,
               imageWidth * imageHeight * imageChannels * sizeof(float),
               cudaMemcpyHostToDevice));
    wbCheck(cudaMemcpy(deviceMaskData,
               hostMaskData,
               maskRows * maskColumns * sizeof(float),
               cudaMemcpyHostToDevice));

    dim3 dimGrid(((imageWidth-1)/TILE_WIDTH)+1, ((imageHeight-1)/TILE_WIDTH)+1,1);
    dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
    convolution<<<dimGrid,dimBlock>>>(deviceInputImageData, deviceMaskData, deviceOutputImageData,
                                       imageChannels, imageWidth, imageHeight);

 cudaError_t err1 = cudaPeekAtLastError();
 cudaDeviceSynchronize();
 printf( "Got CUDA error ... %s \n", cudaGetErrorString(err1));

    cudaMemcpy(hostOutputImageData,
               deviceOutputImageData,
               imageWidth * imageHeight * imageChannels * sizeof(float),
               cudaMemcpyDeviceToHost);

    //printf("\n Output from Device:\n");
#if 0  //comment this to run the portion of code
    for(int i=0;i<imageWidth*imageHeight;i++)
    {
        if(i>0 && (i%imageWidth==0))
            printf("\n");
      printf("%0.2f \t",*(hostOutputImageData+i));

      }
#endif

    //Convolution on Host
    for(int i=0;i<imageWidth;i++)
          {
           for(int j=0;j<imageHeight;j++)
           {
            outputImageOnHost[(i*imageWidth)+j]=convolution_2D_OnHost(hostInputImageData,hostMaskData,imageWidth,imageHeight,i,j);
           }
          }

    //printf("\n Output from Host:\n");
#if 0  //comment this to run the portion of code
    for(int i=0;i<imageWidth*imageHeight;i++)
      {
      if(i>0 && (i%imageWidth==0))
       printf("\n");
      printf("%0.2f \t",*(outputImageOnHost+i));

      }
#endif


        for(int i=0;i<imageWidth*imageHeight;i++)
        {
         if(outputImageOnHost[i]!=hostOutputImageData[i])
         {
           printf("\nMismatch at (Row,Col) = [%d][%d], hostComputed[]: %0.0f And device[]: %0.0f", i / imageWidth, i % imageHeight, outputImageOnHost[i], hostOutputImageData[i]);
            }
        }

    cudaFree(deviceInputImageData);
    cudaFree(deviceOutputImageData);
    cudaFree(deviceMaskData);

    free(hostMaskData);
    free(hostInputImageData);
    free(hostOutputImageData);

    printf("\n Number of Threads Per Block created in code: %d",TILE_WIDTH*TILE_WIDTH);
    printf("\n Number of Blocks Created :%d",(((imageWidth-1)/TILE_WIDTH)+1)*(((imageWidth-1)/TILE_WIDTH)+1));
    printf("No Error");
    return 0;
}

float convolution_2D_OnHost(float * N,float * M,int width, int height,int i,int j)
{
 float Pvalue=0.0;
 int N_start_point_i = i - (Mask_width/2);
 int N_start_point_j = j - (Mask_height/2);

 for(int k=0;k<Mask_width;k++)
 {
  for(int l=0;l<Mask_height;l++)
  {
   if(((N_start_point_i+k)>=0) && ((N_start_point_i+k)<width)&&((N_start_point_j+l)>=0)&&((N_start_point_j+l)<height))
   {
       Pvalue+=N[(N_start_point_i+k)*width+(N_start_point_j+l)]*M[(k*Mask_width)+l];
   }
  }
 }

// return(clamp(Pvalue));
 return((Pvalue));
}

	.file	"tmpxft_00296783_00000000-6_2D_conv_opt.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	_Z21convolution_2D_OnHostPfS_iiii
	.type	_Z21convolution_2D_OnHostPfS_iiii, @function
_Z21convolution_2D_OnHostPfS_iiii:
.LFB2028:
	.cfi_startproc
	endbr64
	movq	%rdi, %r10
	subl	$32, %r8d
	movl	%edx, %edi
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	imull	%r8d, %edi
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r11
	leal	-32(%r9), %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorps	%xmm0, %xmm0
	addl	$32, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%ecx, %ebx
	xorl	%ecx, %ecx
.L3:
	movl	%r12d, %eax
	movl	%ecx, %esi
.L5:
	cmpl	%r8d, %edx
	jle	.L4
	movl	%eax, %ebp
	orl	%r8d, %ebp
	js	.L4
	cmpl	%eax, %ebx
	jle	.L4
	leal	(%rax,%rdi), %r13d
	movslq	%esi, %rbp
	movslq	%r13d, %r13
	movss	(%r10,%r13,4), %xmm1
	mulss	(%r11,%rbp,4), %xmm1
	addss	%xmm1, %xmm0
.L4:
	incl	%eax
	incl	%esi
	cmpl	%r9d, %eax
	jne	.L5
	addl	$64, %ecx
	addl	%edx, %edi
	incl	%r8d
	cmpl	$4096, %ecx
	jne	.L3
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z21convolution_2D_OnHostPfS_iiii, .-_Z21convolution_2D_OnHostPfS_iiii
	.globl	_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii
	.type	_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii, @function
_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rdx, 16(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 12(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movq	%rsi, 32(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rax, 104(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movl	$1, 64(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	48(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z11convolutionPfPKfS_iii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 192
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L13:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii, .-_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii
	.globl	_Z11convolutionPfPKfS_iii
	.type	_Z11convolutionPfPKfS_iii, @function
_Z11convolutionPfPKfS_iii:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii
	.cfi_endproc
.LFE2054:
	.size	_Z11convolutionPfPKfS_iii, .-_Z11convolutionPfPKfS_iii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"Failed to run stmt %d "
.LC3:
	.string	"Got CUDA error ...  %s "
.LC4:
	.string	"Got CUDA error ... %s \n"
.LC5:
	.string	"\nMismatch at (Row,Col) = [%d][%d], hostComputed[]: %0.0f And device[]: %0.0f"
.LC6:
	.string	"\n Number of Threads Per Block created in code: %d"
.LC7:
	.string	"\n Number of Blocks Created :%d"
.LC8:
	.string	"No Error"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$16384, %edi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movq	%rax, %r12
	xorl	%eax, %eax
.L19:
	movl	$0x3f800000, (%r12,%rax,4)
	incq	%rax
	cmpq	$4096, %rax
	jne	.L19
	movl	$16777216, %edi
	call	malloc@PLT
	movl	$16777216, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movl	$16777216, %edi
	movq	%rax, %r13
	call	malloc@PLT
	movq	%rax, %r14
	xorl	%eax, %eax
.L20:
	movl	$0x3f800000, 0(%rbp,%rax,4)
	incq	%rax
	cmpq	$4194304, %rax
	jne	.L20
	leaq	24(%rsp), %rdi
	movl	$16777216, %esi
	call	cudaMalloc@PLT
	movl	$147, %edx
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L42
	leaq	32(%rsp), %rdi
	movl	$16777216, %esi
	call	cudaMalloc@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L23
	movl	$148, %edx
.L42:
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%ebx, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %eax
	jmp	.L18
.L23:
	leaq	40(%rsp), %rdi
	movl	$16384, %esi
	call	cudaMalloc@PLT
	movl	$149, %edx
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L42
	movq	24(%rsp), %rdi
	movl	$16777216, %edx
	movl	$1, %ecx
	movq	%rbp, %rsi
	call	cudaMemcpy@PLT
	movl	$151, %edx
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L42
	movq	40(%rsp), %rdi
	movl	$16384, %edx
	movl	$1, %ecx
	movq	%r12, %rsi
	call	cudaMemcpy@PLT
	movl	$155, %edx
	movl	%eax, %ebx
	testl	%eax, %eax
	jne	.L42
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$137438953504, %rdx
	movabsq	$274877907008, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	movq	32(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$2048, %r9d
	movl	$2048, %r8d
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	call	_Z39__device_stub__Z11convolutionPfPKfS_iiiPfPKfS_iii
.L28:
	call	cudaPeekAtLastError@PLT
	movl	%eax, %ebx
	call	cudaDeviceSynchronize@PLT
	movl	%ebx, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	32(%rsp), %rsi
	movl	$2, %ecx
	movq	%r13, %rdi
	movl	$16777216, %edx
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
.L29:
	movq	%rax, %r15
	movl	%eax, %r8d
	xorl	%ebx, %ebx
	salq	$13, %r15
	addq	%r14, %r15
.L30:
	movl	%ebx, %r9d
	movl	$2048, %ecx
	movl	$2048, %edx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	movq	%rax, 8(%rsp)
	movl	%r8d, 4(%rsp)
	call	_Z21convolution_2D_OnHostPfS_iiii
	movl	4(%rsp), %r8d
	movq	8(%rsp), %rax
	movss	%xmm0, (%r15,%rbx,4)
	incq	%rbx
	cmpq	$2048, %rbx
	jne	.L30
	incq	%rax
	cmpq	$2048, %rax
	jne	.L29
	xorl	%ebx, %ebx
	leaq	.LC5(%rip), %r15
.L31:
	movss	(%r14,%rbx,4), %xmm0
	movss	0(%r13,%rbx,4), %xmm1
	ucomiss	%xmm1, %xmm0
	jp	.L36
	je	.L32
.L36:
	movl	%ebx, %ecx
	movl	%ebx, %edx
	movq	%r15, %rsi
	movl	$2, %edi
	andl	$2047, %ecx
	sarl	$11, %edx
	movb	$2, %al
	cvtss2sd	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm1
	call	__printf_chk@PLT
.L32:
	incq	%rbx
	cmpq	$4194304, %rbx
	jne	.L31
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	movl	$1024, %edx
	leaq	.LC6(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$4096, %edx
	leaq	.LC7(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%eax, %eax
.L18:
	movq	72(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L34
	call	__stack_chk_fail@PLT
.L34:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC9:
	.string	"_Z11convolutionPfPKfS_iii"
	.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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z11convolutionPfPKfS_iii(%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
	.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:

/*
Two Dimensional (2D) Image Convolution in CUDA by Shared & Constant Memory: An Optimized way
After learning the concept of two dimension (2D) Convolution and its implementation in C language;
the next step is to learn to optimize it. As Convolution is one of the most Compute Intensive task in 
Image Processing, it is always better to save time required for it. So, today I am going to share a technique 
to optimize the Convolution process by using CUDA. Here we will use Shared Memory and Constant Memory resources
available in CUDA to get fastest implementation of Convolution.

Here is the strategy, 
We will store convolution mask in Constant Memory as all threads will need the same Mask. Constant Memory is 
the off-chip memory present on GPU card which is read-only memory and faster than Global memory. The Mask is 
passed to kernel with the keyword "__restrict__". Check for it in the kernel definition.
Instead of accessing complete image and performing the convolution operation on it; we will fix TILE size. 
The image pixels are copied to this TILE first and then convolution is performed accordingly. 
This TILE will be stored in the Shared Memory of GPU card which is on-chip memory and faster than 
the Global Memory.
Following is the sample code of the implementation. You can also get it from my GitHub repository.

*/
/*
 * Title: 2D Image Convolution on GPU by using Shared Memory and Constant Memory.
 *
 * Image Size: 2048 X 2048
 * Mask Size: 64 X 64
 * TILE_WIDTH 32
 *
 *
 * */
#include<stdio.h>
#include<hip/hip_runtime_api.h>
#include<hip/hip_runtime.h>
#include<stdlib.h>
#include<assert.h>

#define wbCheck(stmt)  do {                                                    \
        hipError_t err = stmt;                                               \
        if (err != hipSuccess) {                                             \
            printf( "Failed to run stmt %d ", __LINE__);                       \
            printf( "Got CUDA error ...  %s ", hipGetErrorString(err));    \
            return -1;                                                        \
        }                                                                     \
    } while(0)


#define Mask_width  64
#define Mask_height 64
#define Mask_radius_x Mask_width/2
#define Mask_radius_y Mask_height/2
#define TILE_WIDTH 32  //16 X 16 TILE
#define w_x (TILE_WIDTH + Mask_width - 1) //Shared Memory Elements needed to load as per Mask Size
#define w_y (TILE_WIDTH + Mask_height - 1)
#define clamp(x) (min(max((x), 0.0), 1.0))

//@@ INSERT CODE HERE
__global__ void convolution(float *I, const float* __restrict__ M, float *P,int channels, int width, int height)
{
   __shared__ float N_ds[w_y][w_x];
   int k;
   for (k = 0; k < channels; k++)
   {
      //1. Phase to Load Data into Shared Memory. Each Thread loads multiple elements indexed by each Batch loading
    //1.dest: RMO ID 2. destY & destX: Row and Column of Shared Memory
    //3. srcY & srcX: Indexes to fetch data from input Image
    //4. src: RMO index of Input Image

    // First batch loading
      int dest = threadIdx.y * TILE_WIDTH + threadIdx.x,
         destY = dest / w_x, destX = dest % w_x,
         srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius_x,
         srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius_y,
         src = (srcY * width + srcX) * channels + k;
      if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
         N_ds[destY][destX] = I[src];
      else
         N_ds[destY][destX] = 0.0;

        for (int iter=1; iter <= (w_x * w_y) / (TILE_WIDTH*TILE_WIDTH); iter++)
        {
          // Second batch loading
          dest = threadIdx.y * TILE_WIDTH + threadIdx.x + iter*(TILE_WIDTH * TILE_WIDTH);
            destY = dest / w_x, destX = dest % w_x;
            srcY  = blockIdx.y * TILE_WIDTH + destY - Mask_radius_x;
            srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius_y;
            src = (srcY * width + srcX) * channels + k;
            if (destY < w_y && destX < w_x)
            {
                if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
                    N_ds[destY][destX] = I[src];
                else
                    N_ds[destY][destX] = 0.0;
            }
        }
      __syncthreads();

      float accum = 0;
      int y, x;
      for (y = 0; y < Mask_width; y++)
         for (x = 0; x < Mask_width; x++)
            accum += N_ds[threadIdx.y + y][threadIdx.x + x] * M[y * Mask_width + x];
      y = blockIdx.y * TILE_WIDTH + threadIdx.y;
      x = blockIdx.x * TILE_WIDTH + threadIdx.x;
      if (y < height && x < width)
         //P[(y * width + x) * channels + k] = clamp(accum);
       P[(y * width + x) * channels + k] = accum;
      __syncthreads();
   }
}

float convolution_2D_OnHost(float * N,float * M,int width, int height,int i,int j);

int main() {

    int maskRows=Mask_height; // Set it as per requirement of 64 X 32
    int maskColumns=Mask_width;

    int imageChannels=1;
    int imageWidth=2048;
    int imageHeight=2048;

    float * hostInputImageData;
    float * hostOutputImageData;
    float * hostMaskData;
    float * deviceInputImageData;
    float * deviceOutputImageData;
    float * deviceMaskData;
    float * outputImageOnHost;

    hostMaskData = (float *) malloc(sizeof(float)*maskRows*maskColumns);
    for(int i=0;i<maskRows*maskColumns;i++)//To set Mask of size 5*5 which has all values as 1
    {
     hostMaskData[i]=1.0;
    }

//Comment this assert code for 64 X 32
    //assert(maskRows == 5); /* mask height is fixed to 5 in this mp */
    //assert(maskColumns == 5); /* mask width is fixed to 5 in this mp */

    //To store Memory
    hostInputImageData = (float *) malloc(sizeof(float)*imageWidth*imageHeight);
    hostOutputImageData = (float *) malloc(sizeof(float)*imageWidth*imageHeight);
    outputImageOnHost = (float *) malloc(sizeof(float)*imageWidth*imageHeight);
    for(int i=0;i<imageWidth*imageHeight;i++)//To set Image data as 1.0
    {
     hostInputImageData[i]=1.0;
    }

    wbCheck(hipMalloc((void **) &deviceInputImageData, imageWidth * imageHeight * imageChannels * sizeof(float)));
    wbCheck(hipMalloc((void **) &deviceOutputImageData, imageWidth * imageHeight * imageChannels * sizeof(float)));
    wbCheck(hipMalloc((void **) &deviceMaskData, maskRows * maskColumns * sizeof(float)));

    wbCheck(hipMemcpy(deviceInputImageData,
               hostInputImageData,
               imageWidth * imageHeight * imageChannels * sizeof(float),
               hipMemcpyHostToDevice));
    wbCheck(hipMemcpy(deviceMaskData,
               hostMaskData,
               maskRows * maskColumns * sizeof(float),
               hipMemcpyHostToDevice));

    dim3 dimGrid(((imageWidth-1)/TILE_WIDTH)+1, ((imageHeight-1)/TILE_WIDTH)+1,1);
    dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
    convolution<<<dimGrid,dimBlock>>>(deviceInputImageData, deviceMaskData, deviceOutputImageData,
                                       imageChannels, imageWidth, imageHeight);

 hipError_t err1 = hipPeekAtLastError();
 hipDeviceSynchronize();
 printf( "Got CUDA error ... %s \n", hipGetErrorString(err1));

    hipMemcpy(hostOutputImageData,
               deviceOutputImageData,
               imageWidth * imageHeight * imageChannels * sizeof(float),
               hipMemcpyDeviceToHost);

    //printf("\n Output from Device:\n");
#if 0  //comment this to run the portion of code
    for(int i=0;i<imageWidth*imageHeight;i++)
    {
        if(i>0 && (i%imageWidth==0))
            printf("\n");
      printf("%0.2f \t",*(hostOutputImageData+i));

      }
#endif

    //Convolution on Host
    for(int i=0;i<imageWidth;i++)
          {
           for(int j=0;j<imageHeight;j++)
           {
            outputImageOnHost[(i*imageWidth)+j]=convolution_2D_OnHost(hostInputImageData,hostMaskData,imageWidth,imageHeight,i,j);
           }
          }

    //printf("\n Output from Host:\n");
#if 0  //comment this to run the portion of code
    for(int i=0;i<imageWidth*imageHeight;i++)
      {
      if(i>0 && (i%imageWidth==0))
       printf("\n");
      printf("%0.2f \t",*(outputImageOnHost+i));

      }
#endif


        for(int i=0;i<imageWidth*imageHeight;i++)
        {
         if(outputImageOnHost[i]!=hostOutputImageData[i])
         {
           printf("\nMismatch at (Row,Col) = [%d][%d], hostComputed[]: %0.0f And device[]: %0.0f", i / imageWidth, i % imageHeight, outputImageOnHost[i], hostOutputImageData[i]);
            }
        }

    hipFree(deviceInputImageData);
    hipFree(deviceOutputImageData);
    hipFree(deviceMaskData);

    free(hostMaskData);
    free(hostInputImageData);
    free(hostOutputImageData);

    printf("\n Number of Threads Per Block created in code: %d",TILE_WIDTH*TILE_WIDTH);
    printf("\n Number of Blocks Created :%d",(((imageWidth-1)/TILE_WIDTH)+1)*(((imageWidth-1)/TILE_WIDTH)+1));
    printf("No Error");
    return 0;
}

float convolution_2D_OnHost(float * N,float * M,int width, int height,int i,int j)
{
 float Pvalue=0.0;
 int N_start_point_i = i - (Mask_width/2);
 int N_start_point_j = j - (Mask_height/2);

 for(int k=0;k<Mask_width;k++)
 {
  for(int l=0;l<Mask_height;l++)
  {
   if(((N_start_point_i+k)>=0) && ((N_start_point_i+k)<width)&&((N_start_point_j+l)>=0)&&((N_start_point_j+l)<height))
   {
       Pvalue+=N[(N_start_point_i+k)*width+(N_start_point_j+l)]*M[(k*Mask_width)+l];
   }
  }
 }

// return(clamp(Pvalue));
 return((Pvalue));
}

	.text
	.file	"2D_conv_opt.hip"
	.globl	_Z26__device_stub__convolutionPfPKfS_iii # -- Begin function _Z26__device_stub__convolutionPfPKfS_iii
	.type	_Z26__device_stub__convolutionPfPKfS_iii,@function
_Z26__device_stub__convolutionPfPKfS_iii: # @_Z26__device_stub__convolutionPfPKfS_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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z11convolutionPfPKfS_iii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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	_Z26__device_stub__convolutionPfPKfS_iii, .Lfunc_end0-_Z26__device_stub__convolutionPfPKfS_iii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function 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	$40, %rsp
	.cfi_def_cfa_offset 96
	.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	$16384, %edi                    # imm = 0x4000
	callq	malloc
	movq	%rax, %rbx
	xorl	%eax, %eax
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%rbx,%rax,4)      # imm = 0x3F800000
	incq	%rax
	cmpq	$4096, %rax                     # imm = 0x1000
	jne	.LBB1_1
# %bb.2:
	movl	$16777216, %edi                 # imm = 0x1000000
	callq	malloc
	movq	%rax, %r14
	movl	$16777216, %edi                 # imm = 0x1000000
	callq	malloc
	movq	%rax, %r12
	movl	$16777216, %edi                 # imm = 0x1000000
	callq	malloc
	movq	%rax, %r15
	xorl	%eax, %eax
.LBB1_3:                                # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%r14,%rax,4)      # imm = 0x3F800000
	incq	%rax
	cmpq	$4194304, %rax                  # imm = 0x400000
	jne	.LBB1_3
# %bb.4:
	leaq	24(%rsp), %rdi
	movl	$16777216, %esi                 # imm = 0x1000000
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_7
# %bb.5:
	movl	%eax, %ebp
	movl	$.L.str, %edi
	movl	$147, %esi
	jmp	.LBB1_6
.LBB1_7:
	leaq	16(%rsp), %rdi
	movl	$16777216, %esi                 # imm = 0x1000000
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_9
# %bb.8:
	movl	%eax, %ebp
	movl	$.L.str, %edi
	movl	$148, %esi
	jmp	.LBB1_6
.LBB1_9:
	leaq	8(%rsp), %rdi
	movl	$16384, %esi                    # imm = 0x4000
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB1_11
# %bb.10:
	movl	%eax, %ebp
	movl	$.L.str, %edi
	movl	$149, %esi
	jmp	.LBB1_6
.LBB1_11:
	movq	24(%rsp), %rdi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_13
# %bb.12:
	movl	%eax, %ebp
	movl	$.L.str, %edi
	movl	$154, %esi
	jmp	.LBB1_6
.LBB1_13:
	movq	8(%rsp), %rdi
	movl	$16384, %edx                    # imm = 0x4000
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB1_15
# %bb.14:
	movl	%eax, %ebp
	movl	$.L.str, %edi
	movl	$158, %esi
.LBB1_6:
	xorl	%eax, %eax
	callq	printf
	movl	%ebp, %edi
	callq	hipGetErrorString
	movl	$.L.str.1, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$-1, %ebx
.LBB1_26:
	movl	%ebx, %eax
	addq	$40, %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
.LBB1_15:
	.cfi_def_cfa_offset 96
	movq	%r12, 32(%rsp)                  # 8-byte Spill
	movabsq	$274877907008, %rdi             # imm = 0x4000000040
	movabsq	$137438953504, %rdx             # imm = 0x2000000020
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_17
# %bb.16:
	movq	24(%rsp), %rdi
	movq	8(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	$1, %ecx
	movl	$2048, %r8d                     # imm = 0x800
	movl	$2048, %r9d                     # imm = 0x800
	callq	_Z26__device_stub__convolutionPfPKfS_iii
.LBB1_17:
	callq	hipPeekAtLastError
	movl	%eax, %ebp
	callq	hipDeviceSynchronize
	movl	%ebp, %edi
	callq	hipGetErrorString
	xorl	%r13d, %r13d
	movl	$.L.str.2, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	16(%rsp), %rsi
	movl	$16777216, %edx                 # imm = 0x1000000
	movq	32(%rsp), %rdi                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	%r15, %r12
.LBB1_18:                               # %.preheader144
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_19 Depth 2
	xorl	%ebp, %ebp
.LBB1_19:                               #   Parent Loop BB1_18 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	$2048, %edx                     # imm = 0x800
	movl	$2048, %ecx                     # imm = 0x800
	movl	%r13d, %r8d
	movl	%ebp, %r9d
	callq	_Z21convolution_2D_OnHostPfS_iiii
	movss	%xmm0, (%r12,%rbp,4)
	incq	%rbp
	cmpq	$2048, %rbp                     # imm = 0x800
	jne	.LBB1_19
# %bb.20:                               #   in Loop: Header=BB1_18 Depth=1
	incq	%r13
	addq	$8192, %r12                     # imm = 0x2000
	cmpq	$2048, %r13                     # imm = 0x800
	jne	.LBB1_18
# %bb.21:                               # %.preheader.preheader
	xorl	%r12d, %r12d
	movq	32(%rsp), %r13                  # 8-byte Reload
.LBB1_22:                               # %.preheader
                                        # =>This Inner Loop Header: Depth=1
	movss	(%r15,%r12,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movss	(%r13,%r12,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	ucomiss	%xmm1, %xmm0
	jne	.LBB1_23
	jnp	.LBB1_24
.LBB1_23:                               #   in Loop: Header=BB1_22 Depth=1
	movl	%r12d, %esi
	shrl	$11, %esi
	movl	%r12d, %edx
	andl	$2047, %edx                     # imm = 0x7FF
	cvtss2sd	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm1
	movl	$.L.str.3, %edi
	movb	$2, %al
	callq	printf
.LBB1_24:                               #   in Loop: Header=BB1_22 Depth=1
	incq	%r12
	cmpq	$4194304, %r12                  # imm = 0x400000
	jne	.LBB1_22
# %bb.25:
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r13, %rdi
	callq	free
	xorl	%ebx, %ebx
	movl	$.L.str.4, %edi
	movl	$1024, %esi                     # imm = 0x400
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.5, %edi
	movl	$4096, %esi                     # imm = 0x1000
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.6, %edi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB1_26
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z21convolution_2D_OnHostPfS_iiii # -- Begin function _Z21convolution_2D_OnHostPfS_iiii
	.type	_Z21convolution_2D_OnHostPfS_iiii,@function
_Z21convolution_2D_OnHostPfS_iiii:      # @_Z21convolution_2D_OnHostPfS_iiii
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	addl	$-32, %r8d
	addl	$-32, %r9d
	movslq	%r9d, %rax
	movslq	%ecx, %rcx
	movslq	%r8d, %r8
	movslq	%edx, %rdx
	movq	%r8, %r9
	imulq	%rdx, %r9
	leaq	(,%rax,4), %r10
	leaq	(%r10,%r9,4), %r9
	addq	%rdi, %r9
	leaq	(,%rdx,4), %rdi
	xorps	%xmm0, %xmm0
	xorl	%r10d, %r10d
.LBB2_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_2 Depth 2
	movq	%r10, %r11
	addq	%r8, %r11
	setns	%bl
	cmpq	%rdx, %r11
	setl	%r11b
	andb	%bl, %r11b
	xorl	%ebx, %ebx
.LBB2_2:                                #   Parent Loop BB2_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	testb	%r11b, %r11b
	je	.LBB2_6
# %bb.3:                                #   in Loop: Header=BB2_2 Depth=2
	movq	%rax, %r14
	addq	%rbx, %r14
	js	.LBB2_6
# %bb.4:                                #   in Loop: Header=BB2_2 Depth=2
	cmpq	%rcx, %r14
	jge	.LBB2_6
# %bb.5:                                #   in Loop: Header=BB2_2 Depth=2
	movss	(%r9,%rbx,4), %xmm1             # xmm1 = mem[0],zero,zero,zero
	mulss	(%rsi,%rbx,4), %xmm1
	addss	%xmm1, %xmm0
.LBB2_6:                                #   in Loop: Header=BB2_2 Depth=2
	incq	%rbx
	cmpq	$64, %rbx
	jne	.LBB2_2
# %bb.7:                                #   in Loop: Header=BB2_1 Depth=1
	incq	%r10
	addq	$256, %rsi                      # imm = 0x100
	addq	%rdi, %r9
	cmpq	$64, %r10
	jne	.LBB2_1
# %bb.8:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z21convolution_2D_OnHostPfS_iiii, .Lfunc_end2-_Z21convolution_2D_OnHostPfS_iiii
	.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	$_Z11convolutionPfPKfS_iii, %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	_Z11convolutionPfPKfS_iii,@object # @_Z11convolutionPfPKfS_iii
	.section	.rodata,"a",@progbits
	.globl	_Z11convolutionPfPKfS_iii
	.p2align	3, 0x0
_Z11convolutionPfPKfS_iii:
	.quad	_Z26__device_stub__convolutionPfPKfS_iii
	.size	_Z11convolutionPfPKfS_iii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Failed to run stmt %d "
	.size	.L.str, 23

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Got CUDA error ...  %s "
	.size	.L.str.1, 24

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Got CUDA error ... %s \n"
	.size	.L.str.2, 24

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\nMismatch at (Row,Col) = [%d][%d], hostComputed[]: %0.0f And device[]: %0.0f"
	.size	.L.str.3, 77

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"\n Number of Threads Per Block created in code: %d"
	.size	.L.str.4, 50

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\n Number of Blocks Created :%d"
	.size	.L.str.5, 31

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"No Error"
	.size	.L.str.6, 9

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z11convolutionPfPKfS_iii"
	.size	.L__unnamed_1, 26

	.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 _Z26__device_stub__convolutionPfPKfS_iii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z11convolutionPfPKfS_iii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z11convolutionPfPKfS_iii ; -- Begin function _Z11convolutionPfPKfS_iii
	.globl	_Z11convolutionPfPKfS_iii
	.p2align	8
	.type	_Z11convolutionPfPKfS_iii,@function
_Z11convolutionPfPKfS_iii:              ; @_Z11convolutionPfPKfS_iii
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s4, 1
	s_cbranch_scc1 .LBB0_19
; %bb.1:                                ; %.lr.ph
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v5, v0, 10, 10
	s_clause 0x1
	s_load_b128 s[8:11], s[0:1], 0x0
	s_load_b64 s[12:13], s[0:1], 0x10
	s_lshl_b32 s1, s14, 5
	s_lshl_b32 s0, s15, 5
	s_sub_i32 s16, s1, 32
	v_lshl_add_u32 v6, v5, 5, v1
	v_add_nc_u32_e32 v4, s0, v5
	v_lshlrev_b32_e32 v9, 2, v1
	s_sub_i32 s7, s0, 32
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v0, 0xffff, v6
	v_cmp_gt_i32_e32 vcc_lo, s6, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mad_u32_u24 v5, 0x17c, v5, v9
	v_mul_u32_u24_e32 v0, 0xac77, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v7, 22, v0
	v_mul_lo_u16 v0, 0x5f, v7
	v_add_nc_u32_e32 v10, s7, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u16 v0, v6, v0
	v_and_b32_e32 v8, 0xffff, v0
	v_add_nc_u32_e32 v0, s1, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_gt_i32_e64 s1, s6, v10
	v_add_nc_u32_e32 v1, s16, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_gt_i32_e64 s0, s5, v0
	v_mad_u64_u32 v[2:3], null, v4, s5, v[0:1]
	v_lshlrev_b32_e32 v0, 2, v8
	v_cmp_gt_i32_e64 s2, s5, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_and_b32 s17, s0, vcc_lo
	s_and_b32 s1, s1, s2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_mad_u64_u32 v[3:4], null, v10, s5, v[1:2]
	v_or_b32_e32 v4, v10, v1
	v_mul_lo_u32 v1, v2, s4
	v_mad_u32_u24 v2, 0x17c, v7, v0
	s_mov_b32 s2, 0
	v_cmp_lt_i32_e64 s3, -1, v4
	v_mul_lo_u32 v3, v3, s4
	v_add_nc_u32_e32 v4, 0x400, v6
	s_delay_alu instid0(VALU_DEP_3)
	s_and_b32 s3, s1, s3
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
                                        ;     Child Loop BB0_13 Depth 2
                                        ;       Child Loop BB0_14 Depth 3
	v_mov_b32_e32 v0, 0
	s_and_saveexec_b32 s0, s3
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, s2, v3
	v_ashrrev_i32_e32 v7, 31, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v6, vcc_lo, s8, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v7, vcc_lo
	global_load_b32 v0, v[6:7], off
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_mov_b32 s14, 0
	s_waitcnt vmcnt(0)
	ds_store_b32 v2, v0
.LBB0_5:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_nc_u32_e32 v0, s14, v4
	s_mov_b32 s15, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e32 0x2341, v0
	s_cbranch_execz .LBB0_11
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=2
	v_mul_hi_u32 v6, 0x58ed2309, v0
	s_mov_b32 s18, exec_lo
	v_mov_b32_e32 v9, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v7, v0, v6
	v_lshrrev_b32_e32 v7, 1, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v7, v6
	v_lshrrev_b32_e32 v6, 6, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v7, 0x5f, v6
	v_add_nc_u32_e32 v8, s7, v6
	v_sub_nc_u32_e32 v7, v0, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_lt_i32_e32 -1, v8
	s_cbranch_execz .LBB0_10
; %bb.7:                                ;   in Loop: Header=BB0_5 Depth=2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_mov_b32 v9, 0 :: v_dual_add_nc_u32 v0, s16, v7
	v_cmp_gt_i32_e32 vcc_lo, s6, v8
	v_cmp_gt_i32_e64 s0, s5, v0
	v_cmp_lt_i32_e64 s1, -1, v0
	s_delay_alu instid0(VALU_DEP_2)
	s_and_b32 s0, vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s1, s0, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_5 Depth=2
	v_mad_u64_u32 v[9:10], null, v8, s5, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[10:11], null, v9, s4, s[2:3]
	v_ashrrev_i32_e32 v11, 31, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[8:9], 2, v[10:11]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v8, vcc_lo, s8, v8
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v9, vcc_lo, s9, v9, vcc_lo
	global_load_b32 v9, v[8:9], off
.LBB0_9:                                ; %Flow
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s0
.LBB0_10:                               ; %.sink.split
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s18
	v_mul_lo_u32 v0, 0x17c, v6
	v_lshl_add_u32 v0, v7, 2, v0
	s_waitcnt vmcnt(0)
	ds_store_b32 v0, v9
.LBB0_11:                               ; %Flow129
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s15
	s_addk_i32 s14, 0x400
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmpk_eq_i32 s14, 0x2000
	s_cbranch_scc0 .LBB0_5
; %bb.12:                               ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v0, 0
	v_mov_b32_e32 v6, v5
	s_mov_b32 s18, 0
	s_waitcnt lgkmcnt(0)
	s_mov_b64 s[0:1], s[10:11]
	s_barrier
	buffer_gl0_inv
.LBB0_13:                               ; %.preheader
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_14 Depth 3
	s_mov_b64 s[14:15], s[0:1]
	s_mov_b32 s19, 0
.LBB0_14:                               ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_13 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(SALU_CYCLE_1)
	v_add_nc_u32_e32 v7, s19, v6
	s_load_b32 s20, s[14:15], 0x0
	s_add_i32 s19, s19, 4
	s_add_u32 s14, s14, 4
	s_addc_u32 s15, s15, 0
	ds_load_b32 v7, v7
	s_cmpk_lg_i32 s19, 0x100
	s_waitcnt lgkmcnt(0)
	v_fmac_f32_e32 v0, s20, v7
	s_cbranch_scc1 .LBB0_14
; %bb.15:                               ;   in Loop: Header=BB0_13 Depth=2
	s_add_i32 s18, s18, 1
	v_add_nc_u32_e32 v6, 0x17c, v6
	s_add_u32 s0, s0, 0x100
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s18, 64
	s_cbranch_scc1 .LBB0_13
; %bb.16:                               ;   in Loop: Header=BB0_2 Depth=1
	s_and_saveexec_b32 s0, s17
	s_cbranch_execz .LBB0_18
; %bb.17:                               ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v6, s2, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v7, 31, v6
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v6, vcc_lo, s12, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s13, v7, vcc_lo
	global_store_b32 v[6:7], v0, off
.LBB0_18:                               ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	s_add_i32 s2, s2, 1
	s_waitcnt_vscnt null, 0x0
	s_cmp_lg_u32 s2, s4
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_2
.LBB0_19:                               ; %._crit_edge
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11convolutionPfPKfS_iii
		.amdhsa_group_segment_fixed_size 36100
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 36
		.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 12
		.amdhsa_next_free_sgpr 21
		.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	_Z11convolutionPfPKfS_iii, .Lfunc_end0-_Z11convolutionPfPKfS_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 896
; NumSgprs: 23
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 36100 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 23
; NumVGPRsForWavesPerEU: 12
; 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
      - .actual_access:  read_only
        .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 36100
    .kernarg_segment_align: 8
    .kernarg_segment_size: 36
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11convolutionPfPKfS_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     23
    .sgpr_spill_count: 0
    .symbol:         _Z11convolutionPfPKfS_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
