ELF> @@ @8 @ @@@@@88@8@@@}} }}`}`  }}`}`TT@T@DDPtdLyLy@Ly@QtdRtd}}`}`88/lib64/ld-linux-x86-64.so.2GNUGNƯeEe'uD|{16tG      xIk)  Q>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```````HBH5bu %du @%bu h%Zu h%Ru h%Ju h%Bu h%:u h%2u h%*u hp%"u h`%u h P%u h @% u h 0%u h %t h %t h%t h%t h%t h%t h%t h%t h%t h1I^HHPTI@H@H@HHs HtHÐUHSH=t uK}`Ht H}`HHH9s$fDHHmt }`H_t H9rKt H[fff.UH='q HtHt}`ÐUHSH8H}HuHUHMDE̋uHMHUH]HEAHHH8[UHHEEvUHH H}HuEEUUHH0H}HuUHMEE-UHH0H}HuHUEEUHH H}EEUHH }uEEUHH0H}HuHEHEؾ%HVHEHEHr HEHHHE%HEHEHEHE@H+HEHEu HEHEHEЋE}~(Hr H&@kHEHEEHEHEEHEE%SHH@HEЋHr HMHHǸ}u+HEZHq HUHHǸ|HEHq HUHHǸzWHq HUHMHHǸY6Hq ?@HǸ>Heq HUHHǸ!HEHEUHEHEHEؾ%HyHEH}Hq HEHH/UHSHXHHLLdH%(HE1DžEQHH;uHHHHHHHHHf= H@f=t5H@H@Ho @HǸH@fu Ho H@>7H@HHHHHHbuUtHHLHHH;HUdH3%(tuHX[UHH H}H}HCEHEHE%HtHEHHEHEn n Hcи(`HHXt n HcHn HhHn Hrn EEEHEAHƿ`Hm A(`HA(`Ȁ`A0``UHHm Ht6Hm Ht-Hm H(Hm Hm Hm UHSHhH}EH}u H?m HEHZm EEHSm HtHOm Ht H3m Hu H-m HEHHEf=oEEEHl HE)H]HEHHJEtEEEf=t8}u$Hl H @2EHHEE}tREHH}EHHHMEHHUHEH]EIIHHމEHHHMHEHEHHUHEH]EIIHHމEHEHk k HHH;E@Ef=2HEHEAȀ`H}H Yk HJk HHHk H)‹Gk HHEHH?H}HHHEH k HEHHk H)‹ k HHEHH?H}HHHEȀ}tHHj j HH<Hj j HH Hj ]EHuIIމHj j HH Hj ]EH}HuIIމZHEHOj }tSj HcHAj HHh[UHSH(HEHHEйHHUMH]ЋEAAH߉uYHsi HH([UHHi HzUH@Hi @HUHUHSH(Ѐ`kHDi @R@H.i HD$HD$HD$H$AAy@y@HHz`Hh D$$AA@@HH`Hh D$$AA&y@&y@HH`Hfh D$$AA9y@9y@HHȀ`H*h D$$AA@@HHjH([ÐUHH}uUMHEUHEUPHEUPUHSH8H}HuHUHMDEHEuHMHUH]AHH4H8[UHSH8H}HuHUHMDEHEuHMHUH]AHHH8[UHSH8H}HuHUHMDEHEuHMHUH]AHHH8[UHSH8H}HuHUHMDEH]uHMHUHEAHH5H8[UHHH}HEHÐfffff.Hl$Ld$H-b L%b Ll$Lt$L|$H\$H8L)AIHIkHt1@LLDAHH9rH\$Hl$Ld$Ll$ Lt$(L|$0H8ÐHd Ht H11UHSHHPb Ht}`DHHHuH[ÐHHglobalPrintfBufferprintfBufferPtr%cdiouxXeEfgGaAnpsCorrupt printf buffer data - aborting %%@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@p@@@@@@[%d, %d]: printf buffer overflow No printf headers found at all! PUYX@8 @cudaPrintfExample.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_000074e2_00000000-7_cudaPrintfExample.cpp3.i (/tmp/ccBI#.YjgIua) //----------------------------------------------------------- //----------------------------------------------------------- // 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_000074e2_00000000-6_cudaPrintfExample.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 "cudaPrintfExample.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 __constant459[32] = {0x74,0x68,0x72,0x65,0x61,0x64,0x49,0x64,0x78,0x2e,0x78,0x20,0x25,0x64,0x20,0x74,0x68,0x72,0x65,0x61,0x64,0x49,0x64,0x78,0x2e,0x79,0x20,0x25,0x64,0x20,0x20,0x0}; .const .align 1 .b8 __constant460[29] = {0x62,0x6c,0x6f,0x63,0x6b,0x49,0x64,0x78,0x2e,0x78,0x20,0x25,0x64,0x20,0x62,0x6c,0x6f,0x63,0x6b,0x49,0x64,0x78,0x2e,0x79,0x20,0x25,0x64,0xa,0x0}; .entry _Z9printTestv { .reg .u16 %rh<17>; .reg .u32 %r<165>; .reg .u64 %rd<53>; .reg .pred %p<34>; .loc 17 12 0 $LDWbegin__Z9printTestv: .loc 17 14 0 ld.global.u64 %rd1, [printfBufferPtr]; mov.u64 %rd2, 0; setp.eq.u64 %p1, %rd1, %rd2; @!%p1 bra $Lt_0_38146; .loc 3 164 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_20; $Lt_0_38146: 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_38658; .loc 3 168 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_20; $Lt_0_38658: cvt.u32.u16 %r13, %ntid.x; cvt.u32.u16 %r14, %tid.y; mul.lo.u32 %r15, %r14, %r13; cvt.u32.u16 %r16, %ntid.y; mul.lo.u32 %r17, %r16, %r13; cvt.u32.u16 %r18, %tid.z; mul.lo.u32 %r19, %r18, %r17; add.u32 %r20, %r19, %r15; cvt.u32.u16 %r21, %tid.x; add.u32 %r22, %r21, %r20; ld.global.s32 %r23, [restrictRules+0]; mov.s32 %r24, -1; set.ne.u32.s32 %r25, %r23, %r24; neg.s32 %r26, %r25; ld.global.u32 %r27, [restrictRules+0]; set.ne.u32.u32 %r28, %r27, %r22; neg.s32 %r29, %r28; and.b32 %r30, %r26, %r29; mov.u32 %r31, 0; setp.eq.s32 %p3, %r30, %r31; @%p3 bra $Lt_0_39170; .loc 3 170 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_20; $Lt_0_39170: mov.u16 %rh3, %nctaid.y; mul.wide.u16 %r32, %rh3, %rh1; cvt.u32.u16 %r33, %ntid.z; mul.lo.u32 %r34, %r33, %r17; mul.lo.u32 %r35, %r34, %r32; ld.const.s32 %r36, [printfBufferLength]; div.s32 %r37, %r36, %r35; and.b32 %r38, %r37, -256; mov.u32 %r39, 511; setp.gt.u32 %p4, %r38, %r39; @%p4 bra $Lt_0_39682; .loc 3 188 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_20; $Lt_0_39682: .loc 3 192 0 mul.lo.u32 %r40, %r34, %r3; add.u32 %r41, %r40, %r22; mul.lo.u32 %r42, %r41, %r38; cvt.u64.u32 %rd4, %r42; ld.const.u64 %rd5, [globalPrintfBuffer]; add.u64 %rd6, %rd4, %rd5; ld.global.v2.u16 {%r43,%r44}, [%rd6+0]; ld.global.v2.u32 {%r45,%r46}, [%rd6+8]; mov.u32 %r47, 51216; setp.eq.u32 %p5, %r43, %r47; @%p5 bra $Lt_0_40194; .loc 3 198 0 mov.s32 %r45, %r38; .loc 3 200 0 mov.u32 %r48, 51216; st.global.v2.u16 [%rd6+0], {%r48,%r44}; st.global.u32 [%rd6+4], %r41; mov.u32 %r49, 0; st.global.v2.u32 [%rd6+8], {%r38,%r49}; .loc 3 205 0 st.global.u32 [%rd5+8], %r38; mov.u32 %r46, 0; $Lt_0_40194: .loc 3 214 0 add.u32 %r50, %r46, 256; setp.ge.u32 %p6, %r50, %r45; mov.u32 %r51, 256; selp.u32 %r52, %r51, %r50, %p6; st.global.u32 [%rd6+12], %r52; .loc 3 215 0 cvt.u64.u32 %rd7, %r52; add.u64 %rd3, %rd7, %rd6; $LDWendi__Z19getNextPrintfBufPtrv_212_20: .loc 3 401 0 mov.u64 %rd8, 0; setp.ne.u64 %p7, %rd3, %rd8; @%p7 bra $Lt_0_40706; bra.uni $LDWendi__Z17writePrintfHeaderPcS__212_12; $Lt_0_40706: .loc 3 403 0 cvt.u32.u16 %r53, %tid.x; mov.u64 %rd9, -8; setp.ne.u64 %p8, %rd3, %rd9; @%p8 bra $Lt_0_41218; .loc 3 328 0 mov.u64 %rd10, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_18; $Lt_0_41218: .loc 3 331 0 mov.s32 %r54, 4; st.global.s32 [%rd3+8], %r54; .loc 3 333 0 st.global.u32 [%rd3+16], %r53; .loc 3 335 0 mov.s16 %rh4, 0; st.global.s8 [%rd3+24], %rh4; .loc 3 337 0 add.u64 %rd10, %rd3, 24; $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_18: .loc 3 404 0 add.u64 %rd11, %rd3, 256; cvt.u32.u16 %r55, %tid.y; mov.u64 %rd12, 0; set.eq.u32.u64 %r56, %rd10, %rd12; neg.s32 %r57, %r56; add.u64 %rd13, %rd10, 8; set.le.u32.u64 %r58, %rd11, %rd13; neg.s32 %r59, %r58; or.b32 %r60, %r57, %r59; mov.u32 %r61, 0; setp.eq.s32 %p9, %r60, %r61; @%p9 bra $Lt_0_41730; .loc 3 328 0 mov.u64 %rd14, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_16; $Lt_0_41730: .loc 3 331 0 mov.s32 %r62, 4; st.global.s32 [%rd10+0], %r62; .loc 3 333 0 st.global.u32 [%rd10+8], %r55; .loc 3 335 0 mov.s16 %rh5, 0; st.global.s8 [%rd10+16], %rh5; .loc 3 337 0 add.u64 %rd14, %rd10, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_16: .loc 3 406 0 mov.u64 %rd15, 0; set.eq.u32.u64 %r63, %rd14, %rd15; neg.s32 %r64, %r63; set.ge.u32.u64 %r65, %rd14, %rd11; neg.s32 %r66, %r65; or.b32 %r67, %r64, %r66; mov.u32 %r68, 0; setp.eq.s32 %p10, %r67, %r68; @%p10 bra $Lt_0_43010; .loc 3 261 0 mov.u64 %rd16, 0; bra.uni $LDWendi__Z15cuPrintfStrncpyPcPKciS__212_13; $Lt_0_43010: add.u64 %rd17, %rd14, 8; mov.s32 %r69, 0; mov.u64 %rd18, __constant459; $Lt_0_43266: // 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_49410; .loc 3 276 0 add.s32 %r69, %r69, 1; .loc 3 277 0 cvt.u64.s32 %rd20, %r69; 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 %rh6, [%rd23+-1]; st.global.s8 [%rd22+-1], %rh6; .loc 3 279 0 ld.const.s8 %r70, [%rd23+-1]; mov.u32 %r71, 0; setp.eq.s32 %p12, %r70, %r71; @%p12 bra $Lt_0_49410; mov.u32 %r72, 256; setp.ne.s32 %p13, %r69, %r72; @%p13 bra $Lt_0_43266; add.u64 %rd19, %rd14, 264; $Lt_0_49410: $Lt_0_3586: .loc 3 283 0 setp.gt.u64 %p14, %rd11, %rd19; @!%p14 bra $Lt_0_49922; and.b64 %rd24, %rd19, 7; mov.u64 %rd25, 0; setp.eq.u64 %p15, %rd24, %rd25; @%p15 bra $Lt_0_49922; $L_0_36098: .loc 3 285 0 add.s32 %r69, %r69, 1; .loc 3 286 0 add.u64 %rd19, %rd19, 1; mov.s16 %rh7, 0; st.global.s8 [%rd19+-1], %rh7; .loc 3 283 0 setp.gt.u64 %p14, %rd11, %rd19; @!%p14 bra $Lt_0_49922; and.b64 %rd26, %rd19, 7; mov.u64 %rd27, 0; setp.ne.u64 %p16, %rd26, %rd27; @%p16 bra $L_0_36098; $Lt_0_49922: $L_0_36354: .loc 3 288 0 st.global.s32 [%rd14+0], %r69; .loc 3 289 0 mov.u64 %rd28, 0; selp.u64 %rd16, %rd19, %rd28, %p14; $LDWendi__Z15cuPrintfStrncpyPcPKciS__212_13: .loc 3 241 0 mov.u32 %r73, 51217; mov.s64 %rd29, 0; mov.u64 %rd30, 0; setp.ne.u64 %p17, %rd16, %rd30; selp.s64 %rd31, %rd14, %rd29, %p17; cvt.s32.s64 %r74, %rd31; cvt.s32.s64 %r75, %rd3; sub.s32 %r76, %r74, %r75; cvt.u32.u16 %r77, %ctaid.x; mov.u16 %rh8, %nctaid.x; mov.u16 %rh9, %ctaid.y; mul.wide.u16 %r78, %rh8, %rh9; add.u32 %r79, %r77, %r78; cvt.u32.u16 %r13, %ntid.x; cvt.u32.u16 %r80, %tid.z; cvt.u32.u16 %r81, %ntid.y; mul.lo.u32 %r82, %r81, %r13; mul.lo.u32 %r83, %r80, %r82; mul.lo.u32 %r84, %r55, %r13; add.u32 %r85, %r83, %r84; add.u32 %r86, %r53, %r85; st.global.v4.u16 [%rd3+0], {%r73,%r76,%r79,%r86}; $LDWendi__Z17writePrintfHeaderPcS__212_12: .loc 17 15 0 @!%p1 bra $Lt_0_43778; .loc 3 164 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_9; $Lt_0_43778: mov.u16 %rh1, %nctaid.x; mov.u16 %rh10, %ctaid.y; mul.wide.u16 %r87, %rh10, %rh1; cvt.u32.u16 %r88, %ctaid.x; add.u32 %r3, %r88, %r87; ld.global.s32 %r89, [restrictRules+4]; mov.s32 %r90, -1; set.ne.u32.s32 %r91, %r89, %r90; neg.s32 %r92, %r91; ld.global.u32 %r93, [restrictRules+4]; set.ne.u32.u32 %r94, %r93, %r3; neg.s32 %r95, %r94; and.b32 %r96, %r92, %r95; mov.u32 %r97, 0; setp.eq.s32 %p18, %r96, %r97; @%p18 bra $Lt_0_44290; .loc 3 168 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_9; $Lt_0_44290: cvt.u32.u16 %r13, %ntid.x; cvt.u32.u16 %r98, %tid.y; mul.lo.u32 %r99, %r98, %r13; cvt.u32.u16 %r100, %ntid.y; mul.lo.u32 %r17, %r100, %r13; cvt.u32.u16 %r101, %tid.z; mul.lo.u32 %r102, %r101, %r17; add.u32 %r103, %r102, %r99; cvt.u32.u16 %r104, %tid.x; add.u32 %r22, %r104, %r103; ld.global.s32 %r105, [restrictRules+0]; mov.s32 %r106, -1; set.ne.u32.s32 %r107, %r105, %r106; neg.s32 %r108, %r107; ld.global.u32 %r109, [restrictRules+0]; set.ne.u32.u32 %r110, %r109, %r22; neg.s32 %r111, %r110; and.b32 %r112, %r108, %r111; mov.u32 %r113, 0; setp.eq.s32 %p19, %r112, %r113; @%p19 bra $Lt_0_44802; .loc 3 170 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_9; $Lt_0_44802: mov.u16 %rh11, %nctaid.y; mul.wide.u16 %r114, %rh11, %rh1; cvt.u32.u16 %r115, %ntid.z; mul.lo.u32 %r34, %r115, %r17; mul.lo.u32 %r116, %r34, %r114; ld.const.s32 %r117, [printfBufferLength]; div.s32 %r118, %r117, %r116; and.b32 %r38, %r118, -256; mov.u32 %r119, 511; setp.gt.u32 %p20, %r38, %r119; @%p20 bra $Lt_0_45314; .loc 3 188 0 mov.u64 %rd3, 0; bra.uni $LDWendi__Z19getNextPrintfBufPtrv_212_9; $Lt_0_45314: .loc 3 192 0 mul.lo.u32 %r120, %r34, %r3; add.u32 %r41, %r120, %r22; mul.lo.u32 %r121, %r41, %r38; cvt.u64.u32 %rd32, %r121; ld.const.u64 %rd5, [globalPrintfBuffer]; add.u64 %rd6, %rd32, %rd5; ld.global.v2.u16 {%r122,%r44}, [%rd6+0]; ld.global.v2.u32 {%r45,%r46}, [%rd6+8]; mov.u32 %r123, 51216; setp.eq.u32 %p21, %r122, %r123; @%p21 bra $Lt_0_45826; .loc 3 198 0 mov.s32 %r45, %r38; .loc 3 200 0 mov.u32 %r124, 51216; st.global.v2.u16 [%rd6+0], {%r124,%r44}; st.global.u32 [%rd6+4], %r41; mov.u32 %r125, 0; st.global.v2.u32 [%rd6+8], {%r38,%r125}; .loc 3 205 0 st.global.u32 [%rd5+8], %r38; mov.u32 %r46, 0; $Lt_0_45826: .loc 3 214 0 add.u32 %r126, %r46, 256; setp.ge.u32 %p22, %r126, %r45; mov.u32 %r127, 256; selp.u32 %r128, %r127, %r126, %p22; st.global.u32 [%rd6+12], %r128; .loc 3 215 0 cvt.u64.u32 %rd33, %r128; add.u64 %rd3, %rd33, %rd6; $LDWendi__Z19getNextPrintfBufPtrv_212_9: .loc 3 401 0 mov.u64 %rd34, 0; setp.ne.u64 %p23, %rd3, %rd34; @%p23 bra $Lt_0_46338; bra.uni $LDWendi__Z17writePrintfHeaderPcS__212_1; $Lt_0_46338: .loc 3 403 0 cvt.u32.u16 %r129, %ctaid.x; mov.u64 %rd35, -8; setp.ne.u64 %p24, %rd3, %rd35; @%p24 bra $Lt_0_46850; .loc 3 328 0 mov.u64 %rd10, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_7; $Lt_0_46850: .loc 3 331 0 mov.s32 %r130, 4; st.global.s32 [%rd3+8], %r130; .loc 3 333 0 st.global.u32 [%rd3+16], %r129; .loc 3 335 0 mov.s16 %rh12, 0; st.global.s8 [%rd3+24], %rh12; .loc 3 337 0 add.u64 %rd10, %rd3, 24; $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_7: .loc 3 404 0 add.u64 %rd11, %rd3, 256; cvt.u32.u16 %r131, %ctaid.y; mov.u64 %rd36, 0; set.eq.u32.u64 %r132, %rd10, %rd36; neg.s32 %r133, %r132; add.u64 %rd37, %rd10, 8; set.le.u32.u64 %r134, %rd11, %rd37; neg.s32 %r135, %r134; or.b32 %r136, %r133, %r135; mov.u32 %r137, 0; setp.eq.s32 %p25, %r136, %r137; @%p25 bra $Lt_0_47362; .loc 3 328 0 mov.u64 %rd14, 0; bra.uni $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_5; $Lt_0_47362: .loc 3 331 0 mov.s32 %r138, 4; st.global.s32 [%rd10+0], %r138; .loc 3 333 0 st.global.u32 [%rd10+8], %r131; .loc 3 335 0 mov.s16 %rh13, 0; st.global.s8 [%rd10+16], %rh13; .loc 3 337 0 add.u64 %rd14, %rd10, 16; $LDWendi__Z7copyArgIjEPcS0_RT_S0__212_5: .loc 3 406 0 mov.u64 %rd38, 0; set.eq.u32.u64 %r139, %rd14, %rd38; neg.s32 %r140, %r139; set.ge.u32.u64 %r141, %rd14, %rd11; neg.s32 %r142, %r141; or.b32 %r143, %r140, %r142; mov.u32 %r144, 0; setp.eq.s32 %p26, %r143, %r144; @%p26 bra $Lt_0_48642; .loc 3 261 0 mov.u64 %rd16, 0; bra.uni $LDWendi__Z15cuPrintfStrncpyPcPKciS__212_2; $Lt_0_48642: add.u64 %rd17, %rd14, 8; mov.s32 %r69, 0; mov.u64 %rd39, __constant460; $Lt_0_48898: // Loop body line 261, nesting depth: 1, iterations: 256 mov.s64 %rd19, %rd17; .loc 3 274 0 setp.le.u64 %p27, %rd11, %rd17; @%p27 bra $Lt_0_50690; .loc 3 276 0 add.s32 %r69, %r69, 1; .loc 3 277 0 cvt.u64.s32 %rd40, %r69; add.u64 %rd41, %rd14, %rd40; add.u64 %rd42, %rd41, 8; mov.s64 %rd17, %rd42; mov.s64 %rd19, %rd42; add.u64 %rd43, %rd40, %rd39; ld.const.s8 %rh14, [%rd43+-1]; st.global.s8 [%rd42+-1], %rh14; .loc 3 279 0 ld.const.s8 %r145, [%rd43+-1]; mov.u32 %r146, 0; setp.eq.s32 %p28, %r145, %r146; @%p28 bra $Lt_0_50690; mov.u32 %r147, 256; setp.ne.s32 %p29, %r69, %r147; @%p29 bra $Lt_0_48898; add.u64 %rd19, %rd14, 264; $Lt_0_50690: $Lt_0_770: .loc 3 283 0 setp.gt.u64 %p14, %rd11, %rd19; @!%p14 bra $Lt_0_51202; and.b64 %rd44, %rd19, 7; mov.u64 %rd45, 0; setp.eq.u64 %p30, %rd44, %rd45; @%p30 bra $Lt_0_51202; $L_0_37122: .loc 3 285 0 add.s32 %r69, %r69, 1; .loc 3 286 0 add.u64 %rd19, %rd19, 1; mov.s16 %rh15, 0; st.global.s8 [%rd19+-1], %rh15; .loc 3 283 0 setp.gt.u64 %p14, %rd11, %rd19; @!%p14 bra $Lt_0_51202; and.b64 %rd46, %rd19, 7; mov.u64 %rd47, 0; setp.ne.u64 %p31, %rd46, %rd47; @%p31 bra $L_0_37122; $Lt_0_51202: $L_0_37378: .loc 3 288 0 st.global.s32 [%rd14+0], %r69; .loc 3 289 0 mov.u64 %rd48, 0; selp.u64 %rd16, %rd19, %rd48, %p14; $LDWendi__Z15cuPrintfStrncpyPcPKciS__212_2: .loc 3 241 0 mov.u32 %r148, 51217; mov.s64 %rd49, 0; mov.u64 %rd50, 0; setp.ne.u64 %p32, %rd16, %rd50; selp.s64 %rd51, %rd14, %rd49, %p32; cvt.s32.s64 %r149, %rd51; cvt.s32.s64 %r150, %rd3; sub.s32 %r151, %r149, %r150; cvt.u32.u16 %r152, %nctaid.x; mul.lo.u32 %r153, %r152, %r131; add.u32 %r154, %r129, %r153; cvt.u32.u16 %r13, %ntid.x; cvt.u32.u16 %r155, %tid.x; cvt.u32.u16 %r156, %tid.y; mul.lo.u32 %r157, %r156, %r13; cvt.u32.u16 %r158, %tid.z; cvt.u32.u16 %r159, %ntid.y; mul.lo.u32 %r160, %r159, %r13; mul.lo.u32 %r161, %r158, %r160; add.u32 %r162, %r157, %r161; add.u32 %r163, %r155, %r162; st.global.v4.u16 [%rd3+0], {%r148,%r151,%r154,%r163}; $LDWendi__Z17writePrintfHeaderPcS__212_1: .loc 17 16 0 exit; $LDWend__Z9printTestv: } // _Z9printTestv P` 8cudaPrintfExample.cuELF3@  @8@ @ o\ 6$a4J I}M U u+}.shstrtab.strtab.symtab.nv.global.init.nv.global.text._Z9printTestv.nv.info._Z9printTestv.nv.constant1._Z9printTestv.nv.constant14.rel.nv.constant14.nv.constant0_Z9printTestvrestrictRulesprintfBufferPtrprintfBufferLengthglobalPrintfBuffer__constant460__constant459     -@S,a ǀ' Ѐ|0d |0dG )-p! Ѐ ЀL H#0GAl N`!0Rd)-pB  0 D@ @0)Ѐ@ ` @ ǀ'ЀЀ  0GAl0Rd`)-`F@ 0H# F`G!J@  @`G 0 `G$0 Ѐ@0Ad)-` @`G 0 `G  @ ` 0 `! Ѐ  Ѐ0@d@1W"  Ѐ  Ѐ    Ѐ !Ѐ   0؇d @$ ) ЇЀ-@0|0؇d|0ؒd0؇@d0ؒ@d10  Ѐ 1Ѐ  !AЀ @0 !GA!B @0' @00d0d%Ѐ|0d |0d!0GdЀ ЀGЀ@ G   Ѐ  Ѐ! 9!GBЀ  @00d 0d%Ѐ |0d|0d!0dЀ ЀG Gp   p !BP @-@00؇d0d0XQd $P-  !0- !B!@0)%? 5@0'؇ Ѐ $P0GAl ! !ЇB%@0P 0d0dЀ 0dG G !@!@!!C!? @0Ѐ 0d0dЀ0dG G - % P!@ &P   - Ѐ B % 0D@ |0؇d @@|0ؒd 00  ``  L . 0 H`  G  Ѐ)-! Ѐ ЀL H#0GAl N`!0Bd)-B  0 D@ @0)Ѐ@ ` @ ǀ'ЀЀ  0GAl0Bd)-F@ 0H# F`G!J@  @`G 0 `G$0 Ѐ@0Ad )- @`G 0 `G  @ ` 0 `! Ѐ  Ѐ0؇@d0@1G0  Ѐ  Ѐ    Ѐ !Ѐ   0ȇd @$ ) Ѐ-@0|0ȇd|0Ȃd00ȇ@d0Ȃ@d1L p  Ѐ 1Ѐ  !AЀ @0 !GA!B @0 @00d0d!Ѐ|0d |0d0GdЀ ЀGGN @Ѐ  Ѐ! 9!GBЀ  @00d 0d%Ѐ |0d|0d!0dЀ ЀG G   !B@-@00ȇd0d0HAd $P-  !0-X !B!@0)%? 5@0ȇ Ѐ $P0GAl! !B%@0P 0d0dЀ 0dG G @@@@@@!!C!? @0Ѐ 0d0dЀ0dG G - % P@0P   - Ѐ B % 0D@  @|0ȇd 0|0Ȃd H@ Ѐ@ ` 0 @ H`!  G0    Ѐ0 G GD!G%GD%>  G)@)`)0)`)@ )GD% %G G% @% `G%0% `G!0 0d Ѐ0 0G, Ѐ |0Gl 0  $ threadIdx.x %d threadIdx.y %d blockIdx.x %d blockIdx.y %d  ` $qqu_Z9printTestvprintfBufferLengthrestrictRules;ݓ<\|?dԖ<\E|Od*<o\|>dt 4zPRx @ $0AC DtEAC hdAC !AC (AC %AC AC AC $4AC DAC Rd}8AC LAC oAC UAC JAC AC $ AC DjkAC JdEAC l EAC l0EAC lUEAC ]zAC zRx h$4`Q_@F\ȝ@!+9 p @ @@oP@@x@  ``@0@0 o@oo@`}` @ @ @ @ @ @ @ @ @. @> @N @^ @n @~ @ @ @ @ @ @ @ @@@CbF@@GCC: (Ubuntu 4.4.3-4ubuntu5) 4.4.3.symtab.strtab.shstrtab.interp.note.ABI-tag.note.gnu.build-id.gnu.hash.dynsym.dynstr.gnu.version.gnu.version_r.rela.dyn.rela.plt.init.text.fini.rodata.eh_frame_hdr.eh_frame.ctors.dtors.jcr.dynamic.got.got.plt.data.nvFatBinSegment.bss.comment8@8#T@T 1t@t$H@DoP@P(N x@xpV@^o@4ko@Pz0@00`@` p @p  @ p @ @@] Ly@Ly(z@(zt}`}}`}}`}}`}``` Ѐ`Ѐ`X0#  G 8@T@t@@P@x@@@ @ 0@ `@ p @ @ @@@Ly@(z@}`}`}`}````Ѐ``, @}`*}`8}`EP @[`j`x @}`}@}`p@ @E)`A`Y`lȀ`) @B @!c @( @%` ` (`0`2 @4L @k8`@@@Ѐ`@`@kMHy@c`y`}`}`}` `@ @  "F@0"4K@Q@r@P@ `#"@B @Q"v@EE@L}`@ `"@EWs@ H`"@E ` @?T @8c"E@E @`@o@ p @call_gmon_startcrtstuff.c__CTOR_LIST____DTOR_LIST____JCR_LIST____do_global_dtors_auxcompleted.7382dtor_idx.7384frame_dummy__CTOR_END____FRAME_END____JCR_END____do_global_ctors_auxtmpxft_000074e2_00000000-1_cudaPrintfExample.cudafe1.cpp_ZL18cudaMemcpyToSymbolPcPKvmm14cudaMemcpyKind_ZL18globalPrintfBuffer_ZL18printfBufferLength_ZL13restrictRules_ZL15printfBufferPtr_ZL19getNextPrintfBufPtrv_ZL17writePrintfHeaderPcS__ZL15cuPrintfStrncpyPcPKciS__ZL7copyArgPcPKcS__ZL9printf_fp_ZL15printfbuf_start_ZL16printfbuf_device_ZL13printfbuf_len_ZL16outputPrintfDataPcS__ZL15doPrintfDisplayiiPcS_S_S__ZL20__cudaFatCubinHandle_ZL26__cudaUnregisterBinaryUtilvfatbinData_ZL15__fatDeviceText_ZZ27__device_stub__Z9printTestvvE3__f_ZL89__sti____cudaRegisterAll_52_tmpxft_000074e2_00000000_4_cudaPrintfExample_cpp1_ii_51e07f2fv_ZL16CUPRINTF_MAX_LEN_GLOBAL_OFFSET_TABLE___dso_handle__init_array_end__init_array_start_DYNAMICdata_start__libc_csu_fini_startcudaFree__cudaRegisterVar__gmon_start___Jv_RegisterClasses_ZN4dim3C1Ejjjexit@@GLIBC_2.2.5__cudaRegisterFunction_fini_Z27__device_stub__Z9printTestvv__libc_start_main@@GLIBC_2.2.5cudaMemcpyToSymbol__cxa_atexit@@GLIBC_2.2.5cudaMemset_IO_stdin_usedatexitcudaConfigureCall__cudaUnregisterFatBinary__data_startcudaLaunch_Z10cudaLaunchIcE9cudaErrorPT__Z8cuPrintfPKc_Z18cudaMemcpyToSymbolI19cuPrintfRestrictionE9cudaErrorRKT_PKvmm14cudaMemcpyKindcudaPrintfEndfputs@@GLIBC_2.2.5__DTOR_END____libc_csu_init__cudaRegisterFatBinarystrchr@@GLIBC_2.2.5__bss_start_Z18cudaMemcpyToSymbolIPVcE9cudaErrorRKT_PKvmm14cudaMemcpyKind__stack_chk_fail@@GLIBC_2.4_Z9printTestv_end_Z18cudaMemcpyToSymbolIiE9cudaErrorRKT_PKvmm14cudaMemcpyKindstrcspn@@GLIBC_2.2.5fwrite@@GLIBC_2.2.5cudaMemcpycudaMemcpyFromSymbolcudaMalloc_edata__gxx_personality_v0@@CXXABI_1.3fprintf@@GLIBC_2.2.5cudaPrintfInit_Z20cudaMemcpyFromSymbolIPVcE9cudaErrorPvRKT_mm14cudaMemcpyKind_Z16cuPrintfRestrictiistdout@@GLIBC_2.2.5cudaPrintfDisplaymain_init