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
カテゴリー:CUDA
記事作成日:2019-12-15