ELF> @@0@8 @ @@@@@88@8@@@TT ``0 ``TT@T@DDPtddd@d@QtdRtd``88/lib64/ld-linux-x86-64.so.2GNUGNU`?` ]%pmH     xIk)x  dQ>chpCTi3w ` @\`libcudart.so.4__gmon_start___Jv_RegisterClassescudaMalloccudaMemcpyToSymbol__cudaRegisterFunction__cudaRegisterFatBinarycudaLaunchcudaMemcpy__cudaUnregisterFatBinarycudaConfigureCallcudaMemset__cudaRegisterVarcudaFreecudaMemcpyFromSymbollibstdc++.so.6__gxx_personality_v0libm.so.6libgcc_s.so.1libc.so.6__stack_chk_failstrcspnstdoutfputsfwritestrchrfprintf__libc_start_main__cxa_atexitCXXABI_1.3GLIBC_2.4GLIBC_2.2.5 ӯk9ii ui `````` `(`0` 8` @` H` P` X```h`p`x`````````HZ5H5 % @% h% h% h% h% h% h% h% hp% h`% h P% h @% h 0% h % h % h% h%z h%r h%j h%b h%Z h%R h%J h%B hp1I^HHPTI`@Hp@Hǀ@HH! HtHÐUHSH=( uK`H" H`HHH9s$fDHH `H H9r H[fff.UH= HtHt`ÐUHSH8H}HuHUHMDE̋uHMHUH]HEAHHH8[UHHEE~UHH H}HuEE]UHH0H}HuUHMEE5UHH0H}HuHUEEUHH H}EEUHH }uEEUHH0H}HuHEHEؾ%H^HEHEH HEHHHE%HEHEHEHE@H3HEHEu HEHEHEЋE}~(H H&@kHEHEEHEHEEHEE%SH@HEЋH HMHHǸ}u+HEZHs HUHHǸ|HEHN HUHHǸWH1 HUHMHHǸa6H @HǸFH HUHHǸ)HEHEUHEHEHEؾ%HHEH}H HEHH7UHSHXHHLLdH%(HE1DžEQHH;uHHHHHHHHHf= H@f=t5H@H@H !@HǸH@fu He H!@F7H@HHHHHHbuUtHHTHHH;HUdH3%(t}HX[UHH H}H}HCEHEHE%HtHEHHEHEi c Hcи8`HH`t 7 HcH% HpH H EEEHEAHƿ `Hm A8`HA8``A@``UHH` Ht6H\ Ht-HP H H= H6 H' UHSHhH}EH}u Hϼ HEH EEH HtH߼ Ht Hü Hu H HEHHEf=oEEEH HE)H]HEHHREtEEEf=t8}u$H& H !@2EHHEE}tREHH}EHHHMEHHUHEH]EIIHHމEHHHMHEHEHHUHEH]EIIHHމEHEHG I HHH;E@Ef=2HEHEA`HH Hں HHغ H)‹׺ HHEHH?H}HHHEH HEHH H)‹ HHEHH?H}HHHEȀ}tHHl n HH<HY [ HH HF ]EHuIIމH$ & HH H ]EH}HuIIމZHEH߹ }t HcHѹ HHh[UHSHxEE!@("@F]HEйH+HEHHUЋMH]EAAH߉uyHܸ H!@X"@HEH]HEHHUMH]EAAH߉uHL HP!@e"@&]HEH ]HEHHUMH]EAAH߉luYH HoHx[UHH۷ H;UH帖@Hٷ @HUHUH@H @HUHUH@Hu @H\UHUHSH(`H9 b@@H# HD$HD$HD$H$AA@@HH@HӶ HD$HD$HD$H$AA@@HH@H HD$HD$HD$H$AA"@"@HHG`H3 D$$AA@@HH`H D$$AA?@?@HH `H D$$AAR@R@HHc`H D$$AA@@HH'H([UHH}uUMHEUHEUPHEUPUHSH8H}HuHUHMDEHEuHMHUH]AHHH8[UHSH8H}HuHUHMDEHEuHMHUH]AHHH8[UHSH8H}HuHUHMDEHEuHMHUH]AHHxH8[UHSH8H}HuHUHMDEH]uHMHUHEAHHH8[UHHH}HEHdÐfffff.Hl$Ld$H-C L%< Ll$Lt$L|$H\$H8L)AIHI#Ht1@LLDAHH9rH\$Hl$Ld$Ll$ Lt$(L|$0H8ÐHٱ Ht H1j1UHSHH Ht`DHHHuH[ÐHoHglobalPrintfBufferprintfBufferPtr%cdiouxXeEfgGaAnpsCorrupt printf buffer data - aborting %%@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@@,@,@,@@,@@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@,@t@,@,@,@,@,@,@,@,@@,@t@t@@@@,@t@,@,@,@,@,@t@t@,@,@@,@t@,@,@t@[%d, %d]: printf buffer overflow No printf headers found at all! -----------------------------------------This is an example of schduling N threads This is an example of schduling M blocks This is an example of schduling M blocks and N threads PU`m8 @threadVsBlockScheduling.cu .version 1.4 .target sm_10, map_f64_to_f32 // compiled with /pkgs/cuda/open64/lib//be // nvopencc 4.0 built on 2011-05-12 //----------------------------------------------------------- // Compiling /tmp/tmpxft_0000751b_00000000-7_threadVsBlockScheduling.cpp3.i (/tmp/ccBI#.eFRvSe) //----------------------------------------------------------- //----------------------------------------------------------- // Options: //----------------------------------------------------------- // Target:ptx, ISA:sm_10, Endian:little, Pointer Size:64 // -O3 (Optimization level) // -g0 (Debug level) // -m2 (Report advisories) //----------------------------------------------------------- .file 1 "" .file 2 "/tmp/tmpxft_0000751b_00000000-6_threadVsBlockScheduling.cudafe2.gpu" .file 3 "cuPrintf.cu" .file 4 "/usr/lib/gcc/x86_64-linux-gnu/4.4.3/include/stddef.h" .file 5 "/usr/local/cuda/bin/../include/crt/device_runtime.h" .file 6 "/usr/local/cuda/bin/../include/host_defines.h" .file 7 "/usr/local/cuda/bin/../include/builtin_types.h" .file 8 "/usr/local/cuda/bin/../include/device_types.h" .file 9 "/usr/local/cuda/bin/../include/driver_types.h" .file 10 "/usr/local/cuda/bin/../include/surface_types.h" .file 11 "/usr/local/cuda/bin/../include/texture_types.h" .file 12 "/usr/local/cuda/bin/../include/vector_types.h" .file 13 "/usr/local/cuda/bin/../include/device_launch_parameters.h" .file 14 "/usr/local/cuda/bin/../include/crt/storage_class.h" .file 15 "/usr/include/bits/types.h" .file 16 "/usr/include/time.h" .file 17 "threadVsBlockScheduling.cu" .file 18 "/usr/local/cuda/bin/../include/common_functions.h" .file 19 "/usr/local/cuda/bin/../include/math_functions.h" .file 20 "/usr/local/cuda/bin/../include/math_constants.h" .file 21 "/usr/local/cuda/bin/../include/device_functions.h" .file 22 "/usr/local/cuda/bin/../include/sm_11_atomic_functions.h" .file 23 "/usr/local/cuda/bin/../include/sm_12_atomic_functions.h" .file 24 "/usr/local/cuda/bin/../include/sm_13_double_functions.h" .file 25 "/usr/local/cuda/bin/../include/sm_20_atomic_functions.h" .file 26 "/usr/local/cuda/bin/../include/sm_20_intrinsics.h" .file 27 "/usr/local/cuda/bin/../include/surface_functions.h" .file 28 "/usr/local/cuda/bin/../include/texture_fetch_functions.h" .file 29 "/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h" .global .align 8 .b8 restrictRules[8]; .const .u64 globalPrintfBuffer = 0; .const .s32 printfBufferLength = 0; .global .u64 printfBufferPtr = 0; .const .align 1 .b8 __constant461[36] = {0x74,0x68,0x72,0x65,0x61,0x64,0x49,0x64,0x78,0x20,0x3c,0x78,0x3a,0x25,0x64,0x3e,0x20,0x62,0x6c,0x6f,0x63,0x6b,0x49,0x64,0x78,0x20,0x20,0x3c,0x78,0x3a,0x25,0x64,0x3e,0xa,0x20,0x0}; .entry _Z23printTest_1_dimensionalv { .reg .u16 %rh<17>; .reg .u32 %r<83>; .reg .u64 %rd<33>; .reg .pred %p<19>; .loc 17 11 0 $LDWbegin__Z23printTest_1_dimensionalv: .loc 17 13 0 ld.global.u64 %rd1, [printfBufferPtr]; mov.u64 %rd2, 0; setp.ne.u64 %p1, %rd1, %rd2; @%p1 bra $Lt_0_19202; .loc 3 164 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_214_9; $Lt_0_19202: mov.u16 %rh1, %nctaid.x; mov.u16 %rh2, %ctaid.y; mul.wide.u16 %r1, %rh2, %rh1; cvt.u32.u16 %r2, %ctaid.x; add.u32 %r3, %r2, %r1; ld.global.s32 %r4, [restrictRules+4]; mov.s32 %r5, -1; set.ne.u32.s32 %r6, %r4, %r5; neg.s32 %r7, %r6; ld.global.u32 %r8, [restrictRules+4]; set.ne.u32.u32 %r9, %r8, %r3; neg.s32 %r10, %r9; and.b32 %r11, %r7, %r10; mov.u32 %r12, 0; setp.eq.s32 %p2, %r11, %r12; @%p2 bra $Lt_0_19714; .loc 3 168 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_214_9; $Lt_0_19714: mov.u16 %rh3, %ntid.x; mov.u16 %rh4, %tid.y; mul.wide.u16 %r13, %rh4, %rh3; mov.u16 %rh5, %ntid.y; mul.wide.u16 %r14, %rh5, %rh3; cvt.u32.u16 %r15, %tid.z; mul.lo.u32 %r16, %r15, %r14; add.u32 %r17, %r16, %r13; cvt.u32.u16 %r18, %tid.x; add.u32 %r19, %r18, %r17; ld.global.s32 %r20, [restrictRules+0]; mov.s32 %r21, -1; set.ne.u32.s32 %r22, %r20, %r21; neg.s32 %r23, %r22; ld.global.u32 %r24, [restrictRules+0]; set.ne.u32.u32 %r25, %r24, %r19; neg.s32 %r26, %r25; and.b32 %r27, %r23, %r26; mov.u32 %r28, 0; setp.eq.s32 %p3, %r27, %r28; @%p3 bra $Lt_0_20226; .loc 3 170 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_214_9; $Lt_0_20226: mov.u16 %rh6, %nctaid.y; mul.wide.u16 %r29, %rh6, %rh1; cvt.u32.u16 %r30, %ntid.z; mul.lo.u32 %r31, %r30, %r14; mul.lo.u32 %r32, %r29, %r31; ld.const.s32 %r33, [printfBufferLength]; div.s32 %r34, %r33, %r32; and.b32 %r35, %r34, -256; mov.u32 %r36, 511; setp.gt.u32 %p4, %r35, %r36; @%p4 bra $Lt_0_20738; .loc 3 188 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_214_9; $Lt_0_20738: .loc 3 192 0 mul.lo.u32 %r37, %r31, %r3; add.u32 %r38, %r37, %r19; ld.const.u64 %rd4, [globalPrintfBuffer]; mul.lo.u32 %r39, %r38, %r35; cvt.u64.u32 %rd5, %r39; add.u64 %rd6, %rd5, %rd4; ld.global.v2.u16 {%r40,%r41}, [%rd6+0]; ld.global.v2.u32 {%r42,%r43}, [%rd6+8]; mov.u32 %r44, 51216; setp.eq.u32 %p5, %r40, %r44; @%p5 bra $Lt_0_21250; .loc 3 198 0 mov.s32 %r42, %r35; .loc 3 200 0 mov.u32 %r45, 51216; st.global.v2.u16 [%rd6+0], {%r45,%r41}; st.global.u32 [%rd6+4], %r38; mov.u32 %r46, 0; st.global.v2.u32 [%rd6+8], {%r35,%r46}; .loc 3 205 0 st.global.u32 [%rd4+8], %r35; mov.u32 %r43, 0; $Lt_0_21250: .loc 3 214 0 add.u32 %r47, %r43, 256; setp.ge.u32 %p6, %r47, %r42; mov.u32 %r48, 256; selp.u32 %r49, %r48, %r47, %p6; st.global.u32 [%rd6+12], %r49; .loc 3 215 0 cvt.u64.u32 %rd7, %r49; add.u64 %rd3, %rd7, %rd6; $LDWendi__Z19getNextPrintfBufPtrv_214_9: .loc 3 401 0 mov.u64 %rd8, 0; setp.ne.u64 %p7, %rd3, %rd8; @%p7 bra $Lt_0_21762; bra.uni $LDWendi__Z17writePrintfHeaderPcS__214_1; $Lt_0_21762: .loc 3 403 0 cvt.u32.u16 %r50, %tid.x; mov.u64 %rd9, -8; setp.ne.u64 %p8, %rd3, %rd9; @%p8 bra $Lt_0_22274; .loc 3 328 0 mov.u64 %rd10, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__214_7; $Lt_0_22274: .loc 3 331 0 mov.s32 %r51, 4; st.global.s32 [%rd3+8], %r51; .loc 3 333 0 st.global.u32 [%rd3+16], %r50; .loc 3 335 0 mov.s16 %rh7, 0; st.global.s8 [%rd3+24], %rh7; .loc 3 337 0 add.u64 %rd10, %rd3, 24; $LDWendi__Z7copyArgIjEPcS0_RT_S0__214_7: .loc 3 404 0 add.u64 %rd11, %rd3, 256; cvt.u32.u16 %r52, %ctaid.x; mov.u64 %rd12, 0; set.eq.u32.u64 %r53, %rd10, %rd12; neg.s32 %r54, %r53; add.u64 %rd13, %rd10, 8; set.le.u32.u64 %r55, %rd11, %rd13; neg.s32 %r56, %r55; or.b32 %r57, %r54, %r56; mov.u32 %r58, 0; setp.eq.s32 %p9, %r57, %r58; @%p9 bra $Lt_0_22786; .loc 3 328 0 mov.u64 %rd14, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__214_5; $Lt_0_22786: .loc 3 331 0 mov.s32 %r59, 4; st.global.s32 [%rd10+0], %r59; .loc 3 333 0 st.global.u32 [%rd10+8], %r52; .loc 3 335 0 mov.s16 %rh8, 0; st.global.s8 [%rd10+16], %rh8; .loc 3 337 0 add.u64 %rd14, %rd10, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__214_5: .loc 3 406 0 mov.u64 %rd15, 0; set.eq.u32.u64 %r60, %rd14, %rd15; neg.s32 %r61, %r60; set.ge.u32.u64 %r62, %rd14, %rd11; neg.s32 %r63, %r62; or.b32 %r64, %r61, %r63; mov.u32 %r65, 0; setp.eq.s32 %p10, %r64, %r65; @%p10 bra $Lt_0_24066; .loc 3 261 0 mov.u64 %rd16, 0; bra.uni $LDWendi__Z15cuPrintfStrncpyPcPKciS__214_2; $Lt_0_24066: add.u64 %rd17, %rd14, 8; mov.s32 %r66, 0; mov.u64 %rd18, __constant461; $Lt_0_24322: // Loop body line 261, nesting depth: 1, iterations: 256 mov.s64 %rd19, %rd17; .loc 3 274 0 setp.le.u64 %p11, %rd11, %rd17; @%p11 bra $Lt_0_24834; .loc 3 276 0 add.s32 %r66, %r66, 1; .loc 3 277 0 cvt.u64.s32 %rd20, %r66; add.u64 %rd21, %rd14, %rd20; add.u64 %rd22, %rd21, 8; mov.s64 %rd17, %rd22; mov.s64 %rd19, %rd22; add.u64 %rd23, %rd20, %rd18; ld.const.s8 %rh9, [%rd23+-1]; st.global.s8 [%rd22+-1], %rh9; .loc 3 279 0 ld.const.s8 %r67, [%rd23+-1]; mov.u32 %r68, 0; setp.eq.s32 %p12, %r67, %r68; @%p12 bra $Lt_0_24834; mov.u32 %r69, 256; setp.ne.s32 %p13, %r66, %r69; @%p13 bra $Lt_0_24322; add.u64 %rd19, %rd14, 264; $Lt_0_24834: $Lt_0_770: .loc 3 283 0 setp.gt.u64 %p14, %rd11, %rd19; @!%p14 bra $Lt_0_25346; and.b64 %rd24, %rd19, 7; mov.u64 %rd25, 0; setp.eq.u64 %p15, %rd24, %rd25; @%p15 bra $Lt_0_25346; $L_0_18178: .loc 3 285 0 add.s32 %r66, %r66, 1; .loc 3 286 0 add.u64 %rd19, %rd19, 1; mov.s16 %rh10, 0; st.global.s8 [%rd19+-1], %rh10; .loc 3 283 0 setp.gt.u64 %p14, %rd11, %rd19; @!%p14 bra $Lt_0_25346; and.b64 %rd26, %rd19, 7; mov.u64 %rd27, 0; setp.ne.u64 %p16, %rd26, %rd27; @%p16 bra $L_0_18178; $Lt_0_25346: $L_0_18434: .loc 3 288 0 st.global.s32 [%rd14+0], %r66; .loc 3 289 0 mov.u64 %rd28, 0; selp.u64 %rd16, %rd19, %rd28, %p14; $LDWendi__Z15cuPrintfStrncpyPcPKciS__214_2: .loc 3 241 0 mov.u32 %r70, 51217; mov.s64 %rd29, 0; mov.u64 %rd30, 0; setp.ne.u64 %p17, %rd16, %rd30; selp.s64 %rd31, %rd14, %rd29, %p17; cvt.s32.s64 %r71, %rd31; cvt.s32.s64 %r72, %rd3; sub.s32 %r73, %r71, %r72; mov.u16 %rh11, %nctaid.x; mov.u16 %rh12, %ctaid.y; mul.wide.u16 %r74, %rh11, %rh12; add.u32 %r75, %r52, %r74; mov.u16 %rh13, %ntid.x; mov.u16 %rh14, %tid.y; mul.wide.u16 %r76, %rh14, %rh13; cvt.u32.u16 %r77, %tid.z; mov.u16 %rh15, %ntid.y; mul.wide.u16 %r78, %rh15, %rh13; mul.lo.u32 %r79, %r77, %r78; add.u32 %r80, %r76, %r79; add.u32 %r81, %r50, %r80; st.global.v4.u16 [%rd3+0], {%r70,%r73,%r75,%r81}; $LDWendi__Z17writePrintfHeaderPcS__214_1: .loc 17 15 0 exit; $LDWend__Z23printTest_1_dimensionalv: } // _Z23printTest_1_dimensionalv .const .align 1 .b8 __constant463[48] = {0x74,0x68,0x72,0x65,0x61,0x64,0x49,0x64,0x78,0x20,0x5b,0x78,0x3a,0x25,0x64,0x2c,0x20,0x79,0x3a,0x25,0x64,0x5d,0x20,0x62,0x6c,0x6f,0x63,0x6b,0x49,0x64,0x78,0x20,0x20,0x5b,0x78,0x3a,0x25,0x64,0x2c,0x20,0x79,0x3a,0x25,0x64,0x5d,0xa,0x20,0x0}; .entry _Z23printTest_2_dimensionalv { .reg .u16 %rh<14>; .reg .u32 %r<102>; .reg .u64 %rd<39>; .reg .pred %p<21>; .loc 17 19 0 $LDWbegin__Z23printTest_2_dimensionalv: .loc 17 21 0 ld.global.u64 %rd1, [printfBufferPtr]; mov.u64 %rd2, 0; setp.ne.u64 %p1, %rd1, %rd2; @%p1 bra $Lt_1_22274; .loc 3 164 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_215_13; $Lt_1_22274: mov.u16 %rh1, %nctaid.x; mov.u16 %rh2, %ctaid.y; mul.wide.u16 %r1, %rh2, %rh1; cvt.u32.u16 %r2, %ctaid.x; add.u32 %r3, %r2, %r1; ld.global.s32 %r4, [restrictRules+4]; mov.s32 %r5, -1; set.ne.u32.s32 %r6, %r4, %r5; neg.s32 %r7, %r6; ld.global.u32 %r8, [restrictRules+4]; set.ne.u32.u32 %r9, %r8, %r3; neg.s32 %r10, %r9; and.b32 %r11, %r7, %r10; mov.u32 %r12, 0; setp.eq.s32 %p2, %r11, %r12; @%p2 bra $Lt_1_22786; .loc 3 168 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_215_13; $Lt_1_22786: mov.u16 %rh3, %ntid.x; mov.u16 %rh4, %tid.y; mul.wide.u16 %r13, %rh4, %rh3; mov.u16 %rh5, %ntid.y; mul.wide.u16 %r14, %rh5, %rh3; cvt.u32.u16 %r15, %tid.z; mul.lo.u32 %r16, %r15, %r14; add.u32 %r17, %r16, %r13; cvt.u32.u16 %r18, %tid.x; add.u32 %r19, %r18, %r17; ld.global.s32 %r20, [restrictRules+0]; mov.s32 %r21, -1; set.ne.u32.s32 %r22, %r20, %r21; neg.s32 %r23, %r22; ld.global.u32 %r24, [restrictRules+0]; set.ne.u32.u32 %r25, %r24, %r19; neg.s32 %r26, %r25; and.b32 %r27, %r23, %r26; mov.u32 %r28, 0; setp.eq.s32 %p3, %r27, %r28; @%p3 bra $Lt_1_23298; .loc 3 170 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_215_13; $Lt_1_23298: mov.u16 %rh6, %nctaid.y; mul.wide.u16 %r29, %rh6, %rh1; cvt.u32.u16 %r30, %ntid.z; mul.lo.u32 %r31, %r30, %r14; mul.lo.u32 %r32, %r29, %r31; ld.const.s32 %r33, [printfBufferLength]; div.s32 %r34, %r33, %r32; and.b32 %r35, %r34, -256; mov.u32 %r36, 511; setp.gt.u32 %p4, %r35, %r36; @%p4 bra $Lt_1_23810; .loc 3 188 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_215_13; $Lt_1_23810: .loc 3 192 0 mul.lo.u32 %r37, %r31, %r3; add.u32 %r38, %r37, %r19; ld.const.u64 %rd4, [globalPrintfBuffer]; mul.lo.u32 %r39, %r38, %r35; cvt.u64.u32 %rd5, %r39; add.u64 %rd6, %rd5, %rd4; ld.global.v2.u16 {%r40,%r41}, [%rd6+0]; ld.global.v2.u32 {%r42,%r43}, [%rd6+8]; mov.u32 %r44, 51216; setp.eq.u32 %p5, %r40, %r44; @%p5 bra $Lt_1_24322; .loc 3 198 0 mov.s32 %r42, %r35; .loc 3 200 0 mov.u32 %r45, 51216; st.global.v2.u16 [%rd6+0], {%r45,%r41}; st.global.u32 [%rd6+4], %r38; mov.u32 %r46, 0; st.global.v2.u32 [%rd6+8], {%r35,%r46}; .loc 3 205 0 st.global.u32 [%rd4+8], %r35; mov.u32 %r43, 0; $Lt_1_24322: .loc 3 214 0 add.u32 %r47, %r43, 256; setp.ge.u32 %p6, %r47, %r42; mov.u32 %r48, 256; selp.u32 %r49, %r48, %r47, %p6; st.global.u32 [%rd6+12], %r49; .loc 3 215 0 cvt.u64.u32 %rd7, %r49; add.u64 %rd3, %rd7, %rd6; $LDWendi__Z19getNextPrintfBufPtrv_215_13: .loc 3 420 0 mov.u64 %rd8, 0; setp.ne.u64 %p7, %rd3, %rd8; @%p7 bra $Lt_1_24834; bra.uni $LDWendi__Z17writePrintfHeaderPcS__215_1; $Lt_1_24834: .loc 3 422 0 cvt.u32.u16 %r50, %tid.x; mov.u64 %rd9, -8; setp.ne.u64 %p8, %rd3, %rd9; @%p8 bra $Lt_1_25346; .loc 3 328 0 mov.u64 %rd10, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_11; $Lt_1_25346: .loc 3 331 0 mov.s32 %r51, 4; st.global.s32 [%rd3+8], %r51; .loc 3 333 0 st.global.u32 [%rd3+16], %r50; .loc 3 335 0 mov.s16 %rh7, 0; st.global.s8 [%rd3+24], %rh7; .loc 3 337 0 add.u64 %rd10, %rd3, 24; $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_11: .loc 3 423 0 add.u64 %rd11, %rd3, 256; cvt.u32.u16 %r52, %tid.y; mov.u64 %rd12, 0; set.eq.u32.u64 %r53, %rd10, %rd12; neg.s32 %r54, %r53; add.u64 %rd13, %rd10, 8; set.le.u32.u64 %r55, %rd11, %rd13; neg.s32 %r56, %r55; or.b32 %r57, %r54, %r56; mov.u32 %r58, 0; setp.eq.s32 %p9, %r57, %r58; @%p9 bra $Lt_1_25858; .loc 3 328 0 mov.u64 %rd14, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_9; $Lt_1_25858: .loc 3 331 0 mov.s32 %r59, 4; st.global.s32 [%rd10+0], %r59; .loc 3 333 0 st.global.u32 [%rd10+8], %r52; .loc 3 335 0 mov.s16 %rh8, 0; st.global.s8 [%rd10+16], %rh8; .loc 3 337 0 add.u64 %rd14, %rd10, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_9: .loc 3 424 0 cvt.u32.u16 %r60, %ctaid.x; mov.u64 %rd15, 0; set.eq.u32.u64 %r61, %rd14, %rd15; neg.s32 %r62, %r61; add.u64 %rd16, %rd14, 8; set.le.u32.u64 %r63, %rd11, %rd16; neg.s32 %r64, %r63; or.b32 %r65, %r62, %r64; mov.u32 %r66, 0; setp.eq.s32 %p10, %r65, %r66; @%p10 bra $Lt_1_26370; .loc 3 328 0 mov.u64 %rd17, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_7; $Lt_1_26370: .loc 3 331 0 mov.s32 %r67, 4; st.global.s32 [%rd14+0], %r67; .loc 3 333 0 st.global.u32 [%rd14+8], %r60; .loc 3 335 0 mov.s16 %rh9, 0; st.global.s8 [%rd14+16], %rh9; .loc 3 337 0 add.u64 %rd17, %rd14, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_7: .loc 3 425 0 cvt.u32.u16 %r68, %ctaid.y; mov.u64 %rd18, 0; set.eq.u32.u64 %r69, %rd17, %rd18; neg.s32 %r70, %r69; add.u64 %rd19, %rd17, 8; set.le.u32.u64 %r71, %rd11, %rd19; neg.s32 %r72, %r71; or.b32 %r73, %r70, %r72; mov.u32 %r74, 0; setp.eq.s32 %p11, %r73, %r74; @%p11 bra $Lt_1_26882; .loc 3 328 0 mov.u64 %rd20, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_5; $Lt_1_26882: .loc 3 331 0 mov.s32 %r75, 4; st.global.s32 [%rd17+0], %r75; .loc 3 333 0 st.global.u32 [%rd17+8], %r68; .loc 3 335 0 mov.s16 %rh10, 0; st.global.s8 [%rd17+16], %rh10; .loc 3 337 0 add.u64 %rd20, %rd17, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__215_5: .loc 3 427 0 mov.u64 %rd21, 0; set.eq.u32.u64 %r76, %rd20, %rd21; neg.s32 %r77, %r76; set.ge.u32.u64 %r78, %rd20, %rd11; neg.s32 %r79, %r78; or.b32 %r80, %r77, %r79; mov.u32 %r81, 0; setp.eq.s32 %p12, %r80, %r81; @%p12 bra $Lt_1_28162; .loc 3 261 0 mov.u64 %rd22, 0; bra.uni $LDWendi__Z15cuPrintfStrncpyPcPKciS__215_2; $Lt_1_28162: add.u64 %rd23, %rd20, 8; mov.s32 %r82, 0; mov.u64 %rd24, __constant463; $Lt_1_28418: // Loop body line 261, nesting depth: 1, iterations: 256 mov.s64 %rd25, %rd23; .loc 3 274 0 setp.le.u64 %p13, %rd11, %rd23; @%p13 bra $Lt_1_28930; .loc 3 276 0 add.s32 %r82, %r82, 1; .loc 3 277 0 cvt.u64.s32 %rd26, %r82; add.u64 %rd27, %rd20, %rd26; add.u64 %rd28, %rd27, 8; mov.s64 %rd23, %rd28; mov.s64 %rd25, %rd28; add.u64 %rd29, %rd26, %rd24; ld.const.s8 %rh11, [%rd29+-1]; st.global.s8 [%rd28+-1], %rh11; .loc 3 279 0 ld.const.s8 %r83, [%rd29+-1]; mov.u32 %r84, 0; setp.eq.s32 %p14, %r83, %r84; @%p14 bra $Lt_1_28930; mov.u32 %r85, 256; setp.ne.s32 %p15, %r82, %r85; @%p15 bra $Lt_1_28418; add.u64 %rd25, %rd20, 264; $Lt_1_28930: $Lt_1_770: .loc 3 283 0 setp.gt.u64 %p16, %rd11, %rd25; @!%p16 bra $Lt_1_29442; and.b64 %rd30, %rd25, 7; mov.u64 %rd31, 0; setp.eq.u64 %p17, %rd30, %rd31; @%p17 bra $Lt_1_29442; $L_1_21250: .loc 3 285 0 add.s32 %r82, %r82, 1; .loc 3 286 0 add.u64 %rd25, %rd25, 1; mov.s16 %rh12, 0; st.global.s8 [%rd25+-1], %rh12; .loc 3 283 0 setp.gt.u64 %p16, %rd11, %rd25; @!%p16 bra $Lt_1_29442; and.b64 %rd32, %rd25, 7; mov.u64 %rd33, 0; setp.ne.u64 %p18, %rd32, %rd33; @%p18 bra $L_1_21250; $Lt_1_29442: $L_1_21506: .loc 3 288 0 st.global.s32 [%rd20+0], %r82; .loc 3 289 0 mov.u64 %rd34, 0; selp.u64 %rd22, %rd25, %rd34, %p16; $LDWendi__Z15cuPrintfStrncpyPcPKciS__215_2: .loc 3 241 0 mov.u32 %r86, 51217; mov.s64 %rd35, 0; mov.u64 %rd36, 0; setp.ne.u64 %p19, %rd22, %rd36; selp.s64 %rd37, %rd20, %rd35, %p19; cvt.s32.s64 %r87, %rd37; cvt.s32.s64 %r88, %rd3; sub.s32 %r89, %r87, %r88; cvt.u32.u16 %r90, %nctaid.x; mul.lo.u32 %r91, %r90, %r68; add.u32 %r92, %r60, %r91; cvt.u32.u16 %r93, %ntid.x; cvt.u32.u16 %r94, %tid.z; cvt.u32.u16 %r95, %ntid.y; mul.lo.u32 %r96, %r95, %r93; mul.lo.u32 %r97, %r94, %r96; mul.lo.u32 %r98, %r52, %r93; add.u32 %r99, %r97, %r98; add.u32 %r100, %r50, %r99; st.global.v4.u16 [%rd3+0], {%r86,%r89,%r92,%r100}; $LDWendi__Z17writePrintfHeaderPcS__215_1: .loc 17 23 0 exit; $LDWend__Z23printTest_2_dimensionalv: } // _Z23printTest_2_dimensionalv .const .align 1 .b8 __constant465[60] = {0x74,0x68,0x72,0x65,0x61,0x64,0x49,0x64,0x78,0x20,0x5b,0x78,0x3a,0x25,0x64,0x2c,0x20,0x79,0x3a,0x25,0x64,0x2c,0x20,0x7a,0x3a,0x25,0x64,0x5d,0x20,0x62,0x6c,0x6f,0x63,0x6b,0x49,0x64,0x78,0x20,0x20,0x5b,0x78,0x3a,0x25,0x64,0x2c,0x20,0x79,0x3a,0x25,0x64,0x2c,0x20,0x7a,0x3a,0x25,0x64,0x5d,0xa,0x20,0x0}; .entry _Z23printTest_3_dimensionalv { .reg .u16 %rh<16>; .reg .u32 %r<117>; .reg .u64 %rd<45>; .reg .pred %p<23>; .loc 17 27 0 $LDWbegin__Z23printTest_3_dimensionalv: .loc 17 29 0 ld.global.u64 %rd1, [printfBufferPtr]; mov.u64 %rd2, 0; setp.ne.u64 %p1, %rd1, %rd2; @%p1 bra $Lt_2_25346; .loc 3 164 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_216_17; $Lt_2_25346: mov.u16 %rh1, %nctaid.x; mov.u16 %rh2, %ctaid.y; mul.wide.u16 %r1, %rh2, %rh1; cvt.u32.u16 %r2, %ctaid.x; add.u32 %r3, %r2, %r1; ld.global.s32 %r4, [restrictRules+4]; mov.s32 %r5, -1; set.ne.u32.s32 %r6, %r4, %r5; neg.s32 %r7, %r6; ld.global.u32 %r8, [restrictRules+4]; set.ne.u32.u32 %r9, %r8, %r3; neg.s32 %r10, %r9; and.b32 %r11, %r7, %r10; mov.u32 %r12, 0; setp.eq.s32 %p2, %r11, %r12; @%p2 bra $Lt_2_25858; .loc 3 168 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_216_17; $Lt_2_25858: mov.u16 %rh3, %ntid.x; mov.u16 %rh4, %tid.y; mul.wide.u16 %r13, %rh4, %rh3; mov.u16 %rh5, %ntid.y; mul.wide.u16 %r14, %rh5, %rh3; cvt.u32.u16 %r15, %tid.z; mul.lo.u32 %r16, %r15, %r14; add.u32 %r17, %r16, %r13; cvt.u32.u16 %r18, %tid.x; add.u32 %r19, %r18, %r17; ld.global.s32 %r20, [restrictRules+0]; mov.s32 %r21, -1; set.ne.u32.s32 %r22, %r20, %r21; neg.s32 %r23, %r22; ld.global.u32 %r24, [restrictRules+0]; set.ne.u32.u32 %r25, %r24, %r19; neg.s32 %r26, %r25; and.b32 %r27, %r23, %r26; mov.u32 %r28, 0; setp.eq.s32 %p3, %r27, %r28; @%p3 bra $Lt_2_26370; .loc 3 170 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_216_17; $Lt_2_26370: mov.u16 %rh6, %nctaid.y; mul.wide.u16 %r29, %rh6, %rh1; cvt.u32.u16 %r30, %ntid.z; mul.lo.u32 %r31, %r30, %r14; mul.lo.u32 %r32, %r29, %r31; ld.const.s32 %r33, [printfBufferLength]; div.s32 %r34, %r33, %r32; and.b32 %r35, %r34, -256; mov.u32 %r36, 511; setp.gt.u32 %p4, %r35, %r36; @%p4 bra $Lt_2_26882; .loc 3 188 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_216_17; $Lt_2_26882: .loc 3 192 0 mul.lo.u32 %r37, %r31, %r3; add.u32 %r38, %r37, %r19; ld.const.u64 %rd4, [globalPrintfBuffer]; mul.lo.u32 %r39, %r38, %r35; cvt.u64.u32 %rd5, %r39; add.u64 %rd6, %rd5, %rd4; ld.global.v2.u16 {%r40,%r41}, [%rd6+0]; ld.global.v2.u32 {%r42,%r43}, [%rd6+8]; mov.u32 %r44, 51216; setp.eq.u32 %p5, %r40, %r44; @%p5 bra $Lt_2_27394; .loc 3 198 0 mov.s32 %r42, %r35; .loc 3 200 0 mov.u32 %r45, 51216; st.global.v2.u16 [%rd6+0], {%r45,%r41}; st.global.u32 [%rd6+4], %r38; mov.u32 %r46, 0; st.global.v2.u32 [%rd6+8], {%r35,%r46}; .loc 3 205 0 st.global.u32 [%rd4+8], %r35; mov.u32 %r43, 0; $Lt_2_27394: .loc 3 214 0 add.u32 %r47, %r43, 256; setp.ge.u32 %p6, %r47, %r42; mov.u32 %r48, 256; selp.u32 %r49, %r48, %r47, %p6; st.global.u32 [%rd6+12], %r49; .loc 3 215 0 cvt.u64.u32 %rd7, %r49; add.u64 %rd3, %rd7, %rd6; $LDWendi__Z19getNextPrintfBufPtrv_216_17: .loc 3 443 0 mov.u64 %rd8, 0; setp.ne.u64 %p7, %rd3, %rd8; @%p7 bra $Lt_2_27906; bra.uni $LDWendi__Z17writePrintfHeaderPcS__216_1; $Lt_2_27906: .loc 3 445 0 cvt.u32.u16 %r50, %tid.x; mov.u64 %rd9, -8; setp.ne.u64 %p8, %rd3, %rd9; @%p8 bra $Lt_2_28418; .loc 3 328 0 mov.u64 %rd10, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_15; $Lt_2_28418: .loc 3 331 0 mov.s32 %r51, 4; st.global.s32 [%rd3+8], %r51; .loc 3 333 0 st.global.u32 [%rd3+16], %r50; .loc 3 335 0 mov.s16 %rh7, 0; st.global.s8 [%rd3+24], %rh7; .loc 3 337 0 add.u64 %rd10, %rd3, 24; $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_15: .loc 3 446 0 add.u64 %rd11, %rd3, 256; cvt.u32.u16 %r52, %tid.y; mov.u64 %rd12, 0; set.eq.u32.u64 %r53, %rd10, %rd12; neg.s32 %r54, %r53; add.u64 %rd13, %rd10, 8; set.le.u32.u64 %r55, %rd11, %rd13; neg.s32 %r56, %r55; or.b32 %r57, %r54, %r56; mov.u32 %r58, 0; setp.eq.s32 %p9, %r57, %r58; @%p9 bra $Lt_2_28930; .loc 3 328 0 mov.u64 %rd14, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_13; $Lt_2_28930: .loc 3 331 0 mov.s32 %r59, 4; st.global.s32 [%rd10+0], %r59; .loc 3 333 0 st.global.u32 [%rd10+8], %r52; .loc 3 335 0 mov.s16 %rh8, 0; st.global.s8 [%rd10+16], %rh8; .loc 3 337 0 add.u64 %rd14, %rd10, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_13: .loc 3 447 0 cvt.u32.u16 %r60, %tid.z; mov.u64 %rd15, 0; set.eq.u32.u64 %r61, %rd14, %rd15; neg.s32 %r62, %r61; add.u64 %rd16, %rd14, 8; set.le.u32.u64 %r63, %rd11, %rd16; neg.s32 %r64, %r63; or.b32 %r65, %r62, %r64; mov.u32 %r66, 0; setp.eq.s32 %p10, %r65, %r66; @%p10 bra $Lt_2_29442; .loc 3 328 0 mov.u64 %rd17, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_11; $Lt_2_29442: .loc 3 331 0 mov.s32 %r67, 4; st.global.s32 [%rd14+0], %r67; .loc 3 333 0 st.global.u32 [%rd14+8], %r60; .loc 3 335 0 mov.s16 %rh9, 0; st.global.s8 [%rd14+16], %rh9; .loc 3 337 0 add.u64 %rd17, %rd14, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_11: .loc 3 448 0 cvt.u32.u16 %r68, %ctaid.x; mov.u64 %rd18, 0; set.eq.u32.u64 %r69, %rd17, %rd18; neg.s32 %r70, %r69; add.u64 %rd19, %rd17, 8; set.le.u32.u64 %r71, %rd11, %rd19; neg.s32 %r72, %r71; or.b32 %r73, %r70, %r72; mov.u32 %r74, 0; setp.eq.s32 %p11, %r73, %r74; @%p11 bra $Lt_2_29954; .loc 3 328 0 mov.u64 %rd20, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_9; $Lt_2_29954: .loc 3 331 0 mov.s32 %r75, 4; st.global.s32 [%rd17+0], %r75; .loc 3 333 0 st.global.u32 [%rd17+8], %r68; .loc 3 335 0 mov.s16 %rh10, 0; st.global.s8 [%rd17+16], %rh10; .loc 3 337 0 add.u64 %rd20, %rd17, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_9: .loc 3 449 0 cvt.u32.u16 %r76, %ctaid.y; mov.u64 %rd21, 0; set.eq.u32.u64 %r77, %rd20, %rd21; neg.s32 %r78, %r77; add.u64 %rd22, %rd20, 8; set.le.u32.u64 %r79, %rd11, %rd22; neg.s32 %r80, %r79; or.b32 %r81, %r78, %r80; mov.u32 %r82, 0; setp.eq.s32 %p12, %r81, %r82; @%p12 bra $Lt_2_30466; .loc 3 328 0 mov.u64 %rd23, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_7; $Lt_2_30466: .loc 3 331 0 mov.s32 %r83, 4; st.global.s32 [%rd20+0], %r83; .loc 3 333 0 st.global.u32 [%rd20+8], %r76; .loc 3 335 0 mov.s16 %rh11, 0; st.global.s8 [%rd20+16], %rh11; .loc 3 337 0 add.u64 %rd23, %rd20, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_7: .loc 3 450 0 mov.u64 %rd24, 0; set.eq.u32.u64 %r84, %rd23, %rd24; neg.s32 %r85, %r84; add.u64 %rd25, %rd23, 8; set.le.u32.u64 %r86, %rd11, %rd25; neg.s32 %r87, %r86; or.b32 %r88, %r85, %r87; mov.u32 %r89, 0; setp.eq.s32 %p13, %r88, %r89; @%p13 bra $Lt_2_30978; .loc 3 328 0 mov.u64 %rd26, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_5; $Lt_2_30978: .loc 3 331 0 mov.s32 %r90, 4; st.global.s32 [%rd23+0], %r90; .loc 3 333 0 cvt.u32.u16 %r91, %ctaid.z; st.global.u32 [%rd23+8], %r91; .loc 3 335 0 mov.s16 %rh12, 0; st.global.s8 [%rd23+16], %rh12; .loc 3 337 0 add.u64 %rd26, %rd23, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__216_5: .loc 3 451 0 mov.u64 %rd27, 0; set.eq.u32.u64 %r92, %rd26, %rd27; neg.s32 %r93, %r92; set.ge.u32.u64 %r94, %rd26, %rd11; neg.s32 %r95, %r94; or.b32 %r96, %r93, %r95; mov.u32 %r97, 0; setp.eq.s32 %p14, %r96, %r97; @%p14 bra $Lt_2_32258; .loc 3 261 0 mov.u64 %rd28, 0; bra.uni $LDWendi__Z15cuPrintfStrncpyPcPKciS__216_2; $Lt_2_32258: add.u64 %rd29, %rd26, 8; mov.s32 %r98, 0; mov.u64 %rd30, __constant465; $Lt_2_32514: // Loop body line 261, nesting depth: 1, iterations: 256 mov.s64 %rd31, %rd29; .loc 3 274 0 setp.le.u64 %p15, %rd11, %rd29; @%p15 bra $Lt_2_33026; .loc 3 276 0 add.s32 %r98, %r98, 1; .loc 3 277 0 cvt.u64.s32 %rd32, %r98; add.u64 %rd33, %rd26, %rd32; add.u64 %rd34, %rd33, 8; mov.s64 %rd29, %rd34; mov.s64 %rd31, %rd34; add.u64 %rd35, %rd32, %rd30; ld.const.s8 %rh13, [%rd35+-1]; st.global.s8 [%rd34+-1], %rh13; .loc 3 279 0 ld.const.s8 %r99, [%rd35+-1]; mov.u32 %r100, 0; setp.eq.s32 %p16, %r99, %r100; @%p16 bra $Lt_2_33026; mov.u32 %r101, 256; setp.ne.s32 %p17, %r98, %r101; @%p17 bra $Lt_2_32514; add.u64 %rd31, %rd26, 264; $Lt_2_33026: $Lt_2_770: .loc 3 283 0 setp.gt.u64 %p18, %rd11, %rd31; @!%p18 bra $Lt_2_33538; and.b64 %rd36, %rd31, 7; mov.u64 %rd37, 0; setp.eq.u64 %p19, %rd36, %rd37; @%p19 bra $Lt_2_33538; $L_2_24322: .loc 3 285 0 add.s32 %r98, %r98, 1; .loc 3 286 0 add.u64 %rd31, %rd31, 1; mov.s16 %rh14, 0; st.global.s8 [%rd31+-1], %rh14; .loc 3 283 0 setp.gt.u64 %p18, %rd11, %rd31; @!%p18 bra $Lt_2_33538; and.b64 %rd38, %rd31, 7; mov.u64 %rd39, 0; setp.ne.u64 %p20, %rd38, %rd39; @%p20 bra $L_2_24322; $Lt_2_33538: $L_2_24578: .loc 3 288 0 st.global.s32 [%rd26+0], %r98; .loc 3 289 0 mov.u64 %rd40, 0; selp.u64 %rd28, %rd31, %rd40, %p18; $LDWendi__Z15cuPrintfStrncpyPcPKciS__216_2: .loc 3 241 0 mov.u32 %r102, 51217; mov.s64 %rd41, 0; mov.u64 %rd42, 0; setp.ne.u64 %p21, %rd28, %rd42; selp.s64 %rd43, %rd26, %rd41, %p21; cvt.s32.s64 %r103, %rd43; cvt.s32.s64 %r104, %rd3; sub.s32 %r105, %r103, %r104; cvt.u32.u16 %r106, %nctaid.x; mul.lo.u32 %r107, %r106, %r76; add.u32 %r108, %r68, %r107; cvt.u32.u16 %r109, %ntid.x; cvt.u32.u16 %r110, %ntid.y; mul.lo.u32 %r111, %r110, %r109; mul.lo.u32 %r112, %r60, %r111; mul.lo.u32 %r113, %r52, %r109; add.u32 %r114, %r112, %r113; add.u32 %r115, %r50, %r114; st.global.v4.u16 [%rd3+0], {%r102,%r105,%r108,%r115}; $LDWendi__Z17writePrintfHeaderPcS__216_1: .loc 17 31 0 exit; $LDWend__Z23printTest_3_dimensionalv: } // _Z23printTest_3_dimensionalv XL, 8threadVsBlockScheduling.cuELF3*@  @8@ H  4  <4 60 g)4 Y$* 0** * *+*.shstrtab.strtab.symtab.nv.global.init.nv.global.text._Z23printTest_3_dimensionalv.nv.info._Z23printTest_3_dimensionalv.text._Z23printTest_2_dimensionalv.nv.info._Z23printTest_2_dimensionalv.text._Z23printTest_1_dimensionalv.nv.info._Z23printTest_1_dimensionalv.nv.constant1._Z23printTest_1_dimensionalv.nv.constant1._Z23printTest_2_dimensionalv.nv.constant1._Z23printTest_3_dimensionalv.nv.constant14.rel.nv.constant14.nv.constant0_Z23printTest_3_dimensionalv_Z23printTest_2_dimensionalv_Z23printTest_1_dimensionalvrestrictRulesprintfBufferPtrprintfBufferLength__constant465__constant463globalPrintfBuffer__constant461 0       ;Xfv  `< 00  $ǀ' Ѐ|0ȇd|0Ȃd)-P! Ѐ ЀL H#0GAl N`!0Bd`)-PB# 0 D@ @0Ѐ@ `"A ǀ'ЀЀ  0GAl0Bd@)-@F@ 0H# F`G!J@ @`G 0`G$ Ѐ@0Ad)-@ @`G 0 `G  @ ` 0 `! Ѐ  Ѐ0؇@d@1G  Ѐ  Ѐ    Ѐ !Ѐ   0ȇd @$ ) Ѐ-@0|0ȇd|0Ȃd00ȇ@d0Ȃ@d1  Ѐ 1Ѐ  !AЀ @0 !GA !B@0 @00d 0d!Ѐ |0d |0d0Gd Ѐ ЀGGp -L  p Ѐ -Ѐ !GBЀ  @00d 0d!Ѐ |0d|0d0dЀ ЀGG@P   @ !B @5@00ȇd0d0HAd $P5  G0 !B@0G!?  @0ȇ Ѐ $P0GAl ! !B%@0P 0d0dЀ 0dGG   @  @!!C? @0 Ѐ 0d0dЀ0dGG 5 %P@ P 5 Ѐ  B#  0D@ @|0ȇd0|0Ȃd  `Ѐ@ N#B` 0H`"  G  Ѐ0 G GD!G%GD%>  G)@)`)0)`)@ )GD% %G G% @% `G%0% `G!0 0d Ѐ0 0G, Ѐ |0Gl 0 # ǀ' Ѐ|0ȇd|0Ȃd15P! Ѐ ЀL H#0GAl N`!0Bd`15PB# 0 D@ @0Ѐ@ `"A ǀ'ЀЀ  0GAl0Bd@15@F@ 0H# F`G!J@ @`G 0`G$ Ѐ@0Ad15@ @`G 0 `G  @ ` 0 `! Ѐ  Ѐ0؇@d@1G  Ѐ  Ѐ    Ѐ !Ѐ   0ȇd @$ 1 Ѐ5@0|0ȇd|0Ȃd00ȇ@d0Ȃ@d  Ѐ Ѐ  !AЀ @0 !GA !B@0 @0!0d 0d)Ѐ |0d! |0d%0GdЀ! ЀGЀ@ G    Ѐ Ѐ! !GBЀ  @0!B! @0%0d0d) Ѐ|0d% |0d!0Gd Ѐ! ЀGG 5L   Ѐ 5Ѐ! !GBЀ  @0%!B!|0d) |0d @0! Ѐ) 0d%0d% Ѐ0Gd ЀGG N 0 !% !!Ѐ! Ѐ) I!!GB%Ѐ % @0|0d|0dЀ0d)0d) Ѐ0d ЀG G   !B`@9@00ȇd0d0HAd(,P9  )` !B ))? Ѐ 0@0G @0(,P!0GAl )!B-@0P 0d 0dЀ 0dG G   @  @)!C!?  Ѐ @0% 0d0d% Ѐ0d G G 9 -P@P   9Ѐ B % 0D@ |0ȇd  @ @|0Ȃd 0%H@  0  `%0  `G H`G" 0  G     Ѐ0 G GD!G%GD%>  G)@)`)0)`)@ )GD% %G G% @% `G%0% `G!0 0d Ѐ0 0G, Ѐ |0Gl 0 #  ǀ'Ѐ|0ȇd|0Ȃd15`!Ѐ ЀL H#0GAlN`!0Bd`15`B# 0 D@ @0 Ѐ@ `"A ǀ'ЀЀ  0GAl0BdP15PF@ 0H#F`!J@ @ ` 0 `$P Ѐ@0Ad15P@`0` G@`0`! Ѐ Ѐ0؇@d@1G  Ѐ Ѐ  Ѐ! Ѐ 0ȇd @$ 1 Ѐ5@0|0ȇd|0Ȃd00ȇ@d0Ȃ@d Ѐ Ѐ !AЀ @0!GA !B@0 @0!0d 0d)Ѐ |0d! |0d%0GdЀ! ЀG Ѐ@ G  !%   Ѐ  Ѐ% !!GB Ѐ % @0 !B@00d 0d)Ѐ |0d|0d0Gd Ѐ ЀG 0G0    0   Ѐ  Ѐ  !GB Ѐ @0!B!@0%0d0d) Ѐ|0d%|0d!0Gd Ѐ! ЀG G 5L @    Ѐ 5Ѐ! ! !GBЀ @0%!B!|0d)|0d@0! Ѐ) 0d%0d% Ѐ0Gd ЀGGpN    p!!Ѐ! Ѐ% A !GB!Ѐ @0!|0d)|0d%!B! Ѐ)@0- 0d% 0d- Ѐ% 0Gd% ЀG Gp!%!!Ѐ% !!Ѐ) I!!GB%Ѐ %@0 |0d |0d Ѐ 0d)0d) Ѐ 0d  ЀGG !B@9 @00ȇd0d0HAd(,P9    !B   )?  Ѐ 0 @0 @0(,P!0GAl0)!B-@0P  0d 0d Ѐ  0d GG pp@pp@)!C ?  Ѐ @0% 0d 0d% Ѐ 0d  GG 9 -P@`P9Ѐ B  D@ %@|0ȇd`G% @|0Ȃd0)H @ %0`G 0  `G H`  0  G     Ѐ0G GD!G%GD%>  G)@)`)0)`)@ )GD% %G G% @% `G%0% `G!0 0dЀ0 0G, Ѐ|0Gl 0 # threadIdx blockIdx threadIdx [x:%d, y:%d] blockIdx [x:%d, y:%d] threadIdx [x:%d, y:%d, z:%d] blockIdx [x:%d, y:%d, z:%d] *PP`   ` `p p 0**_Z23printTest_3_dimensionalv_Z23printTest_2_dimensionalv_Z23printTest_1_dimensionalvprintfBufferLengthrestrictRules;O<EO\^O|OOOOP