c++ - 为什么 CUDA 8.0(有时)内存访问错误，而 7.5 却没有？

Question

当一些代码开始给出不同的结果时，我正在升级到 CUDA 8.0。我设法用 MCVE 大致复制了这个问题并解决了我的问题。

#include <cub/cub.cuh> // Tested with cub 1.5.5

#include <stdio.h>

static inline void f(cudaError_t err, const char *file, int line)
{
    if (err != cudaSuccess) {
        fprintf(stderr, "ERROR in file %s, line %d: %s (%d)\n", file, line, cudaGetErrorString(err), err);
        fprintf(stdout, "ERROR in file %s, line %d: %s (%d)\n", file, line, cudaGetErrorString(err), err);
    }
}

#define CHKERR(expr) do {f(expr, __FILE__, __LINE__);} while(0)

template<int dimSize>
__device__ __inline__ void UsedToWork(double *s_arr)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    typedef cub::BlockReduce<double, dimSize> BlockReduce;
    __shared__ typename BlockReduce::TempStorage temp_storage;

    // This following line was the issue
    double r = BlockReduce(temp_storage).Sum(s_arr[idx], dimSize);
    __syncthreads();
    if (idx == 0)
        printf("t0 here %f\n\n", r);
}

template<int size>
__global__ void ShouldWork(double *input)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    __shared__ double s_arr[size];
    if (idx < size)
        s_arr[idx] = input[idx];
    __syncthreads();

    UsedToWork<size>(s_arr);
}

int main()
{
    const int arraySize = 32;
    double h[arraySize] = { 
         1,  2,  3,  4,  5,  6,  7,  8,  9, 10,
        11, 12, 13, 14, 15, 16, 17, 18, 19, 20,
        21, 22, 23, 24, 25, 26, 27, 28, 29, 30,
        31, 32
    };

    double *d = 0;
    cudaError_t cudaStatus;

    CHKERR(cudaMalloc((void**)&d, arraySize * sizeof(double)));
    CHKERR(cudaMemcpy(d, h, arraySize * sizeof(double), cudaMemcpyHostToDevice));

    ShouldWork<32><<<1, arraySize * 2 >>>(d);

    CHKERR(cudaGetLastError());
    CHKERR(cudaDeviceSynchronize());
    CHKERR(cudaFree(d));

    return 0;
}

我将兴趣线替换为

double r = BlockReduce(temp_storage).Sum((idx < dimSize ? s_arr[idx] : 0.), dimSize);

确保如果 idx 大于 dimSize(数组的大小)它不会被访问(遇到非法内存访问 (77) )。虽然这显然是一个错误，但为什么 CUDA 7.5 首先允许内存访问没有问题？只是为了让事情更有趣，如果在内核中我替换了

UsedToWork<size>(s_arr);

调用它的(应该是内联的)定义

typedef cub::BlockReduce<double, size> BlockReduce;
__shared__ typename BlockReduce::TempStorage temp_storage;
double r = BlockReduce(temp_storage).Sum(s_arr[idx], size);
__syncthreads();

CUDA 8.0 没有给我 an illegal memory access was encountered (77) 错误。现在我很困惑。行为不应该至少保持一致吗？

在 Windows 7、VS2013 上编译。在 Titan 上以 369.30 运行。

Answer 1

GPU 有一个运行时内存检查器，其详细信息未公开。这种内存检查工具并不十分精确，但如果发生足够严重的错误(例如，越界访问有足够的余量)，那么运行时内存检查将标记一个错误，停止内核，并声明上下文是损坏。

发生这种情况的具体条件尚未公布，可能因 GPU 架构和 GPU 架构不同，也可能因 CUDA 版本不同，以及其他可能因素而异。

正如评论中推测的那样，如果 GPU 代码触及与其上下文无关的内存，则可能存在运行时错误检查机制。然后，给定的数组越界索引可能依赖于该数组恰好位于上下文内存映射中的位置，以确定特定的越界范围是否实际上会超出上下文。

这样的内存映射很可能在不同的 CUDA 版本、不同的 GPU 架构之间变化，甚至可能取决于特定的编译开关。

为了最好(最严格)的内存访问有效性检查，建议使用 cuda-memcheck 工具。例如，如果在 cuda-memcheck 下运行，通过所有 CUDA 运行时错误检查的代码很可能会失败(并且实际上存在实际编码缺陷)。

没有明确保证 GPU 在正常操作中会检测到无效的内存访问。它当然有一定的能力这样做，但它并不完美。我相信对于我熟悉的操作环境的主机代码也可以做出类似的陈述。