非确定性 CUDA C 内核答案

【问题标题】：Non-deterministic CUDA C kernel非确定性 CUDA C 内核
【发布时间】：2017-07-13 06:03:04
【问题描述】：

我仍然是 CUDA 的初学者，我一直在尝试编写一个简单的内核来在 GPU 上执行并行素筛。最初我用 C 编写了我的代码，但我想研究 GPU 上的加速，所以我重写了它：

41.cu

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

#define B 1024
#define T 256
#define N (B*T)

#define checkCudaErrors(error) {\
    if (error != cudaSuccess) {\
        printf("CUDA Error - %s:%d: '%s'\n",__FILE__,__LINE__,cudaGetErrorString(error));\
        exit(1);\
        }\
}\

__global__ void prime_sieve(int *primes) {
    unsigned int i = threadIdx.x + blockIdx.x * blockDim.x;
    primes[i] = i;
    primes[0] = primes[1] = 0;
    if (i > 1 && i<N) { 
        for (int j=2; j<N/2; j++) {
            if (i*j < N) {
                primes[i*j] = 0;
            }
        }
    }
}

int main() {
    int *h_primes=(int*)malloc(N * sizeof(int));
    int *d_primes;
    checkCudaErrors(cudaMalloc( (void**)&d_primes, N*sizeof(int)));
    checkCudaErrors(cudaMemcpy(d_primes,h_primes,N*sizeof(int),cudaMemcpyHostToDevice));    
    prime_sieve<<<B,T>>>(d_primes);
    checkCudaErrors(cudaMemcpy(h_primes,d_primes,N*sizeof(int),cudaMemcpyDeviceToHost));
    checkCudaErrors(cudaFree(d_primes));

    int size = 0;
    int total = 0;
    for (int i=2; i<N; i++) {
        if (h_primes[i]) {
            size++;
        }
        total++;
    }
    printf("\n");
    printf("Length = %d\tPrimes = %d\n",total,size);
    free(h_primes);
    return 0;
}

我在 Ubuntu 16.04 (4.4.0-83-generic) 上运行该程序，并在 8.0.61 版本下使用 nvcc 41.cu -o 41.o -arch=sm_30 进行编译。该程序在 GeForce GTX 780 Ti 上运行，但每次运行时，总是会产生不确定的结果：

Length = 262142 Primes = 49477
Length = 262142 Primes = 49486
Length = 262142 Primes = 49596
Length = 262142 Primes = 49589

没有错误报告。起初我认为这是一个竞争条件，但cuda-memcheck 没有报告racecheck、initcheck 或synccheck 的任何危险，我想不出我的假设有任何问题。我在想这可能是一个同步问题？

这种非确定性行为仅在我增加代码中看到的块大小和线程大小时才会发生。当我尝试块大小和线程大小为 16 时，没有问题（据我所知）。似乎并非所有线程都有机会执行？我计划在非常大的数组大小（

我在这里做错了什么？

【问题讨论】：

标签： c cuda

【解决方案1】：

有一个巨大的竞争条件

所以prime[i] > 0 表示素数，而prime[i]=0 表示复合。

primes[i] = i; 作为每个线程对primes 的第一次更新执行。请记住这一点。

现在让我们看看当线程16 执行时会发生什么。它也标记primes[16]=16 和16 的所有倍数。类似于以下内容

primes[16] = primes[32] = primes[48]=....=primes[k*16]=0

想象一下，线程 48 在线程 16 完成其工作之后（或者当线程 16 循环中的 j>3 时）被调度。

线程48 设置primes[48] = 48。你丢失了线程16所做的更新。

这是一个竞争条件。

在 CUDA 中编码时，您应确保代码的正确性不依赖于特定的 warp 调度。

您应该将执行顺序视为非确定性的。

【讨论】：