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

最新推荐文章于 2024-03-26 09:01:50 发布

xiewen_bupt

最新推荐文章于 2024-03-26 09:01:50 发布

阅读量1.5k

点赞数

分类专栏： GPGPU-CUDA 文章标签： register local memory cuda

本文链接：https://blog.csdn.net/xiewen_bupt/article/details/50387370

版权

本文通过四个CUDA kernel测试，探讨线程分配的数组是在register还是local memory中。测试显示，静态索引时数组存储在寄存器，动态索引且访问一致时在本地内存，动态索引且访问不一致时会导致本地内存重复读取。

摘要由CSDN通过智能技术生成

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

首先一些定义：

#define BLOCK_SIZE 32
#define GRID_SIZE 1
#define ARRAY_SIZE 32

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

__global__ void kernel1(float *buf) {
    float a[ARRAY_SIZE];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    #pragma unroll
    for (int i = 0; i < 5; ++i) {
        a[i] = buf[tid];
    }

    float sum = 0.f;
    #pragma unroll
    for (int i = 0; i < 5; ++i) {
        //static indexing
        sum += a[i];
    }

    buf[tid] = sum;
}

这里，我们注意第二个for循环中，我们使用了pragma unroll展开循环，编译器将会将i优化为对应的数字0,1,2,3,4，同时，编译器也会将a[ARRAY_SIZE]优化为寄存器。我们可以通过nvvp来看一下sass代码，第41