这是CUDA中的错误吗？（遇到非法内存访问）答案

【问题标题】：Is this a bug in CUDA? (illegal memory access was encountered)这是CUDA中的错误吗？（遇到非法内存访问）
【发布时间】：2017-01-29 17:59:25
【问题描述】：

我正在使用以下 CUDA 内核：

__global__
void sum_worker(int *data, int *sum_ptr)
{
        __shared__ int block_sum;
        int idx = threadIdx.x;
        int thread_sum = 0;

        if (threadIdx.x == 0)
                block_sum = 2;

        for (int i = idx; i < MAX_INDEX; i += blockDim.x)
                thread_sum += data[i];

        __syncthreads();

        atomicAdd(&block_sum, thread_sum);

        __syncthreads();

        if (threadIdx.x == 0)
                *sum_ptr = block_sum;
}

使用以下代码启动它：

sum_worker<<<1, 32>>>(primes_or_zeros, sum_buffer);

它工作正常（没有运行时错误并产生正确的结果）。但是，如果我将i += blockDim.x 更改为i += 32，我会在下次调用cudaDeviceSynchronize() 时出错：

Cuda error 'an illegal memory access was encountered' in primes_gpu.cu at line 97

使用cuda-memcheck运行内核：

========= Invalid __global__ read of size 4
=========     at 0x00000108 in /home/clifford/Work/handicraft/2016/perfmeas/primes_gpu.cu:35:sum_worker(int*, int*)
=========     by thread (31,0,0) in block (0,0,0)
=========     Address 0x703b70d7c is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x472225]
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcudart.so.7.5 [0x146ad]
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcudart.so.7.5 (cudaLaunch + 0x143) [0x2ece3]
=========     Host Frame:./perfmeas [0x17c7]
=========     Host Frame:./perfmeas [0x16b7]
=========     Host Frame:./perfmeas [0x16e2]
=========     Host Frame:./perfmeas [0x153f]
=========     Host Frame:./perfmeas [0xdcd]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf0) [0x20830]
=========     Host Frame:./perfmeas [0xf39]
....

地址 0x703b70d7c 确实超出了data 的范围：数组从 0x703b40000 开始，并且有 MAX_INDEX 个元素。在这个测试中，MAX_INDEX 是 50000。 (0x703b70d7c - 0x703b40000) / 4 = 50015。

为i >= 50000 添加额外的检查使问题神奇地消失：

    for (int i = idx; i < MAX_INDEX; i += 32) {
            if (i >= MAX_INDEX)
                    printf("WTF!\n");
            thread_sum += data[i];
    }

这是 CUDA 中的错误还是我在这里做了一些愚蠢的事情？

我在 Ubuntu 2016.04 上使用 CUDA 7.5。 nvcc --version的输出：

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2015 NVIDIA Corporation
Built on Tue_Aug_11_14:27:32_CDT_2015
Cuda compilation tools, release 7.5, V7.5.17

此测试用例的完整源代码可在此处找到：
http://svn.clifford.at/handicraft/2016/perfmeas

（使用选项-gx 运行。此版本使用i += blockDim.x。将其更改为i += 32 以重现该问题。）

编辑：@njuffa 在 cmets 中说他不想跟踪堆栈溢出的链接，因为他“太害怕 [他的] 计算机可能会捕捉到某些东西”，并且更喜欢他可以直接从堆栈溢出复制和粘贴的测试用例.就这样吧：

#include <string.h>
#include <stdio.h>
#include <stdbool.h>
#include <math.h>

#define MAX_PRIMES 100000
#define MAX_INDEX (MAX_PRIMES/2)

__global__
void primes_worker(int *data)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    if (idx >= MAX_INDEX)
        return;

    int p = 2*idx+1;
    for (int i = 3; i*i <= p; i += 2) {
        if (p % i == 0) {
            data[idx] = 0;
            return;
        }
    }

    data[idx] = idx ? p : 0;
}

__global__
void sum_worker(int *data, int *sum_ptr)
{
    __shared__ int block_sum;
    int idx = threadIdx.x;
    int thread_sum = 0;

    if (threadIdx.x == 0)
        block_sum = 2;

#ifdef ENABLE_BUG
    for (int i = idx; i < MAX_INDEX; i += 32)
        thread_sum += data[i];
#else
    for (int i = idx; i < MAX_INDEX; i += blockDim.x)
        thread_sum += data[i];
#endif

    __syncthreads();

    atomicAdd(&block_sum, thread_sum);

    __syncthreads();

    if (threadIdx.x == 0)
        *sum_ptr = block_sum;
}

int *primes_or_zeros;
int *sum_buffer;

void primes_gpu_init()
{
    cudaError_t err;

    err = cudaMalloc((void**)&primes_or_zeros, sizeof(int)*MAX_INDEX);

    if (err != cudaSuccess)
        printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);

    err = cudaMallocHost((void**)&sum_buffer, sizeof(int));

    if (err != cudaSuccess)
        printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);
}

void primes_gpu_done()
{
    cudaError_t err;

    err = cudaFree(primes_or_zeros);

    if (err != cudaSuccess)
        printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);

    err = cudaFreeHost(sum_buffer);

    if (err != cudaSuccess)
        printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);
}

int primes_gpu()
{
    int num_blocks = (MAX_INDEX + 31) / 32;
    int num_treads = 32;

    primes_worker<<<num_blocks, num_treads>>>(primes_or_zeros);
    sum_worker<<<1, 32>>>(primes_or_zeros, sum_buffer);
    cudaError_t err = cudaDeviceSynchronize();

    if (err != cudaSuccess)
        printf("Cuda error '%s' in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);

    return *sum_buffer;
}

int main()
{
    primes_gpu_init();

    int result = primes_gpu();
    printf("Result: %d\n", result);

    if (result != 454396537) {
        printf("Incorrect result!\n");
        return 1;
    }

    primes_gpu_done();
    return 0;
}

用法：

$ nvcc -o demo demo.cu 
$ ./demo 
Result: 454396537

$ nvcc -D ENABLE_BUG -o demo demo.cu 
$ ./demo 
Cuda error 'an illegal memory access was encountered' in demo.cu at line 99
Result: 0
Incorrect result!

【问题讨论】：

@njuffa 查看编辑。但是，如果您“太害怕 [您的] 计算机可能会捕获某些东西”，那么您可能首先不应该在 SO 上执行来自陌生人的随机代码..
在删除 #include <stdbool> 后，我能够在 Windows 上使用 CUDA 7.5 重现您的观察结果。仍在试图弄清楚发生了什么。作为一个快速实验，尝试通过-Xptxas -O2、-Xptxas -O1、然后-Xptxas -O0 降低ptxas 优化级别。从 SO 编译代码存在残余风险，但至少可以预先（在运行之前）检查它是否有任何可疑之处。
当我使用-Xptxas -O1 时，问题就消失了，因此这可能暗示后端代码生成问题。你在我们这边看到同样的情况吗？我还没能找到机器代码中的问题，因为我设法把自己搞糊涂了，走上了死胡同。
使用最新的可用 CUDA 驱动程序，并强制 JIT 编译（通过使用 -arch=sm_30 编译，但在 sm_50 设备上运行）也会使问题消失，再次暗示 PTXAS 优化存在问题（驱动程序中的 PTXAS 组件比 CUDA 7.5 工具链的 PTXAS 组件更新）。这表明无论确切的问题是什么，它可能已经在 CUDA 8.0 RC 中修复（不确定这是否是您尝试的现实选择）。
问题似乎与循环展开有关。随着循环增加一个变量（threadIdx.x），循环保持滚动。当它是编译时间常数 (32) 时，它会以 4 倍展开。展开的循环具有越界访问。有趣的是，当我通过将#pragma unroll 4 直接放在循环之前显式展开循环时，我得到几乎相同的机器代码，但是它可以正常工作！区别可能在于编译器的哪一部分执行展开、前端或后端。所以这看起来像一个带有展开循环的ptxas 错误。

标签： cuda

【解决方案1】：

TL;DR：观察到的行为很可能是由 CUDA 7.5 工具链的 ptxas 组件中的错误引起的，特别是循环展开器。该错误可能已在公开可用的 CUDA 8.0 RC 中修复。

我能够在带有 Quadro K2200 GPU（sm_50 设备）的 64 位 Windows 7 平台上重现问题中报告的行为。定义了ENABLE_BUG 的生成机器代码（SASS）的主要区别在于循环展开了四倍。这是循环增量从变量 threadIdx.x 更改为编译时常量 32 的直接结果，它允许编译器在编译时计算行程计数。

有趣的是，在中间 PTX 级别，即使增量为32，循环也会滚动：

BB7_4:
ld.global.u32 %r12, [%rd10];
add.s32 %r16, %r12, %r16;
add.s64 %rd10, %rd10, 128;
add.s32 %r15, %r15, 32;
setp.lt.s32     %p3, %r15, 50000;
@%p3 bra BB7_4;

由于循环在机器代码中展开，它必须是应用该转换的 ptxas 展开器。

如果我通过在nvcc 命令行上指定-Xptxas -O1 将ptxas 优化级别降低到-O1，代码将按预期工作。如果我为sm_30 构建代码（在sm_50 设备上运行时导致JIT 编译），则代码在使用最新驱动程序Windows 369.26 运行时按预期工作。这强烈表明在 CUDA 7.5 的 ptxas 组件的展开器中存在一个错误，但该错误已得到修复，因为 CUDA 驱动程序中的 ptxas 组件比 ptxas 组件更新CUDA 7.5 工具链。

将#pragma unroll 4直接放在循环前面也可以解决问题，因为在这种情况下，展开是由编译器的nvvm组件执行的，这意味着展开的循环已经存在于PTX级别：

#if ENABLE_BUG
#pragma unroll 4
    for (int i = idx; i < MAX_INDEX; i += 32)
        thread_sum += data[i];
#else

生成的 PTX：

BB7_5:
.pragma "nounroll";
ld.global.u32 %r34, [%rd14];
add.s32 %r35, %r34, %r45;
ld.global.u32 %r36, [%rd14+128];
add.s32 %r37, %r36, %r35;
ld.global.u32 %r38, [%rd14+256];
add.s32 %r39, %r38, %r37;
ld.global.u32 %r40, [%rd14+384];
add.s32 %r45, %r40, %r39;
add.s64 %rd14, %rd14, 512;
add.s32 %r44, %r44, 128;
setp.lt.s32     %p5, %r44, %r3;
@%p5 bra BB7_5;

【讨论】：