Reputation: 49
I was using the cuda-memcheck's leak-check on a small test file I made to test some functionality I wanted to implement in a program I'm working on, and I found that it is not reporting some very obvious memory leaks on global memory, when I have calls to both cudaMalloc()
(from host code) and malloc()
(from device code). It seems the call to device malloc()
is breaking the functionality of cuda-memcheck
.
I'm running this on NVIDIA GeForce GTX 1050(compute capability 6.1), on Windows 10. I have CUDA v10.2, using the Visual Studio C++ compiler(cl.exe
). A friend of mine also ran this on his Arch Linux system, with CUDA v9.1, and an NVIDIA GeForce MX150(compute capability 6.1), with identical results. Here's the code I used:
#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());
}
Here it is pretty clear that there are 2 memory leaks, as I'm not freeing d_t
, as well as the member pointer b0
, using cudaFree()
. I compiled this using nvcc.exe -G -Xcompiler /Zi -o cuda cuda.cu
, and then ran cuda-memcheck.exe --leak-check full cuda.exe
. The output:
========= 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
When I removed the call to d_t->cleanup()
from the kernel, the output:
========= 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
The 2 byte leak is most probably because d_t->yay
is not freed from the device heap.(I haven't checked explicitly if that is the exact leak though, I'm just guessing)
Now, when I also removed the call to d_t->lol()
, as well as print_yay()
from Tester::print()
(basically remove the code that allocates memory to d_t->yay
from device heap using device malloc()
and the code that reads d_t->yay
), so now the kernel looks like:
__global__ void kernel(Tester * d_t) {
d_t->print();
}
the output:
========= 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
Clearly, it showed the correct leaks.
Also, I noticed something else that's weird. When my kernel is:
__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
which is basically the same as the one above, as d_t->cleanup()
wouldn't do anything anyway, the output:
========= 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
It still stopped showing the leak!
Is this some issue with cuda-memcheck
, or is there something wrong with my code?
Upvotes: 4
Views: 939