CUDA GPU 上的 printf 问题

Question

#include <iostream>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>

#include <iostream>
#include <cmath>
#include <vector>

using namespace std;

#define HOST __host__
#define DEVICE __device__

template<class L>
__global__ void launch_global (L f0) { f0(); }

template<typename T>
struct Array{
    T* __restrict__ data;

    DEVICE HOST inline
    T& operator()(int i)const noexcept {
        return data[i];
    }
};


int main(){

    int nx = 5;

    auto vec = Array<double>{new double[nx]};

    Array<double> *vec1;
    cudaMallocManaged((void**)&vec1, nx * sizeof(double));

    launch_global<<<1,256>>>([=] DEVICE () noexcept{
        int i = blockDim.x*blockIdx.x+threadIdx.x;
        if(i < nx){
             (*vec1)(i) = 1.0;
             printf("Printing here %d %g\n", i,1.0);
        }
    });

    cudaDeviceSynchronize();

return 0;
}

编译：

nvcc -x cu --expt-extended-lambda --expt-relaxed-constexpr --expt-relaxed-constexpr kernel_test.cpp -o out

问题：

问题是打印语句 - printf("Printing here %d %g ", i, 1.0), 没有被执行。但是如果我评论前一行 - vec(i)=1.0，那么打印命令就会执行。有人能帮我弄清楚为什么吗？

我还尝试使用 cuda-memcheck 运行。没有错误。

Answer 1

各种问题

首先，这个：

auto vec = Array<double>{new double[nx]};

通过使用

new

创建一个分配的指针。无论您随后做什么，该指针 将永远无法在设备代码中使用。这不是使用 UM 的正确策略。

接下来，你的分配大小在这里没有意义：

cudaMallocManaged((void**)&vec1, nx * sizeof(double));

vec1

是指向

Array<double>

类型的指针。根据

sizeof(double)

分配尺寸是没有意义的。

如果要使用 UM（托管内存）执行此操作，则需要处理类对象的分配以及对象中嵌入指针的分配。

您通常会感到困惑，因为您正在分配一个

Array<double>

对象的数组，每个对象都有一个指向

double

项目的嵌入式指针。下面的代码修复了这些问题，对我来说似乎可以正常运行：

$ cat t2239.cu
#include <iostream>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>

#include <iostream>
#include <cmath>
#include <vector>

using namespace std;

#define HOST __host__
#define DEVICE __device__

template<class L>
__global__ void launch_global (L f0) { f0(); }

template<typename T>
struct Array{
    T* __restrict__ data;

    DEVICE HOST inline
    T& operator()(int i)const noexcept {
        return data[i];
    }
};


int main(){

    int nx = 5;

    Array<double> *vec;
    cudaMallocManaged((void**)&vec, sizeof(Array<double>));
    cudaMallocManaged((void**)&(vec[0].data), nx*sizeof(double));


    launch_global<<<1,256>>>([=] DEVICE () noexcept{
        for(int i = blockDim.x*blockIdx.x+threadIdx.x, stride = blockDim.x*gridDim.x;
             i < nx; i += stride){
             vec[0](i) = 1.0;
             printf("Printing here %d %g\n", i,1.0);
        }
    });

    cudaDeviceSynchronize();

return 0;
}
$ nvcc -o t2239 t2239.cu  --extended-lambda -lineinfo
$ compute-sanitizer ./t2239
========= COMPUTE-SANITIZER
Printing here 0 1
Printing here 1 1
Printing here 2 1
Printing here 3 1
Printing here 4 1
========= ERROR SUMMARY: 0 errors
$