直接从数组读取时超出了绑定地址

我正在开发一个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_指向共享内存中的对齐(wrt Type)区域。我检查了很多次，这些假设应该是有效的(但请随意要求更多的检查)。

现在，在测试代码时，我发现了一些非常奇怪的东西:

• 当使用operator[]显式赋值，并使用operator[] const读取它们时，不会出现问题。
• 当使用print读取数据时，没有问题。
• 当调用count()时，程序崩溃，Address ADDR is out of bounds由cuda-memcheck报告，Invalid __global__ read of size x (x = sizeof(Type))引起。地址在共享内存缓冲区内，所以它应该是有效的。
• 如果，在count内，我用(*this)[i]代替data_[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 %pn", 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] %cn", 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 [] okn");
    X.data[0] = 1;
    X.data[1] = 0;
    printf("Fill with raw okn");
    X.fill(0); // Crash here
    printf("Fill with method okn");
}
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:nt%sn", cudaGetErrorString(err));
        return 1;
    }
    return 0;
}

编辑:也测试了CUDA 4.2，问题仍然存在。