hostのconstexprな関数をdeviceで使う

CUDA Advent Calendar 2019 15日目の記事です．
今日も小ネタ

例

たとえばこんなコードがあったとしましょう． これを

nvcc text.cu

とするとhost用の関数であるadd_2をdeviceから呼んでいるわけですから当然怒られます．

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

しかし，ここにも書かれている通り試験的にではありますが--expt-relaxed-constexprをつけることでconstexprのhostの関数をdeviceで使えるようになります．
逆も然りで，__device__修飾したconsexpr関数をhost側から呼ぶこともできます．

何が嬉しいの?

世の中はconstexpr関数で満ち溢れています．
誰かがCUDAのカーネル関数内で使うことを想定せずに書いたconstexpr関数をkernel関数内で使えるようになります．
nvccがC++17に対応してくれれば，例えばVecPP/cste_mathとかが使えるようになり便利そうですね．

参考

Constexpr functions and function templates