直接从数组读取时超出地址范围答案

【问题标题】：Out of bound address when directly reading from array直接从数组读取时超出地址范围
【发布时间】：2013-03-31 05:36:52
【问题描述】：

我正在开发一个 CUDA 应用程序，它有一些用于在共享内存中分配和释放数组的例程。

在这个应用程序中（很抱歉，我无法提供），我有一个将一块内存封装为数组的类。这个类有一个count 方法，可以统计匹配某个值的元素个数。

所以，想象一下（这是整个班级的实际部分）

template <class Type>
struct Array {
    // ...

    Type &operator[](int i) { return data_[i]; }
    Type operator[](int i) const { return data_[i]; }

    size_t count(const Type &val) const {
        size_t c = 0;
        for (size_t i = 0; i < len_; ++i)
            if (data_[i] == val)
                ++c;
        return c;
    }

    void print(const char *fmt, const char *sep, const char *end) const {
        for (size_t i = 0; i < len_ - 1; ++i) {
            printf(fmt, data_[i]);
            printf(sep);
        }
        printf(fmt, _data[len_ - 1]);
        printf(end);
    }
private:
    Type *data_;
    size_t len_;
};

假设我正在访问的内存是正确分配的（在运行时分配的共享内存，将维度传递给内核），它足够大以包含数据并且data_ 指向对齐的（wrtType）区域内部共享内存。我检查了多次，这些假设应该是有效的（但请随时询问更多检查）。

现在，在测试代码时，我发现了一些很奇怪的东西：

使用operator[] 显式分配值并使用operator[] const 读取它们时，不会出现问题。
使用print 读取数据时，不会出现任何问题。
调用count()时，程序崩溃，cuda-memcheck报告Address ADDR is out of bounds，由Invalid __global__ read of size x引起(x = sizeof(Type))。 ADDR 位于共享内存缓冲区内，因此它应该是有效的。
如果在count 内部，我将data_[i] 替换为(*this)[i]，程序运行良好，不会发生崩溃。

现在，我完全不知道会发生这种情况，也不知道要检查什么以查看幕后发生的事情……为什么直接阅读会崩溃？为什么不使用operator[]？为什么在print 内阅读（直接？）不会崩溃？

我知道这个问题很难，很抱歉提供有关代码的这些小信息...但是请随时询问详细信息，我会尽力回答。欢迎任何想法或建议，因为这是我试图解决的日子，这是我所能得到的。

我正在使用两种不同的 GPU 来测试此代码，一种具有 2.1 的能力，一种具有 3.5 的能力（后者提供了有关此崩溃的详细信息，而第一个没有）。 CUDA 5.0

编辑：我找到了一个发生此错误的最小示例。奇怪的是，使用 sm_20 和 sm_35 编译时会出现错误，但在 sm_30 上不会出现。我使用的 GPU 的上限为 3.5

/* Compile and run with:
  nvcc -g -G bug.cu -o bug -arch=sm_20 # bug!
  nvcc -g -G bug.cu -o bug -arch=sm_30 # no bug :|
  nvcc -g -G bug.cu -o bug -arch=sm_35 # bug!
  cuda-memcheck bug

Here's the output (skipping the initial rows) I get
Ctor for 0x3fffc10 w/o alloc, data 0x10000c8
Calling NON CONST []
Calling NON CONST []
Fill with [] ok
Fill with raw ok
Kernel launch failed with error:
        unspecified launch failure
========= Invalid __global__ write of size 8
=========     at 0x00000188 in /home/bio/are/AlgoCUDA/bug.cu:26:array<double>::fill(double const &)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x010000c8 is out of bounds
=========     Device Frame:/home/bio/are/AlgoCUDA/bug.cu:49:kernel_bug(unsigned long) (kernel_bug(unsigned long) : 0x8c0)
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/libcuda.so (cuLaunchKernel + 0x3dc) [0xc9edc]
=========     Host Frame:/opt/cuda-5.0/lib64/libcudart.so.5.0 [0x13324]
=========     Host Frame:/opt/cuda-5.0/lib64/libcudart.so.5.0 (cudaLaunch + 0x182) [0x3ac62]
=========     Host Frame:bug [0xbb8]
=========     Host Frame:bug [0xaa7]
=========     Host Frame:bug [0xac4]
=========     Host Frame:bug [0xa07]
=========     Host Frame:/lib/libc.so.6 (__libc_start_main + 0xfd) [0x1ec4d]
=========     Host Frame:bug [0x8c9]
=========
========= Program hit error 4 on CUDA API call to cudaDeviceSynchronize 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/libcuda.so [0x26a180]
=========     Host Frame:/opt/cuda-5.0/lib64/libcudart.so.5.0 (cudaDeviceSynchronize + 0x1dd) [0x441fd]
=========     Host Frame:bug [0xa0c]
=========     Host Frame:/lib/libc.so.6 (__libc_start_main + 0xfd) [0x1ec4d]
=========     Host Frame:bug [0x8c9]
=========
========= ERROR SUMMARY: 2 errors


(cuda-gdb) set cuda memcheck on
(cuda-gdb) run
Starting program: /home/bio/are/AlgoCUDA/bug 
[Thread debugging using libthread_db enabled]
[New Thread 0x7ffff5c25700 (LWP 23793)]
[Context Create of context 0x625870 on Device 0]
[Launch of CUDA Kernel 0 (kernel_bug<<<(1,1,1),(1,1,1)>>>) on Device 0]
Memcheck detected an illegal access to address (@global)0x10000c8

Program received signal CUDA_EXCEPTION_1, Lane Illegal Address.
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 12, warp 0, lane 0]
0x0000000000881928 in array<double>::fill (this=0x3fffc10, v=0x3fffc08) at bug.cu:26
26                              data[i] = v;
*/

#include <stdio.h>

extern __shared__ char totalSharedMemory[];

template <class Type>
struct array {
    // Create an array using a specific buffer
    __device__ __host__ array(size_t len, Type *buffer):
        len(len),
        data(buffer) {
        printf("Ctor for %p w/o alloc, data %p\n", this, data);
    }
    __device__ __host__ Type operator[](int i) const {
        printf("Calling CONST []\n");
        return data[i];
    }
    __device__ __host__ Type &operator[](int i) {
        printf("Calling NON CONST []\n");
        return data[i];
    }
    __device__ __host__ void fill(const Type &v) {
        for (size_t i = 0; i < len; ++i) data[i] = v;
    }
    size_t len;
    Type *data;
};

__global__ void kernel_bug(size_t bytesPerBlock) {
    // This is a test writing to show that filling the memory
    // does not produce any error
    for (size_t i = 0; i < bytesPerBlock; ++i) {
        totalSharedMemory[i] = i % ('z' - 'a' + 1) + 'a';
        printf("[%p] %c\n", totalSharedMemory + i, totalSharedMemory[i]);
    }

    // 200 / 8 = 25 so should be aligned
    array<double> X(2, (double *)(totalSharedMemory + 200));
    X[0] = 2;
    X[1] = 4;
    printf("Fill with [] ok\n");
    X.data[0] = 1;
    X.data[1] = 0;
    printf("Fill with raw ok\n");
    X.fill(0); // Crash here
    printf("Fill with method ok\n");
}

int main(int argc, char **argv) {
    // Total memory required
    size_t bytesPerBlock = 686; // Big enough for 85 doubles
    kernel_bug<<<1, 1, bytesPerBlock>>>(bytesPerBlock);
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        fprintf(stderr, "Kernel launch failed with error:\n\t%s\n", cudaGetErrorString(err));
        return 1;
    }
    return 0;
}

编辑：也使用 CUDA 4.2 进行了测试，问题仍然存在。

【问题讨论】：

您没有显示设置 len_ 的代码，但在打印函数中使用 len_ - 1，但在计数中使用 len_。我怀疑你在设置 len_ 时遇到了一个错误。
删除你的拷贝构造函数和赋值运算符，看看问题是否仍然存在，或者你得到编译错误。
@Muscles，哦，好主意，但实际上这不是问题，因为我的测试用例中的 len_ >1。
这只是一种预感，也许你的记忆被错误的副本释放了。对 cuda 不够熟悉，无法再添加。祝你好运。
@AkiRoss：您可能发现了编译器错误。但是，如果您不能为其他人可以编译的可管理大小的 repro 案例编写和提供代码（听起来这可能是一个非常短的 repro 案例），那么您或其他人怎么知道？

标签： c++ memory memory-management cuda

【解决方案1】：

当您试图查明崩溃时，最好在 X.fill(0);call 中删除从 0 到 0.0 的隐式转换。它是有效的 C++，但 CUDA 在函数调用运算符中分配临时对象时可能会遇到麻烦。确实，浏览他们的文档时，我找不到关于将此类临时人员分配到何处的答案——全球？设备？不过，这可能不是问题，但是……可以肯定。

【讨论】：

我删除了转换，但这并没有解决。无论如何感谢您的想法；）

【解决方案2】：

我能够通过以下方式重现您的问题：

RHEL 5.5、驱动程序 304.54、CUDA 5.0、Quadro 5000 GPU。

我无法重现以下问题：

RHEL 5.5、驱动程序 319.72、CUDA 5.5、Quadro 5000 GPU。

请将您的 CUDA 安装更新到 CUDA 5.5，并将您的驱动程序更新到 319.72 或更高版本。

【讨论】：