线程与块间多cuRAND序列管理问题及CUDA图像生成异常排查

尝试通过单次CUDA kernel调用生成32张64x64位图，需求如下：

• 部分随机决策针对单张图全局生效（逐图随机化）
• 部分随机决策逐像素独立生成（逐像素随机化）

当前实现采用两组curandState数组：

• 32个元素的逐图随机状态数组
• 4096个元素的逐像素随机状态数组

但出现两类异常：

1. 仅使用逐像素噪声：预期生成32张完全相同的随机噪点图，实际得到高度相关但存在差异的图，且img_id越大的图像差异越明显；
2. 仅使用逐图噪声：预期每张图为纯色块，实际部分图像的四个象限颜色不一致，前期图像一致、后期图像差异较大。

怀疑问题源于每张64x64图像由2x2个32x32线程块组成，而cuRAND文档明确指出不同块无法安全操作同一状态，但未给出对应解决方案。以下是代码片段：

__global__ void init_rngs(curandState* per_img_rng_state, curandState* per_pxl_rng_state) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    int img_id = blockIdx.z * blockDim.z;
    int pxl_id = col * 64 + row;

    curand_init(42, img_id, 0, &per_img_rng_state[img_id]);
    curand_init(42, pxl_id, 0, &per_pxl_rng_state[pxl_id]);
}

__global__ void make_images(curandState* per_img_rng_state, curandState* per_pxl_rng_state, unsigned char* image) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    int img_id = blockIdx.z * blockDim.z;
    int pxl_id = col * 64 + row;

    unsigned int per_img_noise = curand(&per_img_rng_state[img_id]);
    unsigned int per_pxl_noise = curand(&per_pxl_rng_state[pxl_id]);

    // 混合噪声示例
    unsigned int density = per_img_noise;
    unsigned int value = per_img_noise ^ per_pxl_noise;
    image[img_id][row][col] = (value >= density) ? 0x00 : 0xFF;

    // 仅逐像素噪声示例
    image[img_id][row][col] = (per_pxl_noise & 1) ? 0x00 : 0xFF;

    // 仅逐图噪声示例
    image[img_id][row][col] = per_img_noise / 16777216;
}

void randomize_images() {
  curandState* per_img_rng_state = nullptr;
  curandState* per_pxl_rng_state = nullptr;
  unsigned char* image = nullptr;

  cudaMalloc(&image, 32*64*64);
  cudaMalloc(&per_img_rng_state, 32 * sizeof(curandState));
  cudaMalloc(&per_pxl_rng_state, 64 * 64 * sizeof(curandState));

  // 2x2x32块网格，每个块32x32线程
  init_rngs<<<{2, 2, 32}, {32, 32}>>>(per_img_rng_state, per_pxl_rng_state);
  make_images<<<{2, 2, 32}, {32, 32}>>>(per_img_rng_state, per_pxl_rng_state, image);
}

1. 逐图状态的多线程竞态：
   每张图对应1个curandState，但生成单张图的4个线程块（2x2）中的所有线程都会读写这个状态。curand()是有状态操作，会修改状态值，多线程（跨块）同时操作同一状态会导致竞态——状态被无序修改，不同线程读取到的状态值混乱，最终同图不同象限的随机值不一致。

2. 逐像素状态的重复初始化：
   初始化per_pxl_rng_state时，每个img_id对应的线程都会初始化同一个pxl_id的状态。由于线程执行顺序不确定，多次初始化同一状态会导致最终状态不可预测，不同img_id的同位置像素使用的随机状态已被多次修改，生成的随机值自然有差异。

3. 内存索引错误：
   CUDA中不支持直接使用image[img_id][row][col]这类多维数组语法，会导致内存访问偏移计算错误，加剧异常表现。

方案1：修正状态管理逻辑 #

核心思路 #

• 逐图状态：仅由单一线程初始化和读取，通过共享内存广播给同图所有线程，避免竞态；
• 逐像素状态：每个像素的状态仅由一个线程初始化，避免重复修改；
• 改用线性内存索引，避免多维数组语法错误。

修正后代码 #

__global__ void init_per_img_rngs(curandState* per_img_rng_state) {
    int img_id = blockIdx.x;
    // 每个img_id仅由一个线程初始化状态
    if (threadIdx.x == 0) {
        curand_init(42, img_id, 0, &per_img_rng_state[img_id]);
    }
}

__global__ void init_per_pxl_rngs(curandState* per_pxl_rng_state) {
    int pxl_id = blockIdx.x;
    // 每个pxl_id仅由一个线程初始化状态
    if (threadIdx.x == 0) {
        curand_init(42, pxl_id, 0, &per_pxl_rng_state[pxl_id]);
    }
}

__global__ void make_images(curandState* per_img_rng_state, curandState* per_pxl_rng_state, unsigned char* image) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    int img_id = blockIdx.z;
    // 线性内存索引：避免多维数组语法错误
    int pxl_idx = img_id * 64 * 64 + row * 64 + col;
    int pxl_global_id = row * 64 + col;

    __shared__ unsigned int per_img_val;
    // 仅当前块的(0,0)线程读取逐图随机值并广播
    if (threadIdx.x == 0 && threadIdx.y == 0) {
        per_img_val = curand(&per_img_rng_state[img_id]);
    }
    __syncthreads(); // 等待共享内存广播完成

    // 逐像素状态仅由当前像素线程访问
    unsigned int per_pxl_val = curand(&per_pxl_rng_state[pxl_global_id]);

    // 仅逐像素噪声：所有图同位置像素生成相同值
    image[pxl_idx] = (per_pxl_val & 1) ? 0x00 : 0xFF;

    // 仅逐图噪声：整张图使用同一值，生成纯色块
    // image[pxl_idx] = per_img_val / 16777216;
}

void randomize_images() {
    curandState* per_img_rng_state = nullptr;
    curandState* per_pxl_rng_state = nullptr;
    unsigned char* image = nullptr;

    cudaMalloc(&image, 32*64*64);
    cudaMalloc(&per_img_rng_state, 32 * sizeof(curandState));
    cudaMalloc(&per_pxl_rng_state, 64 * 64 * sizeof(curandState));

    // 单独初始化逐图状态：32个块，每个块1个线程
    init_per_img_rngs<<<32, 1>>>(per_img_rng_state);
    // 单独初始化逐像素状态：4096个块，每个块1个线程
    init_per_pxl_rngs<<<64*64, 1>>>(per_pxl_rng_state);

    // 生成图像：保持原2x2x32块网格，32x32线程块
    make_images<<<dim3(2,2,32), dim3(32,32)>>>(per_img_rng_state, per_pxl_rng_state, image);
}

方案2：无状态随机数生成（推荐） #

如果不需要复杂的随机序列特性，可使用cuRAND无状态随机函数或自定义哈希，彻底避免状态管理问题，代码更简洁，性能更优。

示例代码 #

// 无状态逐图随机值：基于img_id生成固定值
__device__ unsigned int get_per_img_noise(int img_id) {
    curandStateXORWOW temp_state;
    curand_init(42, img_id, 0, &temp_state);
    return curand(&temp_state);
}

// 无状态逐像素随机值：基于全局像素ID生成固定值
__device__ unsigned int get_per_pxl_noise(int pxl_global_id) {
    curandStateXORWOW temp_state;
    curand_init(42, pxl_global_id, 0, &temp_state);
    return curand(&temp_state);
}

__global__ void make_images(unsigned char* image) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    int img_id = blockIdx.z;
    int pxl_idx = img_id * 64 * 64 + row * 64 + col;
    int pxl_global_id = row * 64 + col;

    unsigned int per_img_val = get_per_img_noise(img_id);
    unsigned int per_pxl_val = get_per_pxl_noise(pxl_global_id);

    // 仅逐像素噪声：所有图同位置像素相同
    image[pxl_idx] = (per_pxl_val & 1) ? 0x00 : 0xFF;

    // 仅逐图噪声：整张图纯色
    // image[pxl_idx] = per_img_val / 16777216;
}

void randomize_images() {
    unsigned char* image = nullptr;
    cudaMalloc(&image, 32*64*64);
    make_images<<<dim3(2,2,32), dim3(32,32)>>>(image);
}

内容的提问来源于stack exchange，提问作者Nate Gaylinn