util_ptx.h

Go to the documentation of this file.

static const char* util_ptx = 
"       .version 1.4\n"
"       .target sm_10, map_f64_to_f32\n"
"       // compiled with /usr/local/cuda/open64/lib//be\n"
"       // nvopencc 3.1 built on 2010-06-07\n"
"\n"
"       //-----------------------------------------------------------\n"
"       // Compiling /tmp/tmpxft_00005c16_00000000-7_util.cpp3.i (/tmp/ccBI#.kEa25g)\n"
"       //-----------------------------------------------------------\n"
"\n"
"       //-----------------------------------------------------------\n"
"       // Options:\n"
"       //-----------------------------------------------------------\n"
"       //  Target:ptx, ISA:sm_10, Endian:little, Pointer Size:64\n"
"       //  -O3 (Optimization level)\n"
"       //  -g0 (Debug level)\n"
"       //  -m2 (Report advisories)\n"
"       //-----------------------------------------------------------\n"
"\n"
"       .file   1       \"<command-line>\"\n"
"       .file   2       \"/tmp/tmpxft_00005c16_00000000-6_util.cudafe2.gpu\"\n"
"       .file   3       \"/usr/lib/gcc/x86_64-linux-gnu/4.4.3/include/stddef.h\"\n"
"       .file   4       \"/usr/local/cuda/bin/../include/crt/device_runtime.h\"\n"
"       .file   5       \"/usr/local/cuda/bin/../include/host_defines.h\"\n"
"       .file   6       \"/usr/local/cuda/bin/../include/builtin_types.h\"\n"
"       .file   7       \"/usr/local/cuda/bin/../include/device_types.h\"\n"
"       .file   8       \"/usr/local/cuda/bin/../include/driver_types.h\"\n"
"       .file   9       \"/usr/local/cuda/bin/../include/surface_types.h\"\n"
"       .file   10      \"/usr/local/cuda/bin/../include/texture_types.h\"\n"
"       .file   11      \"/usr/local/cuda/bin/../include/vector_types.h\"\n"
"       .file   12      \"/usr/local/cuda/bin/../include/device_launch_parameters.h\"\n"
"       .file   13      \"/usr/local/cuda/bin/../include/crt/storage_class.h\"\n"
"       .file   14      \"/usr/include/bits/types.h\"\n"
"       .file   15      \"/usr/include/time.h\"\n"
"       .file   16      \"util.cu\"\n"
"       .file   17      \"/usr/local/cuda/bin/../include/common_functions.h\"\n"
"       .file   18      \"/usr/local/cuda/bin/../include/math_functions.h\"\n"
"       .file   19      \"/usr/local/cuda/bin/../include/math_constants.h\"\n"
"       .file   20      \"/usr/local/cuda/bin/../include/device_functions.h\"\n"
"       .file   21      \"/usr/local/cuda/bin/../include/sm_11_atomic_functions.h\"\n"
"       .file   22      \"/usr/local/cuda/bin/../include/sm_12_atomic_functions.h\"\n"
"       .file   23      \"/usr/local/cuda/bin/../include/sm_13_double_functions.h\"\n"
"       .file   24      \"/usr/local/cuda/bin/../include/sm_20_atomic_functions.h\"\n"
"       .file   25      \"/usr/local/cuda/bin/../include/sm_20_intrinsics.h\"\n"
"       .file   26      \"/usr/local/cuda/bin/../include/surface_functions.h\"\n"
"       .file   27      \"/usr/local/cuda/bin/../include/texture_fetch_functions.h\"\n"
"       .file   28      \"/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h\"\n"
"\n"
"\n"
"       .entry ow_ew_multiply_uint_float (\n"
"               .param .u64 __cudaparm_ow_ew_multiply_uint_float_data1,\n"
"               .param .u64 __cudaparm_ow_ew_multiply_uint_float_data2,\n"
"               .param .u32 __cudaparm_ow_ew_multiply_uint_float_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<8>;\n"
"       .reg .u64 %rd<8>;\n"
"       .reg .f32 %f<5>;\n"
"       .reg .pred %p<3>;\n"
"       .loc    16      75      0\n"
"$LDWbegin_ow_ew_multiply_uint_float:\n"
"       .loc    16      77      0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       ld.param.u32    %r4, [__cudaparm_ow_ew_multiply_uint_float_n];\n"
"       setp.le.u32     %p1, %r4, %r3;\n"
"       @%p1 bra        $Lt_0_1026;\n"
"       .loc    16      68      0\n"
"       cvt.u64.u32     %rd1, %r3;\n"
"       mul.wide.u32    %rd2, %r3, 4;\n"
"       ld.param.u64    %rd3, [__cudaparm_ow_ew_multiply_uint_float_data1];\n"
"       add.u64         %rd4, %rd3, %rd2;\n"
"       ld.global.u32   %r5, [%rd4+0];\n"
"       .loc    16      69      0\n"
"       cvt.rn.f32.u32  %f1, %r5;\n"
"       ld.param.u64    %rd5, [__cudaparm_ow_ew_multiply_uint_float_data2];\n"
"       add.u64         %rd6, %rd5, %rd2;\n"
"       ld.global.f32   %f2, [%rd6+0];\n"
"       mul.f32         %f3, %f1, %f2;\n"
"       cvt.rzi.u32.f32         %r6, %f3;\n"
"       .loc    16      70      0\n"
"       st.global.u32   [%rd4+0], %r6;\n"
"$Lt_0_1026:\n"
"       .loc    16      78      0\n"
"       exit;\n"
"$LDWend_ow_ew_multiply_uint_float:\n"
"       } // ow_ew_multiply_uint_float\n"
"\n"
"       .entry ow_ew_multiply_float (\n"
"               .param .u64 __cudaparm_ow_ew_multiply_float_data1,\n"
"               .param .u64 __cudaparm_ow_ew_multiply_float_data2,\n"
"               .param .u32 __cudaparm_ow_ew_multiply_float_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<6>;\n"
"       .reg .u64 %rd<8>;\n"
"       .reg .f32 %f<5>;\n"
"       .reg .pred %p<3>;\n"
"       .loc    16      81      0\n"
"$LDWbegin_ow_ew_multiply_float:\n"
"       .loc    16      83      0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       ld.param.u32    %r4, [__cudaparm_ow_ew_multiply_float_n];\n"
"       setp.le.u32     %p1, %r4, %r3;\n"
"       @%p1 bra        $Lt_1_1026;\n"
"       .loc    16      70      0\n"
"       cvt.u64.u32     %rd1, %r3;\n"
"       mul.wide.u32    %rd2, %r3, 4;\n"
"       ld.param.u64    %rd3, [__cudaparm_ow_ew_multiply_float_data1];\n"
"       add.u64         %rd4, %rd3, %rd2;\n"
"       ld.global.f32   %f1, [%rd4+0];\n"
"       ld.param.u64    %rd5, [__cudaparm_ow_ew_multiply_float_data2];\n"
"       add.u64         %rd6, %rd5, %rd2;\n"
"       ld.global.f32   %f2, [%rd6+0];\n"
"       mul.f32         %f3, %f1, %f2;\n"
"       st.global.f32   [%rd4+0], %f3;\n"
"$Lt_1_1026:\n"
"       .loc    16      84      0\n"
"       exit;\n"
"$LDWend_ow_ew_multiply_float:\n"
"       } // ow_ew_multiply_float\n"
"\n"
"       .entry ow_ew_multiply_uint (\n"
"               .param .u64 __cudaparm_ow_ew_multiply_uint_data1,\n"
"               .param .u64 __cudaparm_ow_ew_multiply_uint_data2,\n"
"               .param .u32 __cudaparm_ow_ew_multiply_uint_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<9>;\n"
"       .reg .u64 %rd<8>;\n"
"       .reg .pred %p<3>;\n"
"       .loc    16      87      0\n"
"$LDWbegin_ow_ew_multiply_uint:\n"
"       .loc    16      89      0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       ld.param.u32    %r4, [__cudaparm_ow_ew_multiply_uint_n];\n"
"       setp.le.u32     %p1, %r4, %r3;\n"
"       @%p1 bra        $Lt_2_1026;\n"
"       .loc    16      70      0\n"
"       cvt.u64.u32     %rd1, %r3;\n"
"       mul.wide.u32    %rd2, %r3, 4;\n"
"       ld.param.u64    %rd3, [__cudaparm_ow_ew_multiply_uint_data1];\n"
"       add.u64         %rd4, %rd3, %rd2;\n"
"       ld.global.u32   %r5, [%rd4+0];\n"
"       ld.param.u64    %rd5, [__cudaparm_ow_ew_multiply_uint_data2];\n"
"       add.u64         %rd6, %rd5, %rd2;\n"
"       ld.global.u32   %r6, [%rd6+0];\n"
"       mul.lo.u32      %r7, %r5, %r6;\n"
"       st.global.u32   [%rd4+0], %r7;\n"
"$Lt_2_1026:\n"
"       .loc    16      90      0\n"
"       exit;\n"
"$LDWend_ow_ew_multiply_uint:\n"
"       } // ow_ew_multiply_uint\n"
"\n"
"       .entry ow_ew_multiply_int (\n"
"               .param .u64 __cudaparm_ow_ew_multiply_int_data1,\n"
"               .param .u64 __cudaparm_ow_ew_multiply_int_data2,\n"
"               .param .u32 __cudaparm_ow_ew_multiply_int_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<9>;\n"
"       .reg .u64 %rd<8>;\n"
"       .reg .pred %p<3>;\n"
"       .loc    16      93      0\n"
"$LDWbegin_ow_ew_multiply_int:\n"
"       .loc    16      95      0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       ld.param.u32    %r4, [__cudaparm_ow_ew_multiply_int_n];\n"
"       setp.le.u32     %p1, %r4, %r3;\n"
"       @%p1 bra        $Lt_3_1026;\n"
"       .loc    16      70      0\n"
"       cvt.u64.u32     %rd1, %r3;\n"
"       mul.wide.u32    %rd2, %r3, 4;\n"
"       ld.param.u64    %rd3, [__cudaparm_ow_ew_multiply_int_data1];\n"
"       add.u64         %rd4, %rd3, %rd2;\n"
"       ld.global.s32   %r5, [%rd4+0];\n"
"       ld.param.u64    %rd5, [__cudaparm_ow_ew_multiply_int_data2];\n"
"       add.u64         %rd6, %rd5, %rd2;\n"
"       ld.global.s32   %r6, [%rd6+0];\n"
"       mul.lo.s32      %r7, %r5, %r6;\n"
"       st.global.s32   [%rd4+0], %r7;\n"
"$Lt_3_1026:\n"
"       .loc    16      96      0\n"
"       exit;\n"
"$LDWend_ow_ew_multiply_int:\n"
"       } // ow_ew_multiply_int\n"
"\n"
"       .entry detect_changes_float (\n"
"               .param .u64 __cudaparm_detect_changes_float_data,\n"
"               .param .u32 __cudaparm_detect_changes_float_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<8>;\n"
"       .reg .u64 %rd<10>;\n"
"       .reg .f32 %f<8>;\n"
"       .reg .pred %p<5>;\n"
"       .loc    16      132     0\n"
"$LDWbegin_detect_changes_float:\n"
"       .loc    16      134     0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       mov.u32         %r4, 0;\n"
"       setp.ne.u32     %p1, %r3, %r4;\n"
"       @%p1 bra        $Lt_4_3074;\n"
"       .loc    16      114     0\n"
"       mov.f32         %f1, 0f00000000;        // 0\n"
"       ld.param.u64    %rd1, [__cudaparm_detect_changes_float_data];\n"
"       cvt.u64.u32     %rd2, %r3;\n"
"       mul.wide.u32    %rd3, %r3, 4;\n"
"       add.u64         %rd4, %rd1, %rd3;\n"
"       st.global.f32   [%rd4+0], %f1;\n"
"       bra.uni         $LDWendi__Z23detect_changes_dispatchIfEvPT_j_214_1;\n"
"$Lt_4_3074:\n"
"       ld.param.u32    %r5, [__cudaparm_detect_changes_float_n];\n"
"       add.u32         %r6, %r3, 1;\n"
"       setp.le.u32     %p2, %r5, %r6;\n"
"       @%p2 bra        $Lt_4_3586;\n"
"       .loc    16      120     0\n"
"       ld.param.u64    %rd5, [__cudaparm_detect_changes_float_data];\n"
"       cvt.u64.u32     %rd6, %r3;\n"
"       mul.wide.u32    %rd7, %r3, 4;\n"
"       add.u64         %rd8, %rd5, %rd7;\n"
"       mov.f32         %f2, 0f00000000;        // 0\n"
"       mov.f32         %f3, 0f3f800000;        // 1\n"
"       ld.global.f32   %f4, [%rd8+0];\n"
"       ld.global.f32   %f5, [%rd8+4];\n"
"       setp.eq.f32     %p3, %f4, %f5;\n"
"       selp.f32        %f6, %f2, %f3, %p3;\n"
"       st.global.f32   [%rd8+4], %f6;\n"
"$Lt_4_3586:\n"
"$LDWendi__Z23detect_changes_dispatchIfEvPT_j_214_1:\n"
"       .loc    16      135     0\n"
"       exit;\n"
"$LDWend_detect_changes_float:\n"
"       } // detect_changes_float\n"
"\n"
"       .entry detect_changes_int (\n"
"               .param .u64 __cudaparm_detect_changes_int_data,\n"
"               .param .u32 __cudaparm_detect_changes_int_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<13>;\n"
"       .reg .u64 %rd<10>;\n"
"       .reg .pred %p<4>;\n"
"       .loc    16      138     0\n"
"$LDWbegin_detect_changes_int:\n"
"       .loc    16      140     0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       mov.u32         %r4, 0;\n"
"       setp.ne.u32     %p1, %r3, %r4;\n"
"       @%p1 bra        $Lt_5_3074;\n"
"       .loc    16      114     0\n"
"       mov.s32         %r5, 0;\n"
"       ld.param.u64    %rd1, [__cudaparm_detect_changes_int_data];\n"
"       cvt.u64.u32     %rd2, %r3;\n"
"       mul.wide.u32    %rd3, %r3, 4;\n"
"       add.u64         %rd4, %rd1, %rd3;\n"
"       st.global.s32   [%rd4+0], %r5;\n"
"       bra.uni         $LDWendi__Z23detect_changes_dispatchIiEvPT_j_215_1;\n"
"$Lt_5_3074:\n"
"       ld.param.u32    %r6, [__cudaparm_detect_changes_int_n];\n"
"       add.u32         %r7, %r3, 1;\n"
"       setp.le.u32     %p2, %r6, %r7;\n"
"       @%p2 bra        $Lt_5_3586;\n"
"       .loc    16      120     0\n"
"       ld.param.u64    %rd5, [__cudaparm_detect_changes_int_data];\n"
"       cvt.u64.u32     %rd6, %r3;\n"
"       mul.wide.u32    %rd7, %r3, 4;\n"
"       add.u64         %rd8, %rd5, %rd7;\n"
"       ld.global.s32   %r8, [%rd8+0];\n"
"       ld.global.s32   %r9, [%rd8+4];\n"
"       set.ne.u32.s32  %r10, %r8, %r9;\n"
"       neg.s32         %r11, %r10;\n"
"       st.global.s32   [%rd8+4], %r11;\n"
"$Lt_5_3586:\n"
"$LDWendi__Z23detect_changes_dispatchIiEvPT_j_215_1:\n"
"       .loc    16      141     0\n"
"       exit;\n"
"$LDWend_detect_changes_int:\n"
"       } // detect_changes_int\n"
"\n"
"       .entry detect_changes_uint (\n"
"               .param .u64 __cudaparm_detect_changes_uint_data,\n"
"               .param .u32 __cudaparm_detect_changes_uint_n)\n"
"       {\n"
"       .reg .u16 %rh<4>;\n"
"       .reg .u32 %r<13>;\n"
"       .reg .u64 %rd<10>;\n"
"       .reg .pred %p<4>;\n"
"       .loc    16      145     0\n"
"$LDWbegin_detect_changes_uint:\n"
"       .loc    16      147     0\n"
"       mov.u16         %rh1, %ctaid.x;\n"
"       mov.u16         %rh2, %ntid.x;\n"
"       mul.wide.u16    %r1, %rh1, %rh2;\n"
"       cvt.u32.u16     %r2, %tid.x;\n"
"       add.u32         %r3, %r2, %r1;\n"
"       mov.u32         %r4, 0;\n"
"       setp.ne.u32     %p1, %r3, %r4;\n"
"       @%p1 bra        $Lt_6_3074;\n"
"       .loc    16      114     0\n"
"       mov.u32         %r5, 0;\n"
"       ld.param.u64    %rd1, [__cudaparm_detect_changes_uint_data];\n"
"       cvt.u64.u32     %rd2, %r3;\n"
"       mul.wide.u32    %rd3, %r3, 4;\n"
"       add.u64         %rd4, %rd1, %rd3;\n"
"       st.global.u32   [%rd4+0], %r5;\n"
"       bra.uni         $LDWendi__Z23detect_changes_dispatchIjEvPT_j_216_1;\n"
"$Lt_6_3074:\n"
"       ld.param.u32    %r6, [__cudaparm_detect_changes_uint_n];\n"
"       add.u32         %r7, %r3, 1;\n"
"       setp.le.u32     %p2, %r6, %r7;\n"
"       @%p2 bra        $Lt_6_3586;\n"
"       .loc    16      120     0\n"
"       ld.param.u64    %rd5, [__cudaparm_detect_changes_uint_data];\n"
"       cvt.u64.u32     %rd6, %r3;\n"
"       mul.wide.u32    %rd7, %r3, 4;\n"
"       add.u64         %rd8, %rd5, %rd7;\n"
"       ld.global.u32   %r8, [%rd8+0];\n"
"       ld.global.u32   %r9, [%rd8+4];\n"
"       set.ne.u32.u32  %r10, %r8, %r9;\n"
"       neg.s32         %r11, %r10;\n"
"       st.global.u32   [%rd8+4], %r11;\n"
"$Lt_6_3586:\n"
"$LDWendi__Z23detect_changes_dispatchIjEvPT_j_216_1:\n"
"       .loc    16      148     0\n"
"       exit;\n"
"$LDWend_detect_changes_uint:\n"
"       } // detect_changes_uint\n"
"\n"
;