CUDA内核访问二维数组实现康威生命游戏邻域遍历问题求助

核心问题梳理 #

• 内核启动参数顺序错误：CUDA核函数启动配置为<<<网格维度, 线程块维度>>>，你的代码中写反了两个参数，导致所有线程的x/y坐标完全不符合预期。
• Pitch分配的内存未按规范访问：你用cudaMallocPitch分配了对齐内存，但是内核中完全没有使用传入的pitch参数，还是用普通的连续内存索引逻辑，访问位置全部错位。
• 坐标与索引对应逻辑混乱：没有统一行/列与x/y的对应关系，导致邻域偏移时相对位置完全错误。

具体修复步骤 #

1. 统一二维数组访问规范 #

首先约定：y为行索引（垂直方向，范围0_{ROWS-1），`x`为列索引（水平方向，范围0}COLUMNS-1），pitch访问统一用如下工具函数：

__device__ int* get_board_element(int* base, size_t pitch, int y, int x) {
    // pitch是每行的字节数，先转为char*计算行偏移，再转int*计算列偏移
    return (int*)((char*)base + y * pitch) + x;
}

2. 修正核函数启动参数 #

把核函数启动的参数顺序调换，同时修正线程块维度的定义：

dim3 block(BLOCKSIZE_X, BLOCKSIZE_Y);
dim3 grid(divideAndRound(COLUMNS, BLOCKSIZE_X), divideAndRound(ROWS, BLOCKSIZE_Y));
// 原来的错误写法是把block和grid顺序写反
numAliveAround <<<grid, block>>> (d_oldBoard, d_newBoard, COLUMNS, ROWS, pitchOld, pitchNew);
cudaDeviceSynchronize();
determineNextState <<<grid, block>>> (d_oldBoard, d_newBoard, COLUMNS, ROWS, pitchOld, pitchNew);
cudaDeviceSynchronize();

3. 重构邻域计数核函数 #

__global__ void numAliveAround(int *oldBoard, int *newBoard, int xSize, int ySize, size_t pitchOld, size_t pitchNew)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if(x >= xSize || y >= ySize) return;

    int alive_count = 0;
    // 遍历8个邻域方向
    for(int dy = -1; dy <= 1; dy++) {
        for(int dx = -1; dx <= 1; dx++) {
            if(dx == 0 && dy == 0) continue; // 跳过当前细胞本身
            // 边界wrap处理
            int nx = (x + dx + xSize) % xSize;
            int ny = (y + dy + ySize) % ySize;
            alive_count += *get_board_element(oldBoard, pitchOld, ny, nx);
        }
    }
    // 写入邻域存活数到newBoard
    *get_board_element(newBoard, pitchNew, y, x) = alive_count;
}

4. 修正状态更新核函数 #

__global__ void determineNextState(int *oldBoard, int *newBoard, int xSize, int ySize, size_t pitchOld, size_t pitchNew)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if(x >= xSize || y >= ySize) return;

    int cur_state = *get_board_element(oldBoard, pitchOld, y, x);
    int alive_nei = *get_board_element(newBoard, pitchNew, y, x);
    int next_state = 0;

    if(cur_state == 1) {
        next_state = (alive_nei == 2 || alive_nei == 3) ? 1 : 0;
    } else {
        next_state = (alive_nei == 3) ? 1 : 0;
    }

    *get_board_element(newBoard, pitchNew, y, x) = next_state;
}

为什么简单操作看起来正常？ #

类似给所有元素加1的操作，只需要每个元素被访问一次即可，哪怕索引错位、线程坐标混乱，只要每个内存地址刚好被一个线程处理，结果就不会出错。但邻域遍历依赖坐标的相对位置，只要索引逻辑有一点问题，相对位置就会完全错乱，所以会直接失效。

内容的提问来源于stack exchange，提问作者Pavel Smirnov