1

我有一个在调试版本中工作但在发布版本中失败的内核。我怀疑我的读写越界，但 CUDA 内存检查器没有显示任何错误。所以，我做了一个测试，我用以下内核写和读越界：

__global__ void addKernel(int *c, const int *a, const int *b)
{
  int x[1];
  for (int i(0); i < 100; ++i) {
    x[i] = i;
  }
  int t(0);
  for (int i(0); i < 100; ++i) {
    t += x[i];
  }
  c[0] = t;
}

CUDA 内存检查器没有捕获越界写入和读取。通过将循环增加到 1,000,000 轮，我能够检测到越界写入，这导致内核在 i = 502,586 处暂停（2MB 越界）。

这是 CUDA 内存检查器工作的预期粒度吗？我能做些什么来让它检测小的越界写入（大约几个字节？）

4

1 回答 1

1

我认为你只是在这个例子中被优化绊倒了，正如已经建议的那样。

这是我的测试用例：

$ cat t1130.cu
#include <stdio.h>

__global__ void addKernel(int *c)
{
  int x[1];
  for (int i(0); i < 100; ++i) {
    x[i] = i;
#ifdef FORCE
  printf("%d ", i);
#endif
  }
  int t(0);
  for (int i(0); i < 100; ++i) {
    t += x[i];
  }
  c[0] = t;
}

int main(){

  int *d_c;
  cudaMalloc(&d_c, sizeof(int));
  addKernel<<<1,1>>>(d_c);
  cudaDeviceSynchronize();
}
$ nvcc -o t1130 t1130.cu
$ cuda-memcheck ./t1130
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$ nvcc -DFORCE -o t1130 t1130.cu
$ cuda-memcheck ./t1130
========= CUDA-MEMCHECK
========= Invalid __local__ write of size 4
=========     at 0x00000168 in addKernel(int*)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x00fffd10 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x15865d]
=========     Host Frame:./t1130 [0x16ca1]
=========     Host Frame:./t1130 [0x314b3]
=========     Host Frame:./t1130 [0x27a1]
=========     Host Frame:./t1130 [0x269c]
=========     Host Frame:./t1130 [0x26b6]
=========     Host Frame:./t1130 [0x2600]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21d65]
=========     Host Frame:./t1130 [0x2489]
=========
0 1 2 3 ========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/lib64/libcuda.so.1 [0x2f31b3]
=========     Host Frame:./t1130 [0x354a6]
=========     Host Frame:./t1130 [0x2605]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21d65]
=========     Host Frame:./t1130 [0x2489]
=========
========= ERROR SUMMARY: 2 errors
$

因此，如果我们按原样编译您的代码，它会在cuda-memcheck. 但是如果我们强制编译器生成实际的循环，那么就会检测到并报告无效访问。

还解释您的评论，如果我们在上述情况下使用-G而不是，也将检测并报告错误，因为优化被禁用。-DFORCEcuda-memcheck

于 2016-04-01T21:27:34.290 回答