在论坛里面讨论到一个问题,__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 < num; 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<3>;
.reg .u32 %r<16>;
.reg .pred %p<4>;
.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
.loc 14 21 0
ld.local.s8 %rh1, [%r10+0]; // id:20 __cuda___cuda_p_HelloCUDA_168+0x0
st.global.s8 [%r13+0], %rh1; // id:21
add.u32 %r13, %r13, 1; //
add.u32 %r10, %r10, 1; //
setp.ne.s32 %p2, %r10, %r12; //
@%p2 bra $Lt_0_7; //
st.shared.s32 [i], %r7; // id:22 i+0x0
bra.uni $Lt_0_5; //
$Lt_0_9:
st.shared.s32 [i], %r6; // id:22 i+0x0
$Lt_0_5:
.loc 14 23 0
exit; //
$LDWend__Z9HelloCUDAPci:
} // _Z9HelloCUDAPci
Cubin code
architecture {sm_10}
abiversion {1}
modname {cubin}
consts {
name = __constant432
segname = const
segnum = 0
offset = 0
bytes = 12
mem {
0x6c6c6548 0x5543206f 0x00214144
}
}
code {
name = _Z9HelloCUDAPci
lmem = 12
smem = 28 // 我们注意这里的smem 的数量
reg = 3
bar = 0
bincode {
0x10000001 0x2400c780 0xd0000001 0x60c00780
0x10000201 0x2400c780 0xd0000801 0x60c00780
0x10000401 0x2400c780 0x307ccbfd 0x6c20c7c8
0xd0001001 0x60c00780 0x10014003 0x00000280
0x1000f801 0x0403c780 0x1000c805 0x0423c780
0x00000005 0xc0000780 0xd4000009 0x40200780
0x20018001 0x00000003 0xd00e0209 0xa0200780
0x3000cbfd 0x6c2147c8 0x20018205 0x00000003
0x1000a003 0x00000280 0x1000ca01 0x0423c780
0x00000c01 0xe4200780 0x30000003 0x00000780
0x00000c01 0xe43f0781
}
}
这个是一段加了shared memory的也有constant 的ptx代码~还有cubin
从cubin来看,确实有可能是通过global memory进来的,不过一定是分发到各自的.param变量里面去的,因为在每一个thread 里面都是可以修改传进来的参数的;
从这点来看:大体应该是 参数( global memory--> (constant memory or shared memory) 然后broadcast to each thread--> 每一个thread都各自的register里面有了参数
大体应该是这样一个过程;
分享到:
相关推荐
cuda 11.1.0 版本的运算平台,本资源免费下载 由于文件太大,所以分割为多个小文件...cat cuda_11.1.0_455.23.05_linux.tar.gz* > cuda_11.1.0_455.23.05_linux.tar.gz tar -zxf cuda_11.1.0_455.23.05_linux.tar.gz
cuda_8.0.44_win10是WIN10 系统下64位使用GPU做caffe开发时安装的cuda8.0版本
cuda_11.1.1_456.81_win10.exe+cudnn-11.1-windows-x64-v8.0.4.30
cuda_10.0.130_410.48_linux.zip,解压后为cuda_10.0.130_410.48_linux.run
英伟达官方2021.6最新CUDA程序 cuda_11.3.1_465.89_win10百度网盘链接 适用于win10 64,亲测可用
cuda_11.4.1_471.41_win10.exe64位安装包,下载链接详情见txt文件
这个是cuda_9.2.148的网络版本安装包。安装时须联网并且按照操作流程进行。
cuda_11.2.0_460.27.04_linux.run cuda_11.2.0_460.27cuda_cuda_11.2.0_460.27.04_linux.run .04_linux.run cuda_11.2.0_460.27.04_linux.run cuda_11.2.0_460.27.04_linux.run cuda_11.2.0_460.27.04_linux.run ...
CUDA_Toolkit_Reference_Manual 齐全的API参考手册 cuda学习者值得拥有
opencv_cuda_opencvcuda_cuda+opencv_opencv_cudaopencv_cuda_源码.zip
opencv_cuda_opencvcuda_cuda+opencv_opencv_cudaopencv_cuda_源码.rar
cuda_10.2.89_441.22_win10.exe64位安装包,下载链接详情见txt文件
CUDA_C_Getting_Started_Windows Nvidia官方CUDA4.0附带文档
cuda10.0补丁包,cuda10.0补丁包,安装方式sudo sh cuda_10.0.130.1_linux.run
CUDA编程资料,包括:CUDA by Example.An Introduction to General-Purpose GPU Programming、CUDA_C_Programming_Guide和cuda_by_example源码
CUDA_Installation_Guide_Linux.pdf+带书签 CUDA_Installation_Guide_Linux.pdf+仅适用于Linux
CUDA_C_Programming_Guide V10.0 最新版的CUDA C编程指南
CUDA_Getting_Started_2.2_Windows.pdf
cuda_11.0.2_451.48_win10.exe
cuda_9.0.176_win10 cuda_9.0.176_win10 cuda_9.0.176_win10