比如用 __host__ & __device__ 的情况如下：

#include <cstdio>
#include <cuda_runtime.h>__host__ __device__ void say_hello(){printf("Hello, world!\n");
}__global__ void kernel(){say_hello();
}int main(){kernel<<<1, 1>>>();cudaDeviceSynchronize();say_hello();return 0;
}
}

则可以用 constexpr 来替代 __host__ __device，替代后的代码如下：

#include <cstdio>
#include <cuda_runtime.h>constexpr const char* cuthead(const char* p){return p+1;
}__global__ void kernel(){printf(cuthead("Gello, world!\n"));
}int main(){kernel<<<1, 1>>>();cudaDeviceSynchronize();printf(cuthead("Cello, world!\n"));return 0;
}

---

● 这样相当于把 constexpr 函数自动变成修饰符 __host__ __device__ ，从而两边都可以调用；
● 因为 constexpr 通常都是一些可以内联的函数，数学计算表达式之类的，一个个加上太累了，所以产生了这个需求；
● 不过必须指定 --expt-relaxed-constexpr 这个选项才能用这个特性，咱们可以用 CMake 的生成器表达式来实现只对 .cu 文件开启此选项 (不然给到 gcc 就出错了)；

# 这个.cu用nvcc编译就是这样的 
nvcc demo.cu --expt-relaxed-constexpr

● constexpr里面没办法调用 printf，也不能用 __syncthreads 之类的 GPU 特有的函数，因此也不能完全替代 __host__ 和 __device__；