#if GOOGLE_CUDA
#define EIGEN_USE_GPU

__global__ void default_function_kernel0(const float* __restrict__ Data,
    const float* __restrict__ K0,
    const float* __restrict__ K1,
    const float* __restrict__ K2,
    float* __restrict__ Output) {

   float Output_local[1];
  __shared__ float Data_shared[128];
  __shared__ float K0_shared[32];
  __shared__ float K1_shared[64];
  __shared__ float K2_shared[1];
  Output_local[0] = 0.000000e+00f;
  for (int rr_outer = 0; rr_outer < 137; ++rr_outer) {
    for (int rs0_outer = 0; rs0_outer < 2; ++rs0_outer) {
      for (int rs2_outer = 0; rs2_outer < 16; ++rs2_outer) {
        __syncthreads();
        for (int ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner = 0; ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner < 8; ++ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner) {
          Data_shared[(((((int)threadIdx.z) * 32) + (((int)threadIdx.y) * 8)) + ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner)] = Data[((((((((((int)threadIdx.z) * 32) + (((int)threadIdx.y) * 8)) + ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner) / 128) * 4096) + (rs0_outer * 2048)) + (((((((int)threadIdx.z) * 32) + (((int)threadIdx.y) * 8)) + ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner) % 128) * 16)) + rs2_outer)];
        }
        for (int ax0_ax1_ax2_fused_fused_inner_inner_inner = 0; ax0_ax1_ax2_fused_fused_inner_inner_inner < 2; ++ax0_ax1_ax2_fused_fused_inner_inner_inner) {
          K0_shared[(((((int)threadIdx.z) * 8) + (((int)threadIdx.y) * 2)) + ax0_ax1_ax2_fused_fused_inner_inner_inner)] = K0[(((((rs0_outer * 4384) + (((int)threadIdx.z) * 1096)) + (((int)threadIdx.y) * 274)) + (ax0_ax1_ax2_fused_fused_inner_inner_inner * 137)) + rr_outer)];
        }
        for (int ax0_ax1_ax2_fused_fused_inner_inner_inner1 = 0; ax0_ax1_ax2_fused_fused_inner_inner_inner1 < 4; ++ax0_ax1_ax2_fused_fused_inner_inner_inner1) {
          K1_shared[(((((int)threadIdx.z) * 16) + (((int)threadIdx.y) * 4)) + ax0_ax1_ax2_fused_fused_inner_inner_inner1)] = K1[((((((int)threadIdx.z) * 2192) + (((int)threadIdx.y) * 548)) + (ax0_ax1_ax2_fused_fused_inner_inner_inner1 * 137)) + rr_outer)];
        }
        if (((int)threadIdx.y) < (1 - ((int)threadIdx.z))) {
          if (((int)threadIdx.y) < 1) {
            if (((int)threadIdx.y) < ((16 - rs2_outer) - ((int)threadIdx.z))) {
              K2_shared[(((int)threadIdx.y) + ((int)threadIdx.z))] = K2[(((((((int)threadIdx.y) * 548) + (((int)threadIdx.z) * 548)) + (rs2_outer * 548)) + (((int)blockIdx.x) * 137)) + rr_outer)];
            }
          }
        }
        __syncthreads();
        for (int rs0_inner = 0; rs0_inner < 8; ++rs0_inner) {
          for (int rs1_inner = 0; rs1_inner < 16; ++rs1_inner) {
            Output_local[0] = (Output_local[0] + (((Data_shared[((rs0_inner * 16) + rs1_inner)] * K0_shared[((rs0_inner * 4) + ((int)threadIdx.z))]) * K1_shared[((rs1_inner * 4) + ((int)threadIdx.y))]) * K2_shared[0]));
          }
        }
      }
    }
  }
  Output[(((((int)threadIdx.z) * 16) + (((int)threadIdx.y) * 4)) + ((int)blockIdx.x))] = Output_local[0];
}

void DenseCpKernelLauncher(const float* U, const float* K0,
    const float* K1, const float* KC, float* V){

  dim3 gridDim0(4, 1, 1);
  dim3 blockDim0(1, 4, 4);

  default_function_kernel0<<<gridDim0, blockDim0>>>(U, K0, K1, KC, V);
  cudaDeviceSynchronize();

}

#endif

	.file	"tmpxft_002216f1_00000000-6_dense_cp.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.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:


#include <hip/hip_runtime.h>
#if GOOGLE_CUDA
#define EIGEN_USE_GPU

__global__ void default_function_kernel0(const float* __restrict__ Data,
    const float* __restrict__ K0,
    const float* __restrict__ K1,
    const float* __restrict__ K2,
    float* __restrict__ Output) {

   float Output_local[1];
  __shared__ float Data_shared[128];
  __shared__ float K0_shared[32];
  __shared__ float K1_shared[64];
  __shared__ float K2_shared[1];
  Output_local[0] = 0.000000e+00f;
  for (int rr_outer = 0; rr_outer < 137; ++rr_outer) {
    for (int rs0_outer = 0; rs0_outer < 2; ++rs0_outer) {
      for (int rs2_outer = 0; rs2_outer < 16; ++rs2_outer) {
        __syncthreads();
        for (int ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner = 0; ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner < 8; ++ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner) {
          Data_shared[(((((int)threadIdx.z) * 32) + (((int)threadIdx.y) * 8)) + ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner)] = Data[((((((((((int)threadIdx.z) * 32) + (((int)threadIdx.y) * 8)) + ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner) / 128) * 4096) + (rs0_outer * 2048)) + (((((((int)threadIdx.z) * 32) + (((int)threadIdx.y) * 8)) + ax0_ax1_ax2_ax3_fused_fused_fused_inner_inner_inner) % 128) * 16)) + rs2_outer)];
        }
        for (int ax0_ax1_ax2_fused_fused_inner_inner_inner = 0; ax0_ax1_ax2_fused_fused_inner_inner_inner < 2; ++ax0_ax1_ax2_fused_fused_inner_inner_inner) {
          K0_shared[(((((int)threadIdx.z) * 8) + (((int)threadIdx.y) * 2)) + ax0_ax1_ax2_fused_fused_inner_inner_inner)] = K0[(((((rs0_outer * 4384) + (((int)threadIdx.z) * 1096)) + (((int)threadIdx.y) * 274)) + (ax0_ax1_ax2_fused_fused_inner_inner_inner * 137)) + rr_outer)];
        }
        for (int ax0_ax1_ax2_fused_fused_inner_inner_inner1 = 0; ax0_ax1_ax2_fused_fused_inner_inner_inner1 < 4; ++ax0_ax1_ax2_fused_fused_inner_inner_inner1) {
          K1_shared[(((((int)threadIdx.z) * 16) + (((int)threadIdx.y) * 4)) + ax0_ax1_ax2_fused_fused_inner_inner_inner1)] = K1[((((((int)threadIdx.z) * 2192) + (((int)threadIdx.y) * 548)) + (ax0_ax1_ax2_fused_fused_inner_inner_inner1 * 137)) + rr_outer)];
        }
        if (((int)threadIdx.y) < (1 - ((int)threadIdx.z))) {
          if (((int)threadIdx.y) < 1) {
            if (((int)threadIdx.y) < ((16 - rs2_outer) - ((int)threadIdx.z))) {
              K2_shared[(((int)threadIdx.y) + ((int)threadIdx.z))] = K2[(((((((int)threadIdx.y) * 548) + (((int)threadIdx.z) * 548)) + (rs2_outer * 548)) + (((int)blockIdx.x) * 137)) + rr_outer)];
            }
          }
        }
        __syncthreads();
        for (int rs0_inner = 0; rs0_inner < 8; ++rs0_inner) {
          for (int rs1_inner = 0; rs1_inner < 16; ++rs1_inner) {
            Output_local[0] = (Output_local[0] + (((Data_shared[((rs0_inner * 16) + rs1_inner)] * K0_shared[((rs0_inner * 4) + ((int)threadIdx.z))]) * K1_shared[((rs1_inner * 4) + ((int)threadIdx.y))]) * K2_shared[0]));
          }
        }
      }
    }
  }
  Output[(((((int)threadIdx.z) * 16) + (((int)threadIdx.y) * 4)) + ((int)blockIdx.x))] = Output_local[0];
}

void DenseCpKernelLauncher(const float* U, const float* K0,
    const float* K1, const float* KC, float* V){

  dim3 gridDim0(4, 1, 1);
  dim3 blockDim0(1, 4, 4);

  default_function_kernel0<<<gridDim0, blockDim0>>>(U, K0, K1, KC, V);
  hipDeviceSynchronize();

}

#endif

	.text
	.file	"dense_cp.hip"
	.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 __hip_cuid_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
