-
-
Save ericcano/75a37585bdaa683723ce1d218925305f to your computer and use it in GitHub Desktop.
PTX for commit 7cca04e3d3d7e046d81f8b12d1db39318d28ce51
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// | |
// 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