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

当某些代码开始提供不同的结果时，我正在升级到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（数组的大小），则将不会访问它（an illegal memory access was encountered (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上编译。在369.30的Titan上运行。

最佳答案

GPU具有运行时内存检查器，其详细信息尚未发布。此内存检查工具不是十分精确，但是如果发生了严重的错误（例如，以足够的余量进行越界访问），则运行时内存检查将标记错误，暂停内核并声明上下文为损坏了。

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

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

从CUDA版本到CUDA版本，从GPU架构到体系结构，这样的内存映射很有可能会变化，甚至可能取决于特定的编译开关。

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

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