CUDA 中设备类的类型限定符

Question

我目前正在尝试使用仅在设备端使用的类制作一段 CUDA 代码（即主机不需要知道它的存在）。但是我无法为该类计算出正确的限定符（下面的deviceclass）：

__device__ float devicefunction (float *x) {return x[0]+x[1];}

class deviceclass {
    private:
        float _a;

    public:
        deviceclass(float *x) {_a = devicefunction(x);}

        float getvalue () {return _a;}
};    

// Device code
__global__ void VecInit(float* A, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N) {
        deviceclass *test;

        test = new deviceclass(1.0, 2.0);

        A[i] = test->getvalue();
    }
}

// Standard CUDA guff below: Variables
float *h_A, *d_A;

// Host code
int main(int argc, char** argv)
{
    printf("Vector initialization...\n");
    int N = 10000;
    size_t size = N * sizeof(float);

    // Allocate
    h_A = (float*)malloc(size);
    cudaMalloc(&d_A, size);

    printf("Computing...\n");
    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    VecInit<<<blocksPerGrid, threadsPerBlock>>>(d_A, N);

    // Copy result from device memory to host memory
    cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost);

    //...etc
}

将Deviceclass 单独设置为__device__ 会在从全局函数调用时引发错误，但是将其设置为__device__ __host__ 或__global__ 似乎没有必要。有人能指出我正确的方向吗？

Answer 1

原来限定符必须在类的成员函数上，下面是一个完整的版本：

#include <iostream>
#include <stdio.h>
#include <stdlib.h>

using namespace std;

void Cleanup(void);


// Functions to be pointed to
__device__ float Plus (float a, float b) {return a+b;}

class deviceclass {

    private:
        float test;

    public:
        __device__ deviceclass(float a, float b) {
            test = Plus(a,b);
        }

        __device__ float getvalue() {return test;}
};

// Device code
__global__ void VecInit(float* A, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N) {
        deviceclass test(1.0, 2.0);

        A[i] = test.getvalue();
    }
}

// Standard CUDA guff below: Variables
float *h_A, *d_A;

// Host code
int main(int argc, char** argv)
{
    printf("Vector initialization...\n");
    int N = 10000;
    size_t size = N * sizeof(float);

    // Allocate
    h_A = (float*)malloc(size);
    cudaMalloc(&d_A, size);

    printf("Computing...\n");
    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    VecInit<<<blocksPerGrid, threadsPerBlock>>>(d_A, N);

    // Copy result from device memory to host memory
    cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost);



    // Verify result
    int i;
    for (i = 0; i < N; ++i) {
        cout << endl << h_A[i];
    }

    cout << endl;

    Cleanup();
}

void Cleanup(void)
{
    // Free device memory
    if (d_A)
        cudaFree(d_A);

    // Free host memory
    if (h_A)
        free(h_A);

    cudaThreadExit();

    exit(0);
}

Answer 2

我认为Node() 是一个错字。

来自 CUDA C 编程指南，第 3.1.5 节：

但是，设备代码仅完全支持 C++ 的一个子集

和附录 D.6：

为具有计算能力 2.x 及更高版本的设备编译的代码可能会使用 C++ 类...

我认为您的代码使用了不兼容的 C++。