CUDA之编程中线程分配的数组在register中还是local memory中？

问题很简单，当我们在编写KERNEL的时候，分配了一个数组，那么这段数组空间是在register中，还是local memory中呢？通过几个测试，我们可以来看一下：

首先一些定义：

[cpp]  view plain  copy

1. #define BLOCK_SIZE 32  
2. #define GRID_SIZE 1  
3. #define ARRAY_SIZE 32  

第一个kernel测试，静态索引static indexing，代码：

[cpp]  view plain  copy

1. __global__ void kernel1(float *buf) {  
2.     float a[ARRAY_SIZE];  
3.     int tid = threadIdx.x + blockIdx.x * blockDim.x;  
4.   
5.     #pragma unroll  
6.     for (int i = 0; i < 5; ++i) {  
7.         a[i] = buf[tid];  
8.     }  
9.   
10.     float sum = 0.f;  
11.     #pragma unroll  
12.     for (int i = 0; i < 5; ++i) {  
13.         //static indexing  
14.         sum += a[i];  
15.     }  
16.   
17.     buf[tid] = sum;  
18. }  

这里，我们注意第二个for循环中，我们使用了pragma unroll展开循环，编译器将会将i优化为对应的数字0,1,2,3,4，同时，编译器也会将a[ARRAY_SIZE]优化为寄存器。我们可以通过nvvp来看一下sass代码，第41行代码对应的SASS代码只包含了5次FADD，并没有从local memory读数据的指令LDL，因此数组a是在寄存器中存储的。

第二个kernel测试，动态索引dynamic indexing，uniform access，代码：

[cpp]  view plain  copy

1. __global__ void kernel2(float *buf, int idx_begin) {  
2.     float a[ARRAY_SIZE];  
3.     int tid = threadIdx.x + blockIdx.x * blockDim.x;  
4.   
5.     #pragma unroll  
6.     for (int i = 0; i < 5; ++i) {  
7.         a[i] = buf[tid];  
8.     }  
9.   
10.     float sum = 0.f;  
11.     #pragma unroll  
12.     for (int i = 0; i < 5; ++i) {  
13.         //dynamic indexing with uniform access  
14.         //The data in a will be stored in local memory, the access of array a will results a local load instruct per warp(LDL), and replays 0 times  
15.         sum += a[i+idx_begin];  
16.     }  
17.   
18.     buf[tid] = sum;  
19. }  

这个kernel中，数组的索引不是常量，因此编译器将不会将数组a存储到寄存器中，而是分配到了local memory中了。同样，我们可以看一下SASS代码，在FADD之前，都会利用LDL(load local memory)指令从local memory中将数据读到寄存器中。这里我们要说明一下，这里我们只分配了一个block，里面包含了32个线程。因此在GPU的SM跑的时候只有一个warp。32个线程每次读取数据时的index是一致的，因此load指令replay只有一次，这中读取方式成为uniform access。在最后，我们会使用nvprof来看看到底replay了多少次。

第三个kernel测试，动态索引dynamic indexing，non-uniform access，代码：

[cpp]  view plain  copy

1. __global__ void kernel3(float *buf, int *idxBuf) {  
2.     float a[ARRAY_SIZE+5];  
3.     int tid = threadIdx.x + blockIdx.x * blockDim.x;  
4.   
5.     #pragma unroll  
6.     for (int i = 0; i < 5; ++i) {  
7.         a[i] = buf[tid];  
8.     }  
9.   
10.     float sum = 0.f;  
11.   
12.     #pragma unroll  
13.     for (int i = 0; i < 5; ++i) {  
14.         //dynamic indexing with non-uniform access  
15.         //The data in a will be stored in local memory, the access of a[] will results a local load instructs per warp(LDL), and replays 31 times  
16.         sum += a[i+threadIdx.x];  
17.     }  
18.   
19.     buf[tid] = sum;  
20. }  

这个kernel中，同样是动态的索引，因此数组a是在Local memory中分配的。我们可以看一下SASS代码。这里，同样要说明一下，这32个线程的索引不一致，因此从local memory中读取的时候数组不在一个cache line中，因此读操作会replay很多次。这里，32个索引各不相同，因此将会replay31次。

第四个测试，nvprof events：

为了验证我们刚才的推测（包括存储位置，读操作replay次数），这里使用nvprof来查看一些重要的指标：

[cpp]  view plain  copy

1. wenx@LINUX-14-04:~/work/$ nvprof --events local_load,local_store,__local_ld_mem_divergence_replays,__local_st_mem_divergence_replays ./a.out   
2.   
3. ==7893== NVPROF is profiling process 7893, command: ./a.out  
4. ==7893== Profiling application: ./a.out  
5. ==7893== Profiling result:  
6. ==7893== Event result:  
7. Invocations                                Event Name         Min         Max         Avg  
8. Device "Tesla K20c (0)"  
9.     Kernel: kernel3(float*, int*)  
10.           1                                local_load           5           5           5  
11.           1                               local_store           5           5           5  
12.           1         __local_ld_mem_divergence_replays         155         155         155  
13.           1         __local_st_mem_divergence_replays           0           0           0  
14.     Kernel: kernel2(float*, int)  
15.           1                                local_load           5           5           5  
16.           1                               local_store           5           5           5  
17.           1         __local_ld_mem_divergence_replays           0           0           0  
18.           1         __local_st_mem_divergence_replays           0           0           0  
19.     Kernel: kernel1(float*)  
20.           1                                local_load           0           0           0  
21.           1                               local_store           0           0           0  
22.           1         __local_ld_mem_divergence_replays           0           0           0  
23.           1         __local_st_mem_divergence_replays           0           0           0  

测试结果表明：

1. kernel1中，数组的索引是常量，存储到了寄存器中。
2. kernel2中，数组的索引是动态的，存储到本地内存中。由于访问的index一致，因此replay为0。
3. kernel3中，数组的索引是动态的，存储到本地内存中。由于访问的index都不一致（32个不同的index），因此local memory的读操作重复了31*5=155次，其中5是指循环了5次。

转载：http://blog.csdn.net/xiewen_bupt/article/details/50387370