cuda - 在 CUDA 中读取带有线程的数组

Question

我想知道这是否可能，以及从具有 CUDA 线程的数组中读取单元格的最佳方法是什么。为了简化我的意思，这是一个例子：

我有一个数组：{1,2,3,4,5,6,...}，我希望每个线程读取我的数组的 n 个单元，主要取决于它的大小。

我一直在尝试一些事情，但它似乎不起作用，所以如果有人能指出一种（正确的）方法来做到这一点，那就太好了。

谢谢你。

Answer 1

通常，您希望连续线程读取连续数组索引。这样做会导致“合并”内存事务。简单的思考方式是，如果 32 个线程在物理上并行运行，并且它们都做一个加载，那么如果所有 32 个加载都落入同一个缓存行，那么可以执行一次内存访问来填充缓存行，而不是 32 个单独的。

所以你想要做的是让每个线程访问n按线程数跨步的单元格，就像这样（假设输入数据在float数组中data）。

int idx = blockDim.x * blockIdx.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = idx; i < numElements; i += stride) {
  float element = data[i];
  process(element);
}

如果您的算法要求每个线程读取n连续的数据元素，那么您将产生非合并负载，这将更加昂贵。在这种情况下，我会考虑重新设计算法，这样就不需要这种类型的访问。

Answer 2

你需要：

线程必须查看下 n 个数字

所以你可以使用：

#define N 2
#define NTHREAD 1024
#define ARRAYSIZE N*NTHREAD

// develop the kernel as:
__global__ void accessArray(int *array){
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    int startId = tid*N;

    // access thread's stride
    for(int i=0; i<N; i++){
        array[startId+i]=tid;
    }
}
// call the kernel by:
accessArray<<<NTHREAD/256, 256>>>(d_array);

转储array并检查它是否是您希望线程工作的方式。

完整代码：

#include <cuda.h>
#include <stdio.h>


#define N 2
#define NTHREAD 1024
#define ARRAYSIZE N*NTHREAD

// develop the kernel as:
__global__ void accessArray(int *array){
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    int startId = tid*N;

    // access thread's stride
    for(int i=0; i<N; i++){
        array[startId+i]=tid;
    }
}

int  main()
{
    int h_array[ARRAYSIZE];
    int *d_array;
    size_t memsize= ARRAYSIZE * sizeof(float);

    for (int i=0; i< ARRAYSIZE; i++) {
        h_array[i] = 0;
    }

    cudaMalloc(&d_array, memsize);
    cudaMemcpy(d_array, h_array, memsize,  cudaMemcpyHostToDevice);

    accessArray<<<NTHREAD/256, 256>>>(d_array);
    cudaMemcpy(h_array, d_array, memsize,  cudaMemcpyDeviceToHost);

    for (int i=0; i<ARRAYSIZE; i++)
        printf("A[%d] => %d\n",i,h_array[i]);

    cudaFree(d_array);
}