如何在设备和主机上调用 constexpr 函数？

Question

来自文档：

I.4.20.4。 constexpr 函数和函数模板

默认情况下，一个 不能从不兼容的函数调用 constexpr 函数 执行空间。实验性 nvcc 标志 --expt-relaxed-constexpr 删除此限制。当指定此标志时，主机代码可以 调用 __device__ constexpr 函数，设备代码可以调用 __host__ constexpr 函数。

我读过它，但我不明白它是什么意思 - device code can invoke a host constexpr function?这是我的测试：

constexpr int bar(int i)
{
#ifdef __CUDA_ARCH__
  return i;
#else
  return 555;
#endif
}

__global__ void kernel()
{
  int tid = (blockDim.x * blockIdx.x) + threadIdx.x;

  printf("%i\n", bar(tid));
}

int main(int argc, char *[])
{
  static_assert(bar(5) > 0);
  // static_assert(bar(argc) > 0); // compile error

  cout << bar(argc) << endl;

  kernel<<<2, 2>>>();
  gpuErrchk(cudaPeekAtLastError());

  gpuErrchk(cudaDeviceSynchronize());

  return 0;
}

打印出来：

555
0
1
2
3

据我了解，主机调用host函数，而设备调用device函数。 IE。它的行为就像我用 __host__ 和 __device__ 属性声明 bar 一样。添加单个属性（__host__ 或 __device__）没有任何区别。

作为比较，std::initializer_list 的文档更加清晰：

I.4.20.2。 std::initializer_list

默认情况下，CUDA 编译器将 隐式考虑 std::initializer_list 的成员函数 有 __host__ __device__ 执行空间说明符，因此 它们可以直接从设备代码中调用。

这里我没有任何问题。

文档的确切含义是什么？

Answer 1

考虑下面的代码。

#include <algorithm> //std::max

__global__ void kernel(int *array, int n) {
    array[0] = std::max(array[1], array[2]);
}

此代码默认不会编译。

error: calling a constexpr __host__ function("max") from a __global__ function("kernel") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

std::max 是没有__device__ 执行空间说明符的标准主机函数，因此无法从设备代码中调用。

但是，当指定编译器标志--expt-relaxed-constexpr 时，代码仍然会编译。我无法向您详细说明这是如何在内部实现的