作業中のプログラムに実装したい機能をテストするために作成した小さなテストファイルで cuda-memcheck のリークチェックを使用していましたが、グローバルメモリで非常に明白なメモリリークが報告されていないことがわかりました、 cudaMalloc()(ホスト コードから) とmalloc()(デバイス コードから)の両方に呼び出しがある場合。デバイスへの呼び出しがmalloc()の機能を壊しているようですcuda-memcheck。
これを Windows 10 の NVIDIA GeForce GTX 1050 (計算機能 6.1) で実行しています。Visual Studio C++ コンパイラ ( cl.exe) を使用して CUDA v10.2 を使用しています。私の友人も、CUDA v9.1 と NVIDIA GeForce MX150 (計算能力 6.1) を搭載した Arch Linux システムでこれを実行し、同じ結果が得られました。使用したコードは次のとおりです。
#define gpuErrchk(ans){ gpuAssert((ans), __FILE__, __LINE__);}
__host__
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
printf("GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
class intarr {
public:
static const int arr_size = 5;
int arr[arr_size];
__host__ __device__ void print() {
printf("arr = {");
for(int i = 0; i < arr_size; ++i) {
printf("%d", arr[i]);
if(i != arr_size-1){
printf(", ");
}
}
printf("}\n");
}
};
class Tester {
__device__ void print_yay() {
printf("yay = {");
for (int i = 0; i < yay_size; ++i) {
if (yay[i] == 'A' || yay[i] == 'a') printf("%c", yay[i]);
else printf("-");
if (i != yay_size - 1) {
printf(", ");
}
}
printf("}\n");
}
public:
intarr * b0;
char * yay;
const int yay_size;
__host__ Tester() : yay_size(2) {
intarr b;
yay = nullptr;
for(int i = 0; i < intarr::arr_size; ++i) b.arr[i] = i;
gpuErrchk(cudaMalloc(&b0, sizeof(intarr)));
gpuErrchk(cudaMemcpy(b0, &b, sizeof(intarr), cudaMemcpyDefault));
}
__device__ void lol() {
yay = (char *)malloc(2*sizeof(char));
new(yay) char[2]{'a', 'A'};
}
__device__ void print() {
b0->print();
print_yay();
}
__device__ void cleanup() {
if (yay) free(yay);
yay = nullptr;
}
/*__host__ ~Tester() {
if(b0) cudaFree(b0);
b0 = nullptr;
}*/
};
__global__ void kernel(Tester * d_t) {
d_t->lol();
d_t->print();
d_t->cleanup();
}
int main() {
printf("Tester = %zu bytes, intarr = %zu bytes\n", sizeof(Tester), sizeof(intarr));
Tester h_t;
Tester * d_t;
gpuErrchk(cudaMalloc(&d_t, sizeof(Tester)));
gpuErrchk(cudaMemcpy(d_t, &h_t, sizeof(Tester), cudaMemcpyDefault));
kernel<<<1, 1>>>(d_t);
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaGetLastError());
gpuErrchk(cudaDeviceReset());
}
d_tここでは、 を解放していないため、2 つのメモリ リークb0とcudaFree(). を使用してこれをコンパイルしnvcc.exe -G -Xcompiler /Zi -o cuda cuda.cu、実行しましcuda-memcheck.exe --leak-check full cuda.exeた。出力：
========= CUDA-MEMCHECK Tester = 24 bytes, intarr = 20 bytes arr = {0, 1, 2, 3, 4} yay = {a, A} ========= LEAK SUMMARY: 0 bytes leaked in 0 allocations ========= ERROR SUMMARY: 0 errors
d_t->cleanup()への呼び出しをカーネルから削除すると、出力は次のようになります。
========= CUDA-MEMCHECK Tester = 24 bytes, intarr = 20 bytes arr = {0, 1, 2, 3, 4} yay = {a, A} ========= Leaked 2 bytes at 0x7016fff24 on the device heap ========= ========= LEAK SUMMARY: 2 bytes leaked in 1 allocations ========= ERROR SUMMARY: 1 error
2 バイトのリークは、d_t->yayがデバイス ヒープから解放されていないことが原因である可能性が最も高いです。 from (基本的に、デバイスを使用してメモリを from デバイス ヒープに割り当てるコードと、 を読み取るコードを削除します)、カーネルは次のようになります。d_t->lol()print_yay()Tester::print()d_t->yaymalloc()d_t->yay
__global__ void kernel(Tester * d_t) {
d_t->print();
}
出力：
========= CUDA-MEMCHECK Tester = 24 bytes, intarr = 20 bytes arr = {0, 1, 2, 3, 4} ========= Leaked 24 bytes at 0x501200200 ========= Saved host backtrace up to driver entry point at cudaMalloc time ========= Host Frame:C:\WINDOWS\system32\nvcuda.dll (cuMemAlloc_v2 + 0x173) [0x19d7a3] ========= Host Frame:D:\cudatest\cuda.exe (cudart::driverHelper::mallocPtr + 0x3e) [0x4017e] ========= Host Frame:D:\cudatest\cuda.exe (cudart::cudaApiMalloc + 0x3e) [0x1ff1e] ========= Host Frame:D:\cudatest\cuda.exe (cudaMalloc + 0xdd) [0xc31d] ========= Host Frame:D:\cudatest\cuda.exe (cudaMalloc<Tester> + 0x1d) [0x48e2d] ========= Host Frame:D:\cudatest\cuda.exe (main + 0x3b) [0x4894b] ========= Host Frame:D:\cudatest\cuda.exe (__scrt_common_main_seh + 0x10c) [0x4a1b8] ========= Host Frame:C:\WINDOWS\System32\KERNEL32.DLL (BaseThreadInitThunk + 0x14) [0x17bd4] ========= Host Frame:C:\WINDOWS\SYSTEM32\ntdll.dll (RtlUserThreadStart + 0x21) [0x6ce51] ========= ========= Leaked 20 bytes at 0x501200000 ========= Saved host backtrace up to driver entry point at cudaMalloc time ========= Host Frame:C:\WINDOWS\system32\nvcuda.dll (cuMemAlloc_v2 + 0x173) [0x19d7a3] ========= Host Frame:D:\cudatest\cuda.exe (cudart::driverHelper::mallocPtr + 0x3e) [0x4017e] ========= Host Frame:D:\cudatest\cuda.exe (cudart::cudaApiMalloc + 0x3e) [0x1ff1e] ========= Host Frame:D:\cudatest\cuda.exe (cudaMalloc + 0xdd) [0xc31d] ========= Host Frame:D:\cudatest\cuda.exe (cudaMalloc<intarr> + 0x1d) [0x48e5d] ========= Host Frame:D:\cudatest\cuda.exe (Tester::Tester + 0x6d) [0x48edd] ========= Host Frame:D:\cudatest\cuda.exe (main + 0x2b) [0x4893b] ========= Host Frame:D:\cudatest\cuda.exe (__scrt_common_main_seh + 0x10c) [0x4a1b8] ========= Host Frame:C:\WINDOWS\System32\KERNEL32.DLL (BaseThreadInitThunk + 0x14) [0x17bd4] ========= Host Frame:C:\WINDOWS\SYSTEM32\ntdll.dll (RtlUserThreadStart + 0x21) [0x6ce51] ========= ========= LEAK SUMMARY: 44 bytes leaked in 2 allocations ========= ERROR SUMMARY: 2 errors
明らかに、正しいリークが示されました。
また、他にも奇妙なことに気付きました。私のカーネルが次の場合：
__global__ void kernel(Tester * d_t) {
d_t->print();
d_t->cleanup();
}
// Also print_yay() is commented out in Tester::print(), to prevent cuda-memcheck
// from terminating prematurely due to an illegal memory access error
d_t->cleanup()とにかく何もしないので、基本的に上記のものと同じです、出力：
========= CUDA-MEMCHECK Tester = 24 bytes, intarr = 20 bytes arr = {0, 1, 2, 3, 4} ========= LEAK SUMMARY: 0 bytes leaked in 0 allocations ========= ERROR SUMMARY: 0 errors
それでもリークが表示されなくなりました！
これは の問題cuda-memcheckですか、それとも私のコードに何か問題がありますか?