Skip to content

Instantly share code, notes, and snippets.

@ericcano

ericcano/testSoAGPU.cu.ptx Secret

Created July 1, 2021 13:49
Show Gist options
  • Save ericcano/75a37585bdaa683723ce1d218925305f to your computer and use it in GitHub Desktop.
Save ericcano/75a37585bdaa683723ce1d218925305f to your computer and use it in GitHub Desktop.
Download ZIP
PTX for commit 7cca04e3d3d7e046d81f8b12d1db39318d28ce51
Raw
testSoAGPU.cu.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-29618528
// Cuda compilation tools, release 11.2, V11.2.152
// Based on NVVM 7.0.1
//
.version 7.2
.target sm_75
.address_size 64
// .globl _ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE
.weak .global .align 8 .b8 _ZTVN7CppUnit17AdditionalMessageE[32];
.weak .global .align 8 .b8 _ZTVN7CppUnit11TestFactoryE[40];
.weak .global .align 8 .b8 _ZTVN7CppUnit16TestSuiteFactoryI10testSoAGPUEE[40];
.visible .entry _ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE(
.param .align 8 .b8 _ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0[80]
)
{
.reg .pred %p<2>;
.reg .b32 %r<7>;
.reg .f64 %fd<5>;
.reg .b64 %rd<28>;
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:76 __global__ void fillSoA(testSoAGPU::SoA soa) {
.loc 1 76 0
Lfunc_begin0:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:76 __global__ void fillSoA(testSoAGPU::SoA soa) {
.loc 1 76 0
ld.param.u64 %rd10, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+64];
ld.param.u64 %rd9, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+56];
ld.param.u64 %rd8, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+48];
ld.param.u64 %rd7, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+40];
ld.param.u64 %rd6, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+32];
ld.param.u64 %rd5, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+24];
ld.param.u64 %rd3, [_ZN69_GLOBAL__N__45_tmpxft_0000d250_00000000_7_testSoAGPU_cpp1_ii_1eb7fc057fillSoAEN10testSoAGPU3SoAE_param_0+8];
Ltmp0:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:77 size_t i = blockIdx.x * blockDim.x + threadIdx.x;
.loc 1 77 14
mov.u32 %r1, %ntid.x;
mov.u32 %r2, %ctaid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r2, %r1, %r3;
cvt.u64.u32 %rd1, %r4;
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:78 if (i >= soa.nElements()) return;
.loc 1 78 5
setp.le.u64 %p1, %rd3, %rd1;
@%p1 bra LBB0_2;
.loc 1 0 5
cvt.u32.u64 %r5, %rd1;
cvta.to.global.u64 %rd12, %rd10;
cvta.to.global.u64 %rd13, %rd9;
cvta.to.global.u64 %rd14, %rd8;
cvta.to.global.u64 %rd15, %rd7;
cvta.to.global.u64 %rd16, %rd5;
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:79 soa[i].x = 11.0 * i;
.loc 1 79 5
cvt.rn.f64.u32 %fd1, %r5;
mul.f64 %fd2, %fd1, 0d4026000000000000;
Ltmp1:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:79 soa[i].x = 11.0 * i;
.loc 1 79 5
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string0, inlined_at 1 79 5
shl.b64 %rd17, %rd1, 3;
add.s64 %rd18, %rd16, %rd17;
st.global.f64 [%rd18], %fd2;
Ltmp2:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:80 soa[i].y = 22.0 * i;
.loc 1 80 5
mul.f64 %fd3, %fd1, 0d4036000000000000;
cvta.to.global.u64 %rd19, %rd6;
Ltmp3:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:80 soa[i].y = 22.0 * i;
.loc 1 80 5
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string0, inlined_at 1 80 5
add.s64 %rd20, %rd19, %rd17;
st.global.f64 [%rd20], %fd3;
Ltmp4:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:81 soa[i].z = 33.0 * i;
.loc 1 81 5
mul.f64 %fd4, %fd1, 0d4040800000000000;
Ltmp5:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:81 soa[i].z = 33.0 * i;
.loc 1 81 5
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string0, inlined_at 1 81 5
add.s64 %rd21, %rd15, %rd17;
st.global.f64 [%rd21], %fd4;
Ltmp6:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:82 soa[i].colour = i;
.loc 1 82 5
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string1, inlined_at 1 82 5
shl.b64 %rd22, %rd1, 1;
add.s64 %rd23, %rd14, %rd22;
st.global.u16 [%rd23], %rd1;
Ltmp7:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:83 soa[i].value = 0x10001 * i;
.loc 1 83 5
mul.lo.s32 %r6, %r5, 65537;
Ltmp8:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:83 soa[i].value = 0x10001 * i;
.loc 1 83 5
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string2, inlined_at 1 83 5
shl.b64 %rd24, %rd1, 2;
add.s64 %rd25, %rd13, %rd24;
st.global.u32 [%rd25], %r6;
Ltmp9:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:84 soa[i].py = &soa[i].y;
.loc 1 84 5
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string3, inlined_at 1 84 5
add.s64 %rd26, %rd12, %rd17;
Ltmp10:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:84 soa[i].py = &soa[i].y;
.loc 1 84 17
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:25 SOA_HOST_DEVICE T* operator& () { return &col_[idx_]; }
.loc 2 25 69, function_name Linfo_string4, inlined_at 1 84 17
add.s64 %rd27, %rd6, %rd17;
Ltmp11:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h:28 SOA_HOST_DEVICE T& operator= (const T2& v) { col_[idx_] = v; return col_[idx_]; }
.loc 2 28 80, function_name Linfo_string3, inlined_at 1 84 5
st.global.u64 [%rd26], %rd27;
Ltmp12:
LBB0_2:
///data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu:85 }
.loc 1 85 3
ret;
Ltmp13:
Lfunc_end0:
}
.file 1 "/data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/test/testSoAGPU.cu"
.file 2 "/data/user/cano/HGCAL/CMSSW_12_0_X_2021-06-23-2300/src/CUDADataFormats/Common/interface/SoAmacros.h"
.section .debug_str
{
Linfo_string0:
.b8 95,90,78,51,99,109,115,52,99,117,100,97,56,83,111,65,86,97,108,117,101,73,100,69,97,83,73,100,69,69,82,100,82,75,84,95,0
Linfo_string1:
.b8 95,90,78,51,99,109,115,52,99,117,100,97,56,83,111,65,86,97,108,117,101,73,116,69,97,83,73,109,69,69,82,116,82,75,84,95,0
Linfo_string2:
.b8 95,90,78,51,99,109,115,52,99,117,100,97,56,83,111,65,86,97,108,117,101,73,105,69,97,83,73,109,69,69,82,105,82,75,84,95,0
Linfo_string3:
.b8 95,90,78,51,99,109,115,52,99,117,100,97,56,83,111,65,86,97,108,117,101,73,80,100,69,97,83,73,83,50,95,69,69,82,83,50,95,82,75,84
.b8 95,0
Linfo_string4:
.b8 95,90,78,51,99,109,115,52,99,117,100,97,56,83,111,65,86,97,108,117,101,73,100,69,97,100,69,118,0
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment