cuda-memcheck не сообщает правильно при использовании в устройстве mallo c () и cudaMallo c () в одной программе

Question

Я использовал проверку утечек cuda-memcheck на небольшом тестовом файле, который я сделал, чтобы проверить некоторые функции, которые я хотел реализовать в программе, над которой я работаю, и обнаружил, что он не сообщает о некоторых очень очевидных утечках памяти в глобальной памяти, когда у меня есть вызовы как cudaMalloc() (из кода хоста), так и malloc() (из кода устройства). Похоже, что вызов устройства malloc() нарушает функциональность cuda-memcheck.
Я запускаю это на NVIDIA GeForce GTX 1050 (вычислительная мощность 6.1), на Windows 10. У меня CUDA v10.2, с помощью компилятора Visual Studio C ++ (cl.exe). Мой друг также запускал это на своей системе Arch Linux с CUDA v9.1 и NVIDIA GeForce MX150 (вычислительная мощность 6.1) с идентичными результатами. Вот код, который я использовал:

#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());
}

Здесь довольно ясно, что есть 2 утечки памяти, поскольку я не освобождаю d_t, а также указатель на член 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 не освобожден из кучи устройства (я не проверял явно, является ли это точной утечкой, я просто предполагаю)
Теперь, когда я также удалил вызов до d_t->lol(), а также print_yay() из Tester::print() (в основном удалить код, который выделяет память для d_t->yay из кучи устройства с помощью устройства malloc(), и код, который читает 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, или что-то не так с моим кодом?