CUDA __global__ function 参数分析

在论坛里面讨论到一个问题，__global__函数里面传递的参数，到底是怎么传输到每一个thread的，然后做了以下的一些分析；

这个是问题讨论帖子：http://topic.csdn.net/u/20090210/22/2d9ac353-9606-4fa3-9dee-9d41d7fb2b40.html

C/C++ code

__global__ static void HelloCUDA(char* result, int num)
{
__shared__ int i;
i = 0;
char p_HelloCUDA[] = "Hello CUDA!";
for(i = 0; i result[i] = p_HelloCUDA[i];
}
}

PTX code

.const .align 1 .b8 __constant432[12] = {0x48,0x65,0x6c,0x6c,0x6f,0x20,0x43,0x55,0x44,0x41,0x21,0x0};

.entry _Z9HelloCUDAPci
{
.reg .u16 %rh;
.reg .u32 %r;
.reg .pred %p;
.param .u32 __cudaparm__Z9HelloCUDAPci_result;
.param .s32 __cudaparm__Z9HelloCUDAPci_num;
.local .align 4 .b8 __cuda___cuda_p_HelloCUDA_168[12];
.shared .s32 i;
.loc 14 15 0
$LBB1__Z9HelloCUDAPci:
mov.u32 %r1, __constant432; //
mov.u32 %r2, __cuda___cuda_p_HelloCUDA_168; //
ld.const.u32 %r3, [%r1+0]; // id:17 not_variable+0x0
st.local.u32 [%r2+0], %r3; // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
ld.const.u32 %r4, [%r1+4]; // id:17 not_variable+0x0
st.local.u32 [%r2+4], %r4; // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
ld.const.u32 %r5, [%r1+8]; // id:17 not_variable+0x0
st.local.u32 [%r2+8], %r5; // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
.loc 14 20 0
mov.s32 %r6, 0; //
ld.param.s32 %r7, [__cudaparm__Z9HelloCUDAPci_num]; // id:16 __cudaparm__Z9HelloCUDAPci_num+0x0
mov.u32 %r8, 0; //
setp.le.s32 %p1, %r7, %r8; //
@%p1 bra $Lt_0_9; //
mov.s32 %r9, %r7; //
mov.u32 %r10, __cuda___cuda_p_HelloCUDA_168; //
mov.u32 %r11, __cuda___cuda_p_HelloCUDA_168; //
add.u32 %r12, %r7, %r11; //
ld.param.u32 %r13, [__cudaparm__Z9HelloCUDAPci_result]; // id:19 __cudaparm__Z9HelloCUDAPci_result+0x0
mov.s32 %r14, %r9; //
$Lt_0_7:
//<loop> Loop body line 20, nesting depth: 1, estimated iterations: unknown<br> .loc 14 21 0<br> ld.local.s8 %rh1, [%r10+0]; // id:20 __cuda___cuda_p_HelloCUDA_168+0x0<br> st.global.s8 [%r13+0], %rh1; // id:21<br> add.u32 %r13, %r13, 1; // <br> add.u32 %r10, %r10, 1; // <br> setp.ne.s32 %p2, %r10, %r12; // <br> @%p2 bra $Lt_0_7; // <br> st.shared.s32 [i], %r7; // id:22 i+0x0<br> bra.uni $Lt_0_5; // <br>$Lt_0_9:<br> st.shared.s32 [i], %r6; // id:22 i+0x0<br>$Lt_0_5:<br> .loc 14 23 0<br> exit; // <br>$LDWend__Z9HelloCUDAPci:<br> } // _Z9HelloCUDAPci<br><dd> <pre></pre> </dd> <dt>Cubin code <dt> <br>architecture {sm_10}<br>abiversion {1}<br>modname {cubin}<br>consts {<br> name = __constant432<br> segname = const<br> segnum = 0<br> offset = 0<br> bytes = 12<br> mem {<br> 0x6c6c6548 0x5543206f 0x00214144 <br> }<br>}<br>code {<br> name = _Z9HelloCUDAPci<br> lmem = 12<br> smem = 28 <u>// 我们注意这里的smem 的数量<br></u> reg = 3<br> bar = 0<br> bincode {<br> 0x10000001 0x2400c780 0xd0000001 0x60c00780 <br> 0x10000201 0x2400c780 0xd0000801 0x60c00780 <br> 0x10000401 0x2400c780 0x307ccbfd 0x6c20c7c8 <br> 0xd0001001 0x60c00780 0x10014003 0x00000280 <br> 0x1000f801 0x0403c780 0x1000c805 0x0423c780 <br> 0x00000005 0xc0000780 0xd4000009 0x40200780 <br> 0x20018001 0x00000003 0xd00e0209 0xa0200780 <br> 0x3000cbfd 0x6c2147c8 0x20018205 0x00000003 <br> 0x1000a003 0x00000280 0x1000ca01 0x0423c780 <br> 0x00000c01 0xe4200780 0x30000003 0x00000780 <br> 0x00000c01 0xe43f0781 <br> }<br>}<br> </dt> </dt> <dd><pre></pre></dd></loop>

这个是一段加了shared memory的也有constant 的ptx代码~还有cubin

从cubin来看，确实有可能是通过global memory进来的，不过一定是分发到各自的.param变量里面去的，因为在每一个thread 里面都是可以修改传进来的参数的；
从这点来看：大体应该是 参数( global memory--> (constant memory or shared memory) 然后broadcast to each thread--> 每一个thread都各自的register里面有了参数
大体应该是这样一个过程；