/**
 * File              : gpu_conv.cu
 * Author            : Xianglan Piao <lanxlpiao@gmail.com>
 * Date              : 2020.06.16
 * Last Modified Date: 2020.07.31
 * Last Modified By  : Xianglan Piao <lanxlpiao@gmail.com>
 * NOTE:             : cuda conv2d
 */

#include <iostream>

#define ifm_size 8
#define wgt_size 5
#define stride 2
#define padding 2  // For same ifm/ofm size padding = wgt_size / 2
#define ofm_size ((ifm_size + 2 * padding - wgt_size) / stride + 1)

//// constant memory
__constant__ float wgt[wgt_size];

//// global memory
__global__ void cuda_conv1d_naive(float* ifm, float* ofm) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  float temp = 0.0f;
  for (int k = 0; k < wgt_size; k++) {
    int col_offset = col * stride - padding + k;
    if (col_offset >= 0 && col_offset < ifm_size) {
      temp += ifm[col_offset] * wgt[k];
    }
  }

  ofm[col] = temp;
}

//// shared memory
#define BLOCK_size 4
__global__ void cuda_conv1d_shared(float* ifm, float* ofm) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  __shared__ float shared_block[BLOCK_size];

  float temp = 0.0f;
  for (int k = 0; k < wgt_size; k++) {
    int col_offset = col * stride - padding + k;

    shared_block[threadIdx.x] = ifm[col_offset];
    __syncthreads();

    if (col_offset >= 0 && col_offset < ifm_size) {
      temp += shared_block[threadIdx.x] * wgt[k];
    }
  }

  ofm[col] = temp;
}

void initData(float* data, int size) {
  for (int i = 0; i < size; i++) {
    data[i] = i + 1;
  }
}

void print(float* data, dim3 dim) {
  for (int x = 0; x < dim.x; x++) {
    if (dim.y == 0) {
      std::cout << data[x] << ",  ";
    } else {
      for (int y = 0; y < dim.y; y++) {
        std::cout << data[x * dim.x + y] << ",  ";
      }
      std::cout << std::endl;
    }
  }
  std::cout << std::endl;
}

//// constant memory
__constant__ float wgt2d[wgt_size * wgt_size];

//// global memory
__global__ void cuda_conv2d_naive(float* ifm, float* ofm) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  int row = blockIdx.y * blockDim.y + threadIdx.y;

  float temp = 0.0f;
  for (int m = 0; m < wgt_size; m++) {
    for (int n = 0; n < wgt_size; n++) {
      int col_offset = col * stride - padding + m;
      int row_offset = row * stride - padding + n;
      if ((col_offset >= 0 && col_offset < ifm_size) &&
          (row_offset >= 0 && row_offset < ifm_size)) {
        temp +=
            ifm[row_offset * ifm_size + col_offset] * wgt2d[m * wgt_size + n];
      }
    }

    ofm[row * ofm_size + col] = temp;
  }
}

void test() {
  dim3 block(0);
  dim3 grid(0);

  float* h_ifm = NULL;
  float* d_ifm = NULL;
  float* d_ofm = NULL;
  float* h_ofm = NULL;

  //// 1D convolution
  std::cout << "\n--- 1D convolution ---\n" << std::endl;
  std::cout << "ifm: " << std::endl;
  h_ifm = (float*)malloc(ifm_size * sizeof(float));
  initData(h_ifm, ifm_size);
  print(h_ifm, dim3(ifm_size, 0, 0));
  std::cout << std::endl;

  cudaMalloc((void**)&d_ifm, ifm_size * sizeof(float));
  cudaMemcpy(d_ifm, h_ifm, ifm_size * sizeof(float), cudaMemcpyHostToDevice);

  float h_wgt[wgt_size] = {1, 2, 4, 2, 1};
  std::cout << "wgt: " << std::endl;
  print(h_wgt, dim3(wgt_size, 0, 0));
  std::cout << std::endl;
  cudaMemcpyToSymbol(wgt, &h_wgt, sizeof(wgt));

  block.x = BLOCK_size;
  grid.x = ofm_size / block.x;

  cudaMalloc((void**)&d_ofm, ofm_size * sizeof(float));
  h_ofm = (float*)calloc(ofm_size, sizeof(float));

  //// using global memory
  cudaMemset(d_ofm, 0, ofm_size * sizeof(float));
  cuda_conv1d_naive<<<grid, block>>>(d_ifm, d_ofm);
  cudaDeviceSynchronize();

  memset(h_ofm, 0, ofm_size);
  cudaMemcpy(h_ofm, d_ofm, ofm_size * sizeof(float), cudaMemcpyDeviceToHost);
  std::cout << "ofm: " << std::endl;
  print(h_ofm, dim3(ofm_size, 0, 0));
  std::cout << std::endl;

  //// using shared memory
  cudaMemset(d_ofm, 0, ofm_size * sizeof(float));
  cuda_conv1d_shared<<<grid, block>>>(d_ifm, d_ofm);
  cudaDeviceSynchronize();

  memset(h_ofm, 0, ofm_size);
  cudaMemcpy(h_ofm, d_ofm, ofm_size * sizeof(float), cudaMemcpyDeviceToHost);
  std::cout << "ofm: " << std::endl;
  print(h_ofm, dim3(ofm_size, 0, 0));
  std::cout << std::endl;

  //// free
  cudaFree(d_ifm);
  cudaFree(d_ofm);
  free(h_ifm);
  free(h_ofm);

  //// 2D convolution
  std::cout << "\n--- 2D convolution ---\n" << std::endl;
  h_ifm = (float*)calloc(ifm_size * ifm_size, sizeof(float));
  initData(h_ifm, ifm_size * ifm_size);
  std::cout << "ifm: " << std::endl;
  print(h_ifm, dim3(ifm_size, ifm_size, 0));
  std::cout << std::endl;

  cudaMalloc((void**)&d_ifm, ifm_size * ifm_size * sizeof(float));
  cudaMemcpy(d_ifm,
             h_ifm,
             ifm_size * ifm_size * sizeof(float),
             cudaMemcpyHostToDevice);

  float h_wgt2d[wgt_size * wgt_size] = {1, 1, 1, 1, 1,   //
                                        1, 2, 2, 2, 1,   //
                                        1, 2, 4, 2, 1,   //
                                        1, 2, 2, 2, 1,   //
                                        1, 1, 1, 1, 1};  //
  std::cout << "wgt: " << std::endl;
  print(h_wgt2d, dim3(wgt_size, wgt_size, 0));
  std::cout << std::endl;
  cudaMemcpyToSymbol(wgt2d, &h_wgt2d, sizeof(wgt2d));

  block.x = BLOCK_size;
  block.y = BLOCK_size;
  grid.x = ofm_size / block.x;
  grid.y = ofm_size / block.y;

  cudaMalloc((void**)&d_ofm, ofm_size * ofm_size * sizeof(float));
  h_ofm = (float*)calloc(ofm_size * ofm_size, sizeof(float));
  cudaMemset(d_ofm, 0, ofm_size * ofm_size * sizeof(float));

  cuda_conv2d_naive<<<grid, block>>>(d_ifm, d_ofm);
  cudaDeviceSynchronize();

  memset(h_ofm, 0, ofm_size * ofm_size);
  cudaMemcpy(h_ofm,
             d_ofm,
             ofm_size * ofm_size * sizeof(float),
             cudaMemcpyDeviceToHost);
  std::cout << "ofm: " << std::endl;
  print(h_ofm, dim3(ofm_size, ofm_size, 0));
  std::cout << std::endl;

  //// free
  cudaFree(d_ifm);
  cudaFree(d_ofm);
  free(h_ifm);
  free(h_ofm);

  //// reset
  cudaDeviceReset();
}

int main(void) {
  test();
  return 0;
}

	.file	"tmpxft_0031ff7e_00000000-6_test_gpu_conv.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3641:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3641:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4302:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4302:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z8initDataPfi
	.type	_Z8initDataPfi, @function
_Z8initDataPfi:
.LFB3635:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
.L4:
	cmpl	%eax, %esi
	jle	.L7
	leal	1(%rax), %edx
	cvtsi2ssl	%edx, %xmm0
	movss	%xmm0, (%rdi,%rax,4)
	incq	%rax
	jmp	.L4
.L7:
	ret
	.cfi_endproc
.LFE3635:
	.size	_Z8initDataPfi, .-_Z8initDataPfi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	",  "
	.text
	.globl	_Z5printPf4dim3
	.type	_Z5printPf4dim3, @function
_Z5printPf4dim3:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%rsi, 16(%rsp)
	movl	20(%rsp), %r14d
.L9:
	cmpl	%r13d, %ebx
	jnb	.L16
	testl	%r14d, %r14d
	jne	.L13
	leaq	_ZSt4cout(%rip), %rdi
	cvtss2sd	(%r12,%rbx,4), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	jmp	.L11
.L13:
	movl	%ebp, %edx
	leal	(%r14,%rbp), %r15d
.L10:
	movl	%edx, %eax
	leaq	_ZSt4cout(%rip), %rdi
	movl	%edx, 12(%rsp)
	cvtss2sd	(%r12,%rax,4), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	12(%rsp), %edx
	incl	%edx
	cmpl	%r15d, %edx
	jne	.L10
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L11:
	incq	%rbx
	addl	%r13d, %ebp
	jmp	.L9
.L16:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	leaq	_ZSt4cout(%rip), %rdi
	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
	jmp	_ZNSolsEPFRSoS_E.isra.0
	.cfi_endproc
.LFE3636:
	.size	_Z5printPf4dim3, .-_Z5printPf4dim3
	.globl	_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_
	.type	_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_, @function
_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_:
.LFB3663:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L17
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17cuda_conv1d_naivePfS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L17:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3663:
	.size	_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_, .-_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_
	.globl	_Z17cuda_conv1d_naivePfS_
	.type	_Z17cuda_conv1d_naivePfS_, @function
_Z17cuda_conv1d_naivePfS_:
.LFB3664:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_
	.cfi_endproc
.LFE3664:
	.size	_Z17cuda_conv1d_naivePfS_, .-_Z17cuda_conv1d_naivePfS_
	.globl	_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_
	.type	_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_, @function
_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_:
.LFB3665:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L22
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z18cuda_conv1d_sharedPfS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L22:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L24
	call	__stack_chk_fail@PLT
.L24:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3665:
	.size	_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_, .-_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_
	.globl	_Z18cuda_conv1d_sharedPfS_
	.type	_Z18cuda_conv1d_sharedPfS_, @function
_Z18cuda_conv1d_sharedPfS_:
.LFB3666:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_
	.cfi_endproc
.LFE3666:
	.size	_Z18cuda_conv1d_sharedPfS_, .-_Z18cuda_conv1d_sharedPfS_
	.globl	_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_
	.type	_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_, @function
_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_:
.LFB3667:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17cuda_conv2d_naivePfS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L27:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3667:
	.size	_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_, .-_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_
	.globl	_Z17cuda_conv2d_naivePfS_
	.type	_Z17cuda_conv2d_naivePfS_, @function
_Z17cuda_conv2d_naivePfS_:
.LFB3668:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_
	.cfi_endproc
.LFE3668:
	.size	_Z17cuda_conv2d_naivePfS_, .-_Z17cuda_conv2d_naivePfS_
	.section	.rodata.str1.1
.LC2:
	.string	"\n--- 1D convolution ---\n"
.LC3:
	.string	"ifm: "
.LC5:
	.string	"wgt: "
.LC6:
	.string	"ofm: "
.LC7:
	.string	"\n--- 2D convolution ---\n"
	.section	.rodata
	.align 32
.LC0:
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1073741824
	.long	1073741824
	.long	1073741824
	.long	1065353216
	.long	1065353216
	.long	1073741824
	.long	1082130432
	.long	1073741824
	.long	1065353216
	.long	1065353216
	.long	1073741824
	.long	1073741824
	.long	1073741824
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.long	1065353216
	.text
	.globl	_Z4testv
	.type	_Z4testv, @function
_Z4testv:
.LFB3637:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r11d, %r11d
	leaq	.LC2(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	leaq	.LC3(%rip), %r15
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	leaq	.LC5(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	subq	$200, %rsp
	.cfi_def_cfa_offset 256
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	movq	%r11, 8(%rsp)
	leaq	8(%rsp), %r14
	movq	%r11, 16(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r15, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$32, %edi
	call	malloc@PLT
	movl	$8, %esi
	movq	%rax, %rdi
	movq	%rax, %rbp
	call	_Z8initDataPfi
	xorl	%edx, %edx
	movl	$8, %esi
	movl	%ebx, 60(%rsp)
	movq	$8, 52(%rsp)
	leaq	64(%rsp), %rbx
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$32, %esi
	movq	%r14, %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$32, %edx
	call	cudaMemcpy@PLT
	movq	%r13, %rsi
	movabsq	$4611686019492741120, %rax
	leaq	_ZSt4cout(%rip), %rdi
	movq	%rax, 64(%rsp)
	addq	$16777216, %rax
	movq	%rax, 72(%rsp)
	movl	$0x3f800000, 80(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edx, %edx
	movq	%rbx, %rdi
	movl	$5, %esi
	movl	%r12d, 60(%rsp)
	leaq	16(%rsp), %r12
	movq	$5, 52(%rsp)
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%ecx, %ecx
	movl	$1, %r8d
	movq	%rbx, %rsi
	movl	$20, %edx
	leaq	_ZL3wgt(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	$16, %esi
	movq	%r12, %rdi
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movl	$1, 48(%rsp)
	call	cudaMalloc@PLT
	movl	$4, %esi
	movl	$4, %edi
	call	calloc@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	movl	$16, %edx
	movq	%rax, %rbx
	call	cudaMemset@PLT
	movl	$1073741825, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	40(%rsp), %rdi
	movl	48(%rsp), %esi
	salq	$2, %rdx
	movl	$1, %ecx
	movq	%rdx, 28(%rsp)
	movl	$1, 36(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z39__device_stub__Z17cuda_conv1d_naivePfS_PfS_
.L33:
	call	cudaDeviceSynchronize@PLT
	xorl	%r9d, %r9d
	movq	16(%rsp), %rsi
	movq	%rbx, %rdi
	movl	%r9d, (%rbx)
	movl	$2, %ecx
	movl	$16, %edx
	call	cudaMemcpy@PLT
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%r10d, %r10d
	xorl	%edx, %edx
	movl	$4, %esi
	movq	%rbx, %rdi
	movl	%r10d, 60(%rsp)
	movq	$4, 52(%rsp)
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	movl	$16, %edx
	call	cudaMemset@PLT
	movl	36(%rsp), %ecx
	movl	48(%rsp), %esi
	xorl	%r9d, %r9d
	movq	28(%rsp), %rdx
	movq	40(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z40__device_stub__Z18cuda_conv1d_sharedPfS_PfS_
.L34:
	call	cudaDeviceSynchronize@PLT
	xorl	%edx, %edx
	movq	16(%rsp), %rsi
	movq	%rbx, %rdi
	movl	%edx, (%rbx)
	movl	$2, %ecx
	movl	$16, %edx
	call	cudaMemcpy@PLT
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	movq	%rbx, %rdi
	movl	$4, %esi
	movl	%ecx, 60(%rsp)
	movq	$4, 52(%rsp)
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbp, %rdi
	leaq	84(%rsp), %rbp
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	leaq	.LC7(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$4, %esi
	movl	$64, %edi
	call	calloc@PLT
	movl	$64, %esi
	movq	%rax, %rdi
	movq	%rax, %rbx
	call	_Z8initDataPfi
	movq	%r15, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edi, %edi
	xorl	%edx, %edx
	movabsq	$34359738376, %rsi
	movl	%edi, 60(%rsp)
	movq	%rbx, %rdi
	movq	%rsi, 52(%rsp)
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$256, %esi
	movq	%r14, %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$256, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	84(%rsp), %rdi
	movl	$25, %ecx
	leaq	.LC0(%rip), %rsi
	rep movsl
	movq	%r13, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%r8d, %r8d
	xorl	%edx, %edx
	movq	%rbp, %rdi
	movabsq	$21474836485, %rsi
	movl	%r8d, 60(%rsp)
	movq	%rsi, 52(%rsp)
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movq	%rbp, %rsi
	movl	$100, %edx
	leaq	_ZL5wgt2d(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	$64, %esi
	movq	%r12, %rdi
	call	cudaMalloc@PLT
	movl	$4, %esi
	movl	$16, %edi
	call	calloc@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	movl	$64, %edx
	movq	%rax, %rbp
	call	cudaMemset@PLT
	movl	36(%rsp), %ecx
	movl	48(%rsp), %esi
	xorl	%r9d, %r9d
	movl	$4, 32(%rsp)
	movq	40(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	28(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z39__device_stub__Z17cuda_conv2d_naivePfS_PfS_
.L35:
	call	cudaDeviceSynchronize@PLT
	xorps	%xmm0, %xmm0
	movq	16(%rsp), %rsi
	movq	%rbp, %rdi
	movups	%xmm0, 0(%rbp)
	movl	$2, %ecx
	movl	$64, %edx
	call	cudaMemcpy@PLT
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	xorl	%edx, %edx
	xorl	%eax, %eax
	movq	%rbp, %rdi
	movabsq	$17179869188, %rsi
	movl	%eax, 60(%rsp)
	movq	%rsi, 52(%rsp)
	call	_Z5printPf4dim3
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	call	cudaDeviceReset@PLT
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$200, %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
.LFE3637:
	.size	_Z4testv, .-_Z4testv
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3638:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z4testv
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3638:
	.size	main, .-main
	.section	.rodata.str1.1
.LC10:
	.string	"_Z17cuda_conv2d_naivePfS_"
.LC11:
	.string	"_Z18cuda_conv1d_sharedPfS_"
.LC12:
	.string	"_Z17cuda_conv1d_naivePfS_"
.LC13:
	.string	"wgt"
.LC14:
	.string	"wgt2d"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3670:
	.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	.LC10(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z17cuda_conv2d_naivePfS_(%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	.LC11(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z18cuda_conv1d_sharedPfS_(%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	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC12(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z17cuda_conv1d_naivePfS_(%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	.LC13(%rip), %rdx
	movl	$20, %r9d
	leaq	_ZL3wgt(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$100, %r9d
	leaq	_ZL5wgt2d(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.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
.LFE3670:
	.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	_ZL5wgt2d
	.comm	_ZL5wgt2d,100,32
	.local	_ZL3wgt
	.comm	_ZL3wgt,20,16
	.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:

/**
 * File              : gpu_conv.cu
 * Author            : Xianglan Piao <lanxlpiao@gmail.com>
 * Date              : 2020.06.16
 * Last Modified Date: 2020.07.31
 * Last Modified By  : Xianglan Piao <lanxlpiao@gmail.com>
 * NOTE:             : cuda conv2d
 */


#include <hip/hip_runtime.h>
#include <iostream>

#define ifm_size 8
#define wgt_size 5
#define stride 2
#define padding 2  // For same ifm/ofm size padding = wgt_size / 2
#define ofm_size ((ifm_size + 2 * padding - wgt_size) / stride + 1)

//// constant memory
__constant__ float wgt[wgt_size];

//// global memory
__global__ void cuda_conv1d_naive(float* ifm, float* ofm) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  float temp = 0.0f;
  for (int k = 0; k < wgt_size; k++) {
    int col_offset = col * stride - padding + k;
    if (col_offset >= 0 && col_offset < ifm_size) {
      temp += ifm[col_offset] * wgt[k];
    }
  }

  ofm[col] = temp;
}

//// shared memory
#define BLOCK_size 4
__global__ void cuda_conv1d_shared(float* ifm, float* ofm) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  __shared__ float shared_block[BLOCK_size];

  float temp = 0.0f;
  for (int k = 0; k < wgt_size; k++) {
    int col_offset = col * stride - padding + k;

    shared_block[threadIdx.x] = ifm[col_offset];
    __syncthreads();

    if (col_offset >= 0 && col_offset < ifm_size) {
      temp += shared_block[threadIdx.x] * wgt[k];
    }
  }

  ofm[col] = temp;
}

void initData(float* data, int size) {
  for (int i = 0; i < size; i++) {
    data[i] = i + 1;
  }
}

void print(float* data, dim3 dim) {
  for (int x = 0; x < dim.x; x++) {
    if (dim.y == 0) {
      std::cout << data[x] << ",  ";
    } else {
      for (int y = 0; y < dim.y; y++) {
        std::cout << data[x * dim.x + y] << ",  ";
      }
      std::cout << std::endl;
    }
  }
  std::cout << std::endl;
}

//// constant memory
__constant__ float wgt2d[wgt_size * wgt_size];

//// global memory
__global__ void cuda_conv2d_naive(float* ifm, float* ofm) {
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  int row = blockIdx.y * blockDim.y + threadIdx.y;

  float temp = 0.0f;
  for (int m = 0; m < wgt_size; m++) {
    for (int n = 0; n < wgt_size; n++) {
      int col_offset = col * stride - padding + m;
      int row_offset = row * stride - padding + n;
      if ((col_offset >= 0 && col_offset < ifm_size) &&
          (row_offset >= 0 && row_offset < ifm_size)) {
        temp +=
            ifm[row_offset * ifm_size + col_offset] * wgt2d[m * wgt_size + n];
      }
    }

    ofm[row * ofm_size + col] = temp;
  }
}

void test() {
  dim3 block(0);
  dim3 grid(0);

  float* h_ifm = NULL;
  float* d_ifm = NULL;
  float* d_ofm = NULL;
  float* h_ofm = NULL;

  //// 1D convolution
  std::cout << "\n--- 1D convolution ---\n" << std::endl;
  std::cout << "ifm: " << std::endl;
  h_ifm = (float*)malloc(ifm_size * sizeof(float));
  initData(h_ifm, ifm_size);
  print(h_ifm, dim3(ifm_size, 0, 0));
  std::cout << std::endl;

  hipMalloc((void**)&d_ifm, ifm_size * sizeof(float));
  hipMemcpy(d_ifm, h_ifm, ifm_size * sizeof(float), hipMemcpyHostToDevice);

  float h_wgt[wgt_size] = {1, 2, 4, 2, 1};
  std::cout << "wgt: " << std::endl;
  print(h_wgt, dim3(wgt_size, 0, 0));
  std::cout << std::endl;
  hipMemcpyToSymbol(HIP_SYMBOL(wgt), &h_wgt, sizeof(wgt));

  block.x = BLOCK_size;
  grid.x = ofm_size / block.x;

  hipMalloc((void**)&d_ofm, ofm_size * sizeof(float));
  h_ofm = (float*)calloc(ofm_size, sizeof(float));

  //// using global memory
  hipMemset(d_ofm, 0, ofm_size * sizeof(float));
  cuda_conv1d_naive<<<grid, block>>>(d_ifm, d_ofm);
  hipDeviceSynchronize();

  memset(h_ofm, 0, ofm_size);
  hipMemcpy(h_ofm, d_ofm, ofm_size * sizeof(float), hipMemcpyDeviceToHost);
  std::cout << "ofm: " << std::endl;
  print(h_ofm, dim3(ofm_size, 0, 0));
  std::cout << std::endl;

  //// using shared memory
  hipMemset(d_ofm, 0, ofm_size * sizeof(float));
  cuda_conv1d_shared<<<grid, block>>>(d_ifm, d_ofm);
  hipDeviceSynchronize();

  memset(h_ofm, 0, ofm_size);
  hipMemcpy(h_ofm, d_ofm, ofm_size * sizeof(float), hipMemcpyDeviceToHost);
  std::cout << "ofm: " << std::endl;
  print(h_ofm, dim3(ofm_size, 0, 0));
  std::cout << std::endl;

  //// free
  hipFree(d_ifm);
  hipFree(d_ofm);
  free(h_ifm);
  free(h_ofm);

  //// 2D convolution
  std::cout << "\n--- 2D convolution ---\n" << std::endl;
  h_ifm = (float*)calloc(ifm_size * ifm_size, sizeof(float));
  initData(h_ifm, ifm_size * ifm_size);
  std::cout << "ifm: " << std::endl;
  print(h_ifm, dim3(ifm_size, ifm_size, 0));
  std::cout << std::endl;

  hipMalloc((void**)&d_ifm, ifm_size * ifm_size * sizeof(float));
  hipMemcpy(d_ifm,
             h_ifm,
             ifm_size * ifm_size * sizeof(float),
             hipMemcpyHostToDevice);

  float h_wgt2d[wgt_size * wgt_size] = {1, 1, 1, 1, 1,   //
                                        1, 2, 2, 2, 1,   //
                                        1, 2, 4, 2, 1,   //
                                        1, 2, 2, 2, 1,   //
                                        1, 1, 1, 1, 1};  //
  std::cout << "wgt: " << std::endl;
  print(h_wgt2d, dim3(wgt_size, wgt_size, 0));
  std::cout << std::endl;
  hipMemcpyToSymbol(HIP_SYMBOL(wgt2d), &h_wgt2d, sizeof(wgt2d));

  block.x = BLOCK_size;
  block.y = BLOCK_size;
  grid.x = ofm_size / block.x;
  grid.y = ofm_size / block.y;

  hipMalloc((void**)&d_ofm, ofm_size * ofm_size * sizeof(float));
  h_ofm = (float*)calloc(ofm_size * ofm_size, sizeof(float));
  hipMemset(d_ofm, 0, ofm_size * ofm_size * sizeof(float));

  cuda_conv2d_naive<<<grid, block>>>(d_ifm, d_ofm);
  hipDeviceSynchronize();

  memset(h_ofm, 0, ofm_size * ofm_size);
  hipMemcpy(h_ofm,
             d_ofm,
             ofm_size * ofm_size * sizeof(float),
             hipMemcpyDeviceToHost);
  std::cout << "ofm: " << std::endl;
  print(h_ofm, dim3(ofm_size, ofm_size, 0));
  std::cout << std::endl;

  //// free
  hipFree(d_ifm);
  hipFree(d_ofm);
  free(h_ifm);
  free(h_ofm);

  //// reset
  hipDeviceReset();
}

int main(void) {
  test();
  return 0;
}

	.text
	.file	"test_gpu_conv.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z32__device_stub__cuda_conv1d_naivePfS_ # -- Begin function _Z32__device_stub__cuda_conv1d_naivePfS_
	.type	_Z32__device_stub__cuda_conv1d_naivePfS_,@function
_Z32__device_stub__cuda_conv1d_naivePfS_: # @_Z32__device_stub__cuda_conv1d_naivePfS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z17cuda_conv1d_naivePfS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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__cuda_conv1d_naivePfS_, .Lfunc_end0-_Z32__device_stub__cuda_conv1d_naivePfS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z33__device_stub__cuda_conv1d_sharedPfS_ # -- Begin function _Z33__device_stub__cuda_conv1d_sharedPfS_
	.type	_Z33__device_stub__cuda_conv1d_sharedPfS_,@function
_Z33__device_stub__cuda_conv1d_sharedPfS_: # @_Z33__device_stub__cuda_conv1d_sharedPfS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z18cuda_conv1d_sharedPfS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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	_Z33__device_stub__cuda_conv1d_sharedPfS_, .Lfunc_end1-_Z33__device_stub__cuda_conv1d_sharedPfS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z8initDataPfi                  # -- Begin function _Z8initDataPfi
	.type	_Z8initDataPfi,@function
_Z8initDataPfi:                         # @_Z8initDataPfi
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB2_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%esi, %eax
	xorl	%ecx, %ecx
.LBB2_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	leaq	1(%rcx), %rdx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%edx, %xmm0
	movss	%xmm0, (%rdi,%rcx,4)
	movq	%rdx, %rcx
	cmpq	%rdx, %rax
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	retq
.Lfunc_end2:
	.size	_Z8initDataPfi, .Lfunc_end2-_Z8initDataPfi
	.cfi_endproc
                                        # -- End function
	.globl	_Z5printPf4dim3                 # -- Begin function _Z5printPf4dim3
	.type	_Z5printPf4dim3,@function
_Z5printPf4dim3:                        # @_Z5printPf4dim3
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, (%rsp)                    # 8-byte Spill
	testl	%esi, %esi
	je	.LBB3_7
# %bb.1:                                # %.lr.ph
	movq	%rdi, %r14
	movq	(%rsp), %rbp                    # 8-byte Reload
	movq	%rbp, %rax
	shrq	$32, %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	%ebp, %eax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	xorl	%ebx, %ebx
	shrq	$32, %rbp
	xorl	%r12d, %r12d
.LBB3_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_4 Depth 2
	movq	16(%rsp), %r15                  # 8-byte Reload
	movq	%rbx, %r13
	testq	%rbp, %rbp
	jne	.LBB3_4
# %bb.3:                                #   in Loop: Header=BB3_2 Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14,%r12,4), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movl	$.L.str, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_6
.LBB3_4:                                #   Parent Loop BB3_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%r13d, %eax
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r14,%rax,4), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movl	$.L.str, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%r13
	decq	%r15
	jne	.LBB3_4
# %bb.5:                                #   in Loop: Header=BB3_2 Depth=1
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	movl	$_ZSt4cout, %eax
	addq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB3_6:                                #   in Loop: Header=BB3_2 Depth=1
	incq	%r12
	addq	(%rsp), %rbx                    # 8-byte Folded Reload
	cmpq	8(%rsp), %r12                   # 8-byte Folded Reload
	jne	.LBB3_2
.LBB3_7:                                # %._crit_edge
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %ecx
	movq	-24(%rax), %rdi
	addq	%rcx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	addq	$24, %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
	jmp	_ZNSo5flushEv                   # TAILCALL
.Lfunc_end3:
	.size	_Z5printPf4dim3, .Lfunc_end3-_Z5printPf4dim3
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__cuda_conv2d_naivePfS_ # -- Begin function _Z32__device_stub__cuda_conv2d_naivePfS_
	.type	_Z32__device_stub__cuda_conv2d_naivePfS_,@function
_Z32__device_stub__cuda_conv2d_naivePfS_: # @_Z32__device_stub__cuda_conv2d_naivePfS_
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z17cuda_conv2d_naivePfS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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_end4:
	.size	_Z32__device_stub__cuda_conv2d_naivePfS_, .Lfunc_end4-_Z32__device_stub__cuda_conv2d_naivePfS_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z4testv
.LCPI5_0:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40800000                      # float 4
	.long	0x40000000                      # float 2
	.text
	.globl	_Z4testv
	.type	_Z4testv,@function
_Z4testv:                               # @_Z4testv
	.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	$152, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%ebx, %ebx
	movq	%rbx, 8(%rsp)
	movq	%rbx, (%rsp)
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$24, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$32, %edi
	callq	malloc
	movq	%rax, %r14
.LBB5_1:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	leaq	1(%rbx), %rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	movss	%xmm0, (%r14,%rbx,4)
	movq	%rax, %rbx
	cmpq	$8, %rax
	jne	.LBB5_1
# %bb.2:                                # %_Z8initDataPfi.exit
	movabsq	$4294967297, %rbx               # imm = 0x100000001
	movl	$8, %esi
	movq	%r14, %rdi
	callq	_Z5printPf4dim3
	movl	$_ZSt4cout, %r13d
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	leaq	8(%rsp), %r15
	movl	$32, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$32, %edx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movaps	.LCPI5_0(%rip), %xmm0           # xmm0 = [1.0E+0,2.0E+0,4.0E+0,2.0E+0]
	leaq	16(%rsp), %r15
	movaps	%xmm0, (%r15)
	movl	$1065353216, 16(%r15)           # imm = 0x3F800000
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$5, %esi
	movq	%r15, %rdi
	callq	_Z5printPf4dim3
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$wgt, %edi
	movl	$20, %edx
	movq	%r15, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movq	%rsp, %r12
	movl	$16, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	movl	$4, %edi
	movl	$4, %esi
	callq	calloc
	movq	%rax, %r15
	movq	(%r12), %rdi
	movl	$16, %edx
	xorl	%esi, %esi
	callq	hipMemset
	leaq	3(%rbx), %r12
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_4
# %bb.3:
	movq	8(%rsp), %rdi
	movq	(%rsp), %rsi
	callq	_Z32__device_stub__cuda_conv1d_naivePfS_
.LBB5_4:
	callq	hipDeviceSynchronize
	movq	(%rsp), %rsi
	movl	$16, %edx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$4, %esi
	movq	%r15, %rdi
	callq	_Z5printPf4dim3
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%rsp), %rdi
	movl	$16, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r12, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_6
# %bb.5:
	movq	8(%rsp), %rdi
	movq	(%rsp), %rsi
	callq	_Z33__device_stub__cuda_conv1d_sharedPfS_
.LBB5_6:
	callq	hipDeviceSynchronize
	movl	$0, (%r15)
	movq	(%rsp), %rsi
	movl	$16, %edx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %r12d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$4, %esi
	movq	%r15, %rdi
	callq	_Z5printPf4dim3
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$24, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$64, %edi
	movl	$4, %esi
	callq	calloc
	movq	%rax, %r14
	xorl	%eax, %eax
.LBB5_7:                                # %.lr.ph.i57
                                        # =>This Inner Loop Header: Depth=1
	leaq	1(%rax), %rcx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%ecx, %xmm0
	movss	%xmm0, (%r14,%rax,4)
	movq	%rcx, %rax
	cmpq	$64, %rcx
	jne	.LBB5_7
# %bb.8:                                # %_Z8initDataPfi.exit61
	movabsq	$17179869188, %r15              # imm = 0x400000004
	movl	$_ZSt4cout, %ebp
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movabsq	$34359738376, %rsi              # imm = 0x800000008
	movq	%r14, %rdi
	callq	_Z5printPf4dim3
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	leaq	8(%rsp), %r12
	movl	$256, %esi                      # imm = 0x100
	movq	%r12, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$256, %edx                      # imm = 0x100
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	48(%rsp), %r12
	movl	$12, %ecx
	movl	$.L__const._Z4testv.h_wgt2d, %esi
	movq	%r12, %rdi
	rep;movsq (%rsi), %es:(%rdi)
	movl	$1065353216, 96(%r12)           # imm = 0x3F800000
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movabsq	$21474836485, %rsi              # imm = 0x500000005
	movq	%r12, %rdi
	callq	_Z5printPf4dim3
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$wgt2d, %edi
	movl	$100, %edx
	movq	%r12, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movq	%rsp, %r13
	movl	$64, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	movl	$16, %edi
	movl	$4, %esi
	callq	calloc
	movq	%rax, %r12
	movq	(%r13), %rdi
	movl	$64, %edx
	xorl	%esi, %esi
	callq	hipMemset
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	%r15, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_10
# %bb.9:
	movq	8(%rsp), %rdi
	movq	(%rsp), %rsi
	callq	_Z32__device_stub__cuda_conv2d_naivePfS_
.LBB5_10:
	callq	hipDeviceSynchronize
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%r12)
	movq	(%rsp), %rsi
	movl	$64, %edx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%r12, %rdi
	movq	%r15, %rsi
	callq	_Z5printPf4dim3
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbp, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	%r14, %rdi
	callq	free
	movq	%r12, %rdi
	callq	free
	callq	hipDeviceReset
	addq	$152, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z4testv, .Lfunc_end5-_Z4testv
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	_Z4testv
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	main, .Lfunc_end6-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	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17cuda_conv1d_naivePfS_, %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	$_Z18cuda_conv1d_sharedPfS_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17cuda_conv2d_naivePfS_, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$wgt, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movl	$20, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$wgt2d, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movl	$100, %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_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	wgt,@object                     # @wgt
	.local	wgt
	.comm	wgt,20,16
	.type	_Z17cuda_conv1d_naivePfS_,@object # @_Z17cuda_conv1d_naivePfS_
	.section	.rodata,"a",@progbits
	.globl	_Z17cuda_conv1d_naivePfS_
	.p2align	3, 0x0
_Z17cuda_conv1d_naivePfS_:
	.quad	_Z32__device_stub__cuda_conv1d_naivePfS_
	.size	_Z17cuda_conv1d_naivePfS_, 8

	.type	_Z18cuda_conv1d_sharedPfS_,@object # @_Z18cuda_conv1d_sharedPfS_
	.globl	_Z18cuda_conv1d_sharedPfS_
	.p2align	3, 0x0
_Z18cuda_conv1d_sharedPfS_:
	.quad	_Z33__device_stub__cuda_conv1d_sharedPfS_
	.size	_Z18cuda_conv1d_sharedPfS_, 8

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

	.type	wgt2d,@object                   # @wgt2d
	.local	wgt2d
	.comm	wgt2d,100,16
	.type	_Z17cuda_conv2d_naivePfS_,@object # @_Z17cuda_conv2d_naivePfS_
	.section	.rodata,"a",@progbits
	.globl	_Z17cuda_conv2d_naivePfS_
	.p2align	3, 0x0
_Z17cuda_conv2d_naivePfS_:
	.quad	_Z32__device_stub__cuda_conv2d_naivePfS_
	.size	_Z17cuda_conv2d_naivePfS_, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"\n--- 1D convolution ---\n"
	.size	.L.str.1, 25

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"ifm: "
	.size	.L.str.2, 6

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"wgt: "
	.size	.L.str.3, 6

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"ofm: "
	.size	.L.str.4, 6

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\n--- 2D convolution ---\n"
	.size	.L.str.5, 25

	.type	.L__const._Z4testv.h_wgt2d,@object # @__const._Z4testv.h_wgt2d
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const._Z4testv.h_wgt2d:
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40000000                      # float 2
	.long	0x40000000                      # float 2
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40800000                      # float 4
	.long	0x40000000                      # float 2
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40000000                      # float 2
	.long	0x40000000                      # float 2
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.long	0x3f800000                      # float 1
	.size	.L__const._Z4testv.h_wgt2d, 100

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z17cuda_conv1d_naivePfS_"
	.size	.L__unnamed_1, 26

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z18cuda_conv1d_sharedPfS_"
	.size	.L__unnamed_2, 27

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z17cuda_conv2d_naivePfS_"
	.size	.L__unnamed_3, 26

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"wgt"
	.size	.L__unnamed_4, 4

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"wgt2d"
	.size	.L__unnamed_5, 6

	.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__cuda_conv1d_naivePfS_
	.addrsig_sym _Z33__device_stub__cuda_conv1d_sharedPfS_
	.addrsig_sym _Z32__device_stub__cuda_conv2d_naivePfS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym wgt
	.addrsig_sym _Z17cuda_conv1d_naivePfS_
	.addrsig_sym _Z18cuda_conv1d_sharedPfS_
	.addrsig_sym _ZSt4cout
	.addrsig_sym wgt2d
	.addrsig_sym _Z17cuda_conv2d_naivePfS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17cuda_conv1d_naivePfS_ ; -- Begin function _Z17cuda_conv1d_naivePfS_
	.globl	_Z17cuda_conv1d_naivePfS_
	.p2align	8
	.type	_Z17cuda_conv1d_naivePfS_,@function
_Z17cuda_conv1d_naivePfS_:              ; @_Z17cuda_conv1d_naivePfS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x1c
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mov_b32_e32 v3, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	v_mov_b32_e32 v0, v3
	s_mov_b64 s[4:5], 0
	s_delay_alu instid0(VALU_DEP_2)
	v_lshl_add_u32 v2, v1, 1, -2
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s6, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e32 8, v2
	s_cbranch_execz .LBB0_3
; %bb.2:                                ;   in Loop: Header=BB0_1 Depth=1
	v_lshlrev_b64 v[4:5], 2, v[2:3]
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, wgt@rel32@lo+4
	s_addc_u32 s9, s9, wgt@rel32@hi+12
	s_add_u32 s8, s4, s8
	s_addc_u32 s9, s5, s9
	s_load_b32 s7, s[8:9], 0x0
	v_add_co_u32 v4, vcc_lo, s0, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s1, v5, vcc_lo
	global_load_b32 v4, v[4:5], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_fmac_f32_e32 v0, s7, v4
.LBB0_3:                                ;   in Loop: Header=BB0_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s6
	v_add_nc_u32_e32 v2, 1, v2
	s_add_u32 s4, s4, 4
	s_addc_u32 s5, s5, 0
	s_cmp_eq_u32 s4, 20
	s_cbranch_scc0 .LBB0_1
; %bb.4:
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s2, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s3, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17cuda_conv1d_naivePfS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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 6
		.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	_Z17cuda_conv1d_naivePfS_, .Lfunc_end0-_Z17cuda_conv1d_naivePfS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 232
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; 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	_Z18cuda_conv1d_sharedPfS_ ; -- Begin function _Z18cuda_conv1d_sharedPfS_
	.globl	_Z18cuda_conv1d_sharedPfS_
	.p2align	8
	.type	_Z18cuda_conv1d_sharedPfS_,@function
_Z18cuda_conv1d_sharedPfS_:             ; @_Z18cuda_conv1d_sharedPfS_
; %bb.0:
	s_load_b32 s2, s[0:1], 0x1c
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_load_b128 s[0:3], s[0:1], 0x0
	v_lshlrev_b32_e32 v2, 2, v0
	v_lshlrev_b32_e32 v3, 1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, v4, s0
	v_add_co_ci_u32_e32 v5, vcc_lo, s1, v5, vcc_lo
	s_mov_b64 s[0:1], 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_u32 v4, vcc_lo, v0, -8
	v_dual_mov_b32 v0, 0 :: v_dual_add_nc_u32 v3, -2, v3
	v_add_co_ci_u32_e32 v5, vcc_lo, -1, v5, vcc_lo
.LBB1_1:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v6, vcc_lo, v4, s0
	v_add_co_ci_u32_e32 v7, vcc_lo, s1, v5, vcc_lo
	s_mov_b32 s4, exec_lo
	global_load_b32 v6, v[6:7], off
	s_waitcnt vmcnt(0)
	ds_store_b32 v2, v6
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cmpx_gt_u32_e32 8, v3
	s_cbranch_execz .LBB1_3
; %bb.2:                                ;   in Loop: Header=BB1_1 Depth=1
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, wgt@rel32@lo+4
	s_addc_u32 s7, s7, wgt@rel32@hi+12
	s_add_u32 s6, s0, s6
	s_addc_u32 s7, s1, s7
	ds_load_b32 v6, v2
	s_load_b32 s5, s[6:7], 0x0
	s_waitcnt lgkmcnt(0)
	v_fmac_f32_e32 v0, s5, v6
.LBB1_3:                                ;   in Loop: Header=BB1_1 Depth=1
	s_or_b32 exec_lo, exec_lo, s4
	v_add_nc_u32_e32 v3, 1, v3
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	s_cmp_eq_u32 s0, 20
	s_cbranch_scc0 .LBB1_1
; %bb.4:
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_u32 v1, vcc_lo, s2, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s3, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z18cuda_conv1d_sharedPfS_
		.amdhsa_group_segment_fixed_size 16
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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	_Z18cuda_conv1d_sharedPfS_, .Lfunc_end1-_Z18cuda_conv1d_sharedPfS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 304
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 16 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	_Z17cuda_conv2d_naivePfS_ ; -- Begin function _Z17cuda_conv2d_naivePfS_
	.globl	_Z17cuda_conv2d_naivePfS_
	.p2align	8
	.type	_Z17cuda_conv2d_naivePfS_,@function
_Z17cuda_conv2d_naivePfS_:              ; @_Z17cuda_conv2d_naivePfS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x1c
	s_load_b128 s[4:7], s[0:1], 0x0
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v0, v0, 10, 10
	s_mov_b32 s8, 0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s3, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, s14, s2, v[1:2]
	v_mad_u64_u32 v[3:4], null, s15, s3, v[0:1]
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, wgt2d@rel32@lo+4
	s_addc_u32 s3, s3, wgt2d@rel32@hi+12
	v_lshl_add_u32 v0, v3, 2, v2
	v_lshlrev_b32_e32 v2, 1, v2
	v_lshlrev_b32_e32 v6, 4, v3
	v_lshl_add_u32 v4, v3, 1, -2
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
	v_lshlrev_b64 v[7:8], 2, v[0:1]
	v_mov_b32_e32 v1, 0
	v_add_nc_u32_e32 v5, -2, v2
	v_add3_u32 v6, v6, v2, 0xffffffee
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v2, vcc_lo, s6, v7
	v_mov_b32_e32 v7, 0
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v8, vcc_lo
.LBB2_1:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB2_2 Depth 2
	v_add_nc_u32_e32 v0, s8, v5
	v_mov_b32_e32 v8, v4
	s_mov_b64 s[6:7], 0
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_gt_u32_e32 vcc_lo, 8, v0
	v_mov_b32_e32 v0, v6
.LBB2_2:                                ;   Parent Loop BB2_1 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cmp_lt_i32_e64 s0, -1, v8
	v_cmp_gt_i32_e64 s1, 8, v8
	s_and_b32 s0, vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, s1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB2_4
; %bb.3:                                ;   in Loop: Header=BB2_2 Depth=2
	v_lshlrev_b64 v[9:10], 2, v[0:1]
	s_add_u32 s10, s2, s6
	s_addc_u32 s11, s3, s7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v9, s0, s4, v9
	v_add_co_ci_u32_e64 v10, s0, s5, v10, s0
	s_load_b32 s0, s[10:11], 0x0
	global_load_b32 v9, v[9:10], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_fmac_f32_e32 v7, s0, v9
.LBB2_4:                                ;   in Loop: Header=BB2_2 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	v_add_nc_u32_e32 v0, 8, v0
	v_add_nc_u32_e32 v8, 1, v8
	s_add_u32 s6, s6, 4
	s_addc_u32 s7, s7, 0
	s_cmp_eq_u32 s6, 20
	s_cbranch_scc0 .LBB2_2
; %bb.5:                                ;   in Loop: Header=BB2_1 Depth=1
	s_add_i32 s8, s8, 1
	v_add_nc_u32_e32 v6, 1, v6
	s_add_u32 s2, s2, 20
	s_addc_u32 s3, s3, 0
	s_cmp_eq_u32 s8, 5
	global_store_b32 v[2:3], v7, off
	s_cbranch_scc0 .LBB2_1
; %bb.6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17cuda_conv2d_naivePfS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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 11
		.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_end2:
	.size	_Z17cuda_conv2d_naivePfS_, .Lfunc_end2-_Z17cuda_conv2d_naivePfS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 376
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; 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
	.protected	wgt                     ; @wgt
	.type	wgt,@object
	.section	.bss,"aw",@nobits
	.globl	wgt
	.p2align	4, 0x0
wgt:
	.zero	20
	.size	wgt, 20

	.protected	wgt2d                   ; @wgt2d
	.type	wgt2d,@object
	.globl	wgt2d
	.p2align	4, 0x0
wgt2d:
	.zero	100
	.size	wgt2d, 100

	.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 wgt
	.addrsig_sym wgt2d
	.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:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17cuda_conv1d_naivePfS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17cuda_conv1d_naivePfS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .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:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 16
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z18cuda_conv1d_sharedPfS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z18cuda_conv1d_sharedPfS_.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:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17cuda_conv2d_naivePfS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17cuda_conv2d_naivePfS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
