cuda-memcheck とかいう CUDA で valgrind みたいなことができるツールがある。
↓こんなードを書いたとする。
#include <stdio.h> __global__ void f(int *ptr) { int tid = threadIdx.x; ptr[tid] = tid; } int main() { int *dmem; int hmem[32]; memset(hmem, 0, sizeof(hmem)); cudaMalloc(&dmem, 32); cudaMemcpy(dmem, hmem, 32, cudaMemcpyHostToDevice); f<<<1,32>>>(dmem); cudaMemcpy(hmem, dmem, 32, cudaMemcpyDeviceToHost); printf("%d\n", hmem[4]); }
いちおう動いているように見えるが、間違っている。(どこが間違ってるかすぐわかる人偉い)
ビルドするときに、
$ nvcc -G -g hoge.cu
とかする(必須ではない)
んで、
$ cuda-memcheck ./a.out
とかやると
========= Invalid __global__ write of size 4 ========= at 0x000000d0 in /home/w0/test/cuda/memcheck.cu:6:f(int*) ========= by thread (31,0,0) in block (0,0,0) ========= Address 0x70014007c is out of bounds ========= Saved host backtrace up to driver entry point at kernel launch time ========= Host Frame:/lib64/libcuda.so (cuLaunchKernel + 0x331) [0xcd5d1] ========= Host Frame:./a.out [0x1b698] ========= Host Frame:./a.out [0x3ab63] ========= Host Frame:./a.out [0x2937] ========= Host Frame:./a.out [0x280d] ========= Host Frame:./a.out [0x2827] ========= Host Frame:./a.out [0x2750] ========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21b75] ========= Host Frame:./a.out [0x24e9]
メモリアクセスエラーしてるのがわかる。
あと cuda-gdb から set cuda memcheck しても同じようなことをデバッガからできる、とか書いてある。
http://docs.nvidia.com/cuda/cuda-memcheck/index.html
