89 lines
2.2 KiB
C
89 lines
2.2 KiB
C
const char * pair_gpu_dev_kernel =
|
|
" .version 2.3\n"
|
|
" .target sm_20\n"
|
|
" .address_size 64\n"
|
|
" .entry kernel_zero (\n"
|
|
" .param .u64 __cudaparm_kernel_zero_mem,\n"
|
|
" .param .s32 __cudaparm_kernel_zero_numel)\n"
|
|
" {\n"
|
|
" .reg .u32 %r<9>;\n"
|
|
" .reg .u64 %rd<6>;\n"
|
|
" .reg .pred %p<3>;\n"
|
|
" .loc 16 95 0\n"
|
|
"$LDWbegin_kernel_zero:\n"
|
|
" mov.u32 %r1, %ctaid.x;\n"
|
|
" mov.u32 %r2, %ntid.x;\n"
|
|
" mul.lo.u32 %r3, %r1, %r2;\n"
|
|
" mov.u32 %r4, %tid.x;\n"
|
|
" add.u32 %r5, %r4, %r3;\n"
|
|
" ld.param.s32 %r6, [__cudaparm_kernel_zero_numel];\n"
|
|
" setp.le.s32 %p1, %r6, %r5;\n"
|
|
" @%p1 bra $Lt_0_1026;\n"
|
|
" .loc 16 99 0\n"
|
|
" mov.s32 %r7, 0;\n"
|
|
" ld.param.u64 %rd1, [__cudaparm_kernel_zero_mem];\n"
|
|
" cvt.s64.s32 %rd2, %r5;\n"
|
|
" mul.wide.s32 %rd3, %r5, 4;\n"
|
|
" add.u64 %rd4, %rd1, %rd3;\n"
|
|
" st.global.s32 [%rd4+0], %r7;\n"
|
|
"$Lt_0_1026:\n"
|
|
" .loc 16 100 0\n"
|
|
" exit;\n"
|
|
"$LDWend_kernel_zero:\n"
|
|
" }\n"
|
|
" .entry kernel_info (\n"
|
|
" .param .u64 __cudaparm_kernel_info_info)\n"
|
|
" {\n"
|
|
" .reg .u32 %r<16>;\n"
|
|
" .reg .u64 %rd<3>;\n"
|
|
" .loc 16 102 0\n"
|
|
"$LDWbegin_kernel_info:\n"
|
|
" .loc 16 103 0\n"
|
|
" ld.param.u64 %rd1, [__cudaparm_kernel_info_info];\n"
|
|
" mov.s32 %r1, 200;\n"
|
|
" st.global.s32 [%rd1+0], %r1;\n"
|
|
" .loc 16 104 0\n"
|
|
" mov.s32 %r2, 32;\n"
|
|
" st.global.s32 [%rd1+4], %r2;\n"
|
|
" .loc 16 105 0\n"
|
|
" mov.s32 %r3, 32;\n"
|
|
" st.global.s32 [%rd1+8], %r3;\n"
|
|
" .loc 16 106 0\n"
|
|
" mov.s32 %r4, 1;\n"
|
|
" st.global.s32 [%rd1+12], %r4;\n"
|
|
" .loc 16 107 0\n"
|
|
" mov.s32 %r5, 8;\n"
|
|
" st.global.s32 [%rd1+16], %r5;\n"
|
|
" .loc 16 108 0\n"
|
|
" mov.s32 %r6, 64;\n"
|
|
" st.global.s32 [%rd1+20], %r6;\n"
|
|
" .loc 16 109 0\n"
|
|
" mov.s32 %r7, 128;\n"
|
|
" st.global.s32 [%rd1+24], %r7;\n"
|
|
" .loc 16 110 0\n"
|
|
" mov.s32 %r8, 11;\n"
|
|
" st.global.s32 [%rd1+28], %r8;\n"
|
|
" .loc 16 111 0\n"
|
|
" mov.s32 %r9, 8;\n"
|
|
" st.global.s32 [%rd1+32], %r9;\n"
|
|
" .loc 16 112 0\n"
|
|
" mov.s32 %r10, 128;\n"
|
|
" st.global.s32 [%rd1+36], %r10;\n"
|
|
" .loc 16 113 0\n"
|
|
" mov.s32 %r11, 128;\n"
|
|
" st.global.s32 [%rd1+40], %r11;\n"
|
|
" .loc 16 114 0\n"
|
|
" mov.s32 %r12, 128;\n"
|
|
" st.global.s32 [%rd1+44], %r12;\n"
|
|
" .loc 16 115 0\n"
|
|
" mov.s32 %r13, 128;\n"
|
|
" st.global.s32 [%rd1+48], %r13;\n"
|
|
" .loc 16 116 0\n"
|
|
" mov.s32 %r14, 8;\n"
|
|
" st.global.s32 [%rd1+52], %r14;\n"
|
|
" .loc 16 117 0\n"
|
|
" exit;\n"
|
|
"$LDWend_kernel_info:\n"
|
|
" }\n"
|
|
;
|