在金属着色语言（MSL）中，如何在地址空间上编写通用的按引用传递函数？

Question

我正在为 Metal 计算着色器编写一个

Foo

类，并希望用户能够添加

Foo

，无论它们驻留在哪个地址空间。如果我通过引用获取这些值，这意味着我必须声明事物所在地址空间的许多排列：

class Foo {
private:
  uint8_t bunchOfData[160];

public:
  operator+(const thread Foo&) const thread;
  operator+(const constant Foo&) cosnt thread;
  operator+(const thread Foo&) const constant;
  operator+(const constant Foo&) const constant;
  // ... More overloads for local. For device, users should manually unpack to maintain coalescing.
}

1. 有没有办法使用模板来通用化地址空间限定符？
2. 我应该只使用值传递并完全避免整个问题吗？

• Foo 的大小并不简单，每个 160 个字节。
• 当我上次编写计算内核时（即 2009 年 OpenCL 时代），GPU 无法进行函数调用，因此按值传递与按引用传递并没有真正产生区别，因为无论如何一切都是内联的。现在还是这样吗？

3. 如果按引用传递仍然有用，那么传统的做法是让调用者将值复制到

   thread

   空间并仅提供线程-线程变体吗？

我尝试过写：

class Foo {
private:
  uint8_t bunchOfData[160];

public:
  template <typename T> operator+(const T& rhs) const;
}

但这给了我错误

error: reference type must have explicit address space qualifier
      template <typename T> operator+(const T& rhs);

无论如何，我什至都没有正确指定

this

的地址空间...

Answer 1

按引用传递与按值传递绝对相关，除非优化器可以推断出您要执行的操作。 （如果您的类使用默认的复制构造函数等，则可能会出现这种情况）在任何情况下，每个线程不必要地复制 160 字节可能是可测量的开销。

地址空间限定符的模板化很棘手。没有特殊的语法来参数化地址空间本身。令人沮丧的是，可以参数化引用的类型，但可以固定地址空间： template <typename T> void do_foo(threadgroup T* my_t);

至少根据我的经验，这样做比允许地址空间灵活更不常见。

重载的工作方式正如您所期望的那样，对于仅在指针或引用参数的地址空间上有所不同的函数签名。至少，这意味着您可以为您关心的每种地址空间提供实现，但代价是重复。

但是，通过一些模板技巧，我们

可以

编写带有指向具有any地址空间的特定类型的指针的函数：

// define these once:
template <typename CHECK_T, typename TARGET_T> struct is_ptr_to;
template <typename T> struct is_ptr_to<thread T*, T>
{
    static constant const bool value = true;
};
template <typename T> struct is_ptr_to<threadgroup T*, T>
{
    static constant const bool value = true;
};
template <typename T> struct is_ptr_to<device T*, T>
{
    static constant const bool value = true;
};

// function where INT_PTR may be any of thread int*, threadgroup int*, or device int*:
template<typename INT_PTR, metal::enable_if_t<is_ptr_to<INT_PTR, int>::value, bool> = true>
    void set_1(INT_PTR a)
{
    *a = 1;
}

虽然不漂亮，但确实有效。 （也许更擅长模板的人可以想出更紧凑的东西；或者您可以将其包装在宏中。）

不幸的是，你不能对引用类型做完全相同的事情，因为函数模板中按值和按引用之间的歧义被解决为选择按值；一旦您在参数声明中插入 & 符号，Metal 编译器就会抱怨地址空间丢失。我尝试过的所有解决方法都需要显式模板实例化参数或用于包装参数的辅助函数。 （不过，我确实希望有人能证明我错了。）