CUDA 内核调用中的隐式构造函数

Question

我正在尝试将一些 POD 传递给内核，该内核具有一些非 POD 作为参数，并且具有非显式构造函数。其背后的想法是：在主机上分配一些内存，将内存传递给内核，并将内存封装在对象中，而无需用户显式执行该步骤。

构造函数被标记为 __device__ 代码，但传递参数时不会调用它们，我不明白为什么。

我的问题实际上与我应该如何做这件事无关，而是试图了解幕后发生的事情。

这里有一个例子（我使用的是 CUDA 5，GPU 的能力为 2.1，因此是 printf）。

#include <stdio.h>

struct Test {
    __device__ Test() {
        printf("Default\n"),
        _n = 0;
    }
    __device__ Test(int n) {
        printf("Construct %d\n", n);
        _n = n;
    }
    __device__ Test(const Test &t) {
        printf("Copy constr %d\n", t._n);
        _n = t._n;
    }
    __device__ Test &operator=(const Test &t) {
        printf("Assignment %d\n", t._n);
        _n = t._n;
        return *this;
    }
    __device__ int calc() const {
        printf("Calculating %d\n", threadIdx.x + 10 * _n);
        return threadIdx.x + 10 * _n;
    }
    int _n;
};

__global__ void dosome(Test a, Test b) {
    printf("Kernel data %d %d\n", a._n, b._n);
    a.calc();
    b.calc();
}

int main(int argc, char **argv) {
    dosome<<<1, 2>>>(2, 3);
    cudaError_t cudaerr = cudaDeviceSynchronize();
    if (cudaerr != cudaSuccess)
        printf("kernel launch failed with error:\n\t%s\n",cudaGetErrorString(cudaerr));
    return 0;
}

编辑：忘了说，没有打印任何构造函数消息，但打印了 calc 和内核消息。

EDIT2：是否保证 CUDA 在将测试对象复制到设备上之前会初始化它？

Answer 1

您必须像普通方法一样看到构造函数。如果您使用 限定它__host__，那么您将能够将其称为主机端。如果您使用 限定它__device__，您将能够将其称为设备端。如果你用两者来限定它，你就可以在两边调用它。

当您这样做时会发生的情况dosome<<<1, 2>>>(2, 3);是，这两个对象是在explicit主机端隐式构造的（因为您的构造函数不是，所以也许这也让您感到困惑），然后再memcpy传递到设备。该过程中不涉及复制构造函数。

让我们来说明一下：

    __global__ void dosome(Test a, Test b) {
        a.calc();
        b.calc();
    }

    int main(int argc, char **argv) {
        dosome<<<1, 2>>>(2, 3); // Constructors must be at least __host__
        return 0;
    }

// Outputs:
Construct 2 (from the host side)
Construct 3 (from the host side)

现在，如果您将内核更改为采用ints 而不是Test：

__global__ void dosome(int arga, int argb) {
    // Constructors must be at least __device__
    Test a(arga);
    Test b(argb);
    a.calc();
    b.calc();
}

int main(int argc, char **argv) {
    dosome<<<1, 2>>>(2, 3);
    return 0;
}

// Outputs:
Construct 2 (from the device side)
Construct 3 (from the device side)