在 CUDA 中选择性地编译头文件和类函数

Posted 2023-02-16

【中文标题】在 CUDA 中选择性地编译头文件和类函数

【英文标题】：selectively compile headers and class functions in CUDA

【发布时间】：2013-03-28 19:45:38

【问题描述】：

我正在尝试在 CUDA 中使用我的 c++ 类。

我有这样的课程：

#include<string>
#include<stdlib.h>

class exampleClass
int i;
__host__ __device__ exampleClass(int _i):i(_i);
__host__ __device__ void increment()i++;
__host__ __device__ string outputMessage(return itoa(i);

;

我已将其设置在 .cu 文件中并设置为编译 CUDA c/c++

这无法使用 nvcc 编译，因为 cuda 没有字符串。

我想做的是保留 CUDA 唯一的功能，方法是：

#ifndef __CUDA_ARCH__
  #include<string>
#endif
    #include<stdlib.h>

    class exampleClass
    int i;
    __host__ __device__ exampleClass(int _i):i(_i);
    __host__ __device__ void increment()i++;
#ifndef __CUDA_ARCH__
     string outputMessage(return itoa(i);
#endif

    ;

但我知道这不起作用……至少，它对我不起作用。 nvcc 不喜欢包含字符串，也不喜欢需要字符串类型的函数。

抱歉，如果示例不是一流的。总之，我想做的是让核心类成员在 CUDA 上可执行，同时保持在主机端进行花哨的主机操作以进行分析和输出的能力。

更新：我的最终目标是创建一个基类，其中包含多个指向多个多态类的指针类型。这个基类本身将是可派生的。我认为这在 CUDA5.0 中是可能的。我弄错了吗？

【问题讨论】：

在 nvcc 中使用 stl 类（例如字符串或向量）没有问题。你当然不能将字符串发送到内核，或者在内核中使用字符串函数，但这看起来不像你想要做的......你应该能够删除 "设备“装饰。

MadScienceDreams 对我的问题的回答是正确的。然而，问题背后的问题在这里找到了解决方案：***.com/questions/5994005/…

更具体地说，您需要在 CUDA C/C++ 编译器中包含 -dc 标志以允许多个 cu 文件之间的链接。

【参考方案1】：

以下代码构建，虽然我没有运行它：

class exampleClass
int i;
public:
__host__ __device__ exampleClass(int _i):i(_i);
__host__ __device__ void increment()i++;

 __host__ string outputMessage() return "asdf";


;

__global__ void testkernel (                        
    exampleClass *a,
    int IH, int IW)

    const int i = IMUL(blockIdx.x, blockDim.x) + threadIdx.x;
    const int j = IMUL(blockIdx.y, blockDim.y) + threadIdx.y;


    if (i<IW && j<IH) 
    
        const int i_idx = i + IMUL(j, IW);  
        exampleClass* ptr = a+i_idx;
        ptr->increment();
    


__host__ void test_function(exampleClass *a,
    int IH, int IW)

    for (int i = 0; i < IW; i++)
        for (int j = 0; j < IH; j++)
        
            const int i_idx = i + j*IW;
            exampleClass* ptr = a+i_idx;
            cout << ptr->outputMessage();
        

请注意，您必须将类从设备内存移动到主机内存才能正常“工作”。如果您尝试对类做任何花哨的事情（例如多态性），这可能会失败。

【讨论】：

我认为在 CUDA5.0 中允许多态性？我的“exampleClass”是一个有几个派生类的基类。同样，exampleClass 包含几个持有多态类的点。

@PaulD。嗯，从 4.0 开始，这是一个很酷的变化。 CUDA 是为数不多的软件包之一，似乎每年都会增加一些不错的功能！这些类的 VTable 必须同时引用 __host__ 和 __device__ 函数指针，因此类大小会增长得很快。请注意，仍然不支持虚拟继承。

我对 VTable 不熟悉...我不仅仅是自学成才。你能解释一下或发布一个链接吗？同样，nVidia 网站表示自 Toolkit 4.0 起就支持虚拟继承。 developer.nvidia.com/cuda-toolkit-40

代码标准5说"It is not allowed to pass as an argument to a global function an object of a class derived from virtual base classes."，所以没有虚拟继承。它还指出“编译器将虚函数表放置在全局或常量内存中”。虽然我不知道这是怎么发生的（静态负载？）。 [***中的 VTable](en.wikipedia.org/wiki/Virtual_method_table)

我的印象是“虚拟基类”是一种特殊的继承类型：***.com/questions/21558/…