在编写类似的 CUDA 内核时如何不重复自己没有宏？

Posted 2023-02-25

【中文标题】在编写类似的 CUDA 内核时如何不重复自己没有宏？

【英文标题】：How not to repeat myself without macros when writing similar CUDA kernels?

【发布时间】：2013-05-13 08:16:26

【问题描述】：

我有几个 CUDA 内核，它们基本相同，但有一些变化。我想做的是减少所需的代码量。我的第一个想法是使用宏，所以我得到的内核看起来像这样（简化）：

__global__ void kernelA( ... )

   INIT(); // macro to initialize variables

   // do specific stuff for kernelA
   b = a + c;

   END(); // macro to write back the result


__global__ void kernelB( ... )

   INIT(); // macro to initialize variables

   // do specific stuff for kernelB
   b = a - c;

   END(); // macro to write back the result

...

由于宏令人讨厌、丑陋和邪恶，我正在寻找一种更好、更清洁的方法。有什么建议吗？

（switch 语句无法完成这项工作：实际上，相同的部分和内核特定的部分非常交织在一起。需要几个 switch 语句，这会使代码变得非常不可读。此外，函数调用不会初始化所需的变量。）

（这个问题可能也适用于一般 C++，只需将所有 'CUDA kernel' 替换为 'function' 并删除 '__global__' ）

【问题讨论】：

上课？职能？结构编程？

我们需要更多详细信息来帮助您。如果可以选择切换，@talonmies 对this question 的回答可能会对您有所帮助。

OOP 通常会回答“几个 switch 语句 [...] 这会使代码非常难以阅读。”使用“创建对象层次结构并将开关分支拆分为类特化”。为什么那行不通？

谁告诉你宏是肮脏、丑陋和邪恶的？它们是 C 语言最好的特性之一。不要害怕使用它们。嗯，嗯，你也有 C++ 模板。

【参考方案1】：

更新：我在 cmets 中被告知，类和继承不能与 CUDA 很好地混合。因此，只有答案的第一部分适用于 CUDA，而其他部分是对您问题中更一般的 C++ 部分的回答。

对于 CUDA，您将不得不使用纯函数，“C 风格”：

struct KernelVars 
  int a;
  int b;
  int c;
;

__device__ void init(KernelVars& vars) 
  INIT(); //whatever the actual code is


__device__ void end(KernelVars& vars) 
  END(); //whatever the actual code is


__global__ void KernelA(...) 
  KernelVars vars;
  init(vars);
  b = a + c;
  end(vars);

这是通用 C++ 的答案，您可以使用 OOP 技术，如构造函数和析构函数（它们非常适合那些初始化/结束对），或者也可以与其他语言一起使用的模板方法模式：

使用 ctor/dtor 和模板，“C++ 风格”：

class KernelBase 
protected:
  int a, b, c;

public:
  KernelBase() 
    INIT(); //replace by the contents of that macro
     
  ~KernelBase() 
    END();  //replace by the contents of that macro
  
  virtual void run() = 0;
;

struct KernelAdd : KernelBase 
  void run()  b = a + c; 
;

struct KernelSub : KernelBase 
  void run()  b = a - c; 
;

template<class K>
void kernel(...)

  K k;
  k.run();


void kernelA( ... )  kernel<KernelAdd>(); 

使用模板方法模式，通用“OOP风格”

class KernelBase 
  virtual void do_run() = 0;
protected:
  int a, b, c;
public:
  void run()  //the template method
    INIT(); 

    do_run();

    END();
  
;

struct KernelAdd : KernelBase 
  void do_run()  b = a + c; 
;

struct KernelSub : KernelBase 
  void do_run()  b = a - c; 
;

void kernelA(...)

  KernelAdd k;
  k.run();

【讨论】：

遗憾的是没有类，没有ctors，没有dtors，也没有继承。但是您的纯 C 风格几乎是正确的（只有 init 和 end 必须是 `__device__´ 函数）。但是 +1！

@kronos 就像我说的，我不知道 CUDA，所以我不知道它的限制。您可以看到涉及类的解决方案作为问题的“通用 C++”部分的答案。

感谢您的回答，我认为将所有内核变量放在一个结构中并使用 device 函数就可以了。

对不起我的错。在 2.0 之前的计算能力上，我有点卡在这个话题上的接缝。我必须纠正自己。类、ctors 和继承具有次要或主要限制。

【参考方案2】：

您可以将设备函数用作“INIT()”和“END()”的替代方案。

__device__ int init()

    return threadIdx.x + blockIdx.x * blockDim.x;

另一种选择是使用函数模板：

#define ADD 1
#define SUB 2

template <int __op__> __global__ void caluclate(float* a, float* b, float* c)

   // init code ...
switch (__op__)

case ADD:
  c[id] = a[id] + b[id];
break;
case SUB:
  c[id] = a[id] - b[id];
break;
    
    // end code ...

并使用以下方法调用它们：

calcualte<ADD><<<...>>>(a, b, c);

CUDA 编译器完成这项工作，构建不同的函数版本并删除死代码部分以优化性能。