CUDA 中的热方程矩阵 - 非法地址错误

Question

按照官方指南中的this question with reference to the shared memory example，我正在尝试构建热方程矩阵，就像我制作的这张画得不好的图像一样

这是我到目前为止所做的，最小的例子

#define N 32
#define BLOCK_SIZE 16
#define NUM_BLOCKS ((N + BLOCK_SIZE - 1)/ BLOCK_SIZE)

__global__ void heat_matrix(int* A)
{
    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    __shared__ int temp_sm_A[N*N];
    int* temp_A = &temp_sm_A[0]; memset(temp_A, 0, N*N*sizeof(int));

    if (tid < N) //(*)
    {
        #pragma unroll
        for (unsigned int m = 0; m < NUM_BLOCKS; ++m)
        {           
            #pragma unroll
            for (unsigned int e = 0; e < BLOCK_SIZE ; ++e) 
            {
                if ( (tid == 0 && e == 0) || (tid == (N-1) && e == (BLOCK_SIZE-1) ) )
                {
                    temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2;
                    temp_A[tid + (e + BLOCK_SIZE * m) * N + ( tid==0 ? 1 : -1 )] = 1;
                }
                if ( tid == e )
                {
                    temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1;
                    //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N -1));
                    temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2;
                    //printf("temp_A[%d] = -2;\n", (tid + (e + BLOCK_SIZE * m) * N));
                    temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1;
                    //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N +1));
                }
            }
        }
        __syncthreads(); //(**)
        memcpy(A, temp_A, N*N*sizeof(int));
    }
}
int main(){
    int* h_A = (int*)malloc(N*N*sizeof(int)); memset(h_A, 0, N*N*sizeof(int));
    int* d_A; 
    checkCudaErrors(cudaMalloc((void**)&d_A, N*N*sizeof(int)));
    checkCudaErrors(cudaMemcpy(d_A, h_A, N*N*sizeof(int), cudaMemcpyHostToDevice));
    dim3 dim_grid((N/2 + BLOCK_SIZE -1)/ BLOCK_SIZE);
    dim3 dim_block(BLOCK_SIZE);

    heat_matrix <<< dim_grid, dim_block >>> (d_A);
    checkCudaErrors(cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost));
...
}

代码被安排以适应大 N（大于 32）。我利用了块划分。执行时 nvcc 产生

CUDA error at matrix.cu:102 code=77(cudaErrorIllegalAddress) "cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)"

而cuda-memcheck只提供了一个错误（其实还有一个，不过是来自cudasuccess=checkCudaErrors(cudaDeviceReset()); ...）

========= CUDA-MEMCHECK
========= Invalid __shared__ write of size 4
=========     at 0x00000cd0 in heat_matrix(int*)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0xfffffffc is out of bounds
...

我看不出我在代码中哪里做错了。第一个块中的线程 0 怎么会引发非法访问？甚至还有具体的if案例来处理，并没有报出错误所在的代码行。

此外，对于我的代码，是否有比处理所有这些 if 更有效的方法？当然有，但我找不到更好的并行表达式来将案例拆分为第二个 for.

附带说明一下，对我来说 (*) 似乎是不必要的；如果我想跟进其他 GPU 函数调用，则 (**) 是必需的。我说得对吗？

Answer 1

在您的内核中，temp_A 是指向共享内存数组开头的本地指针。考虑到：

N = 32;

BLOCK_SIZE = 16;

m (0,1);

e (0,BLOCK_SIZE)

像 temp_A[tid + (e + BLOCK_SIZE * m) * N] 这样的访问很容易超出 1024 元素长数组的范围。

Answer 2

查看这一行：
```
            temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1;
```
对于在第一次迭代期间 tid 等于零的线程，tid + (e + BLOCK_SIZE * m) * N - 1 计算为索引 -1。这正是 cuda-memcheck 输出所抱怨的（地址由于下溢而回绕）。
稍后将对该行
进行类似的out-of-bounds访问
```
            temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1;
```
当 tid、e 和 m 都取最大值时。
您有多个线程写入同一个内存位置。每个线程应该在每次内循环迭代中写入一个数组元素。不需要写出相邻元素，因为它们已经被自己的线程覆盖了。
初始化 memset() 和主循环内的存储之间存在竞争条件。在 memset().

syncthreads()

对 memset() 和 memcpy() 的调用将导致每个线程执行完整的 set/copy，执行操作 N 次，而不仅仅是一次。
处理这个问题的常用方法是显式写出操作，在块的线程之间分配工作。
然而...
先在共享内存中创建矩阵，然后再将其复制到全局内存没有任何好处。直接写入全局内存中的 A 完全不需要 memset()、memcpy() 和 syncthreads()。
使用仅 16 个线程的块大小会留下一半未使用的资源，因为线程块是以 32 个线程为单位分配的（一个 warp）。

您可能需要 re-read CUDA C 编程指南中有关 Thread Hierarchy 的部分。

CUDA 中的热方程矩阵 - 非法地址错误

Heat equation matrix in CUDA - illegal address error

cuda

matrix

shared-memory

memcpy