c++ - 在编写类似的 CUDA 内核时，如何不在没有宏的情况下重复自己？

Question

我有几个 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 内核”替换为“函数”并删除“__global__”)

Answer 1

更新:我在评论中被告知，类和继承不能与 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 技术，如构造函数和析构函数(它们非常适合那些 init/end 对)，或者也可以与其他语言一起使用的模板方法模式:

使用 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();
}