Using __half with NVRTC and jitify

[qavl]

Hi! I want to use fp16 data type in kernels compiled with NVRTC. However, when I try to do so, I get following errors:

warning: cuda_fp16.h: File not found
---------------------------------------------------
--- JIT compile log for ...---
---------------------------------------------------
error: identifier "__half" is undefined

Can you tell me, what is the best way to use cuda_fp16.h with jitify?

Thank you.

[benbarsdell]

To ensure cuda_fp16.h can be found you'll need to pass the CUDA Toolkit include directory as a flag like this: -I/path/to/cuda/include.

Here's a minimal example (it uses half which is equivalent to __half):

  jitify::JitCache jit_cache;
  jit_cache
      .program(R"(fp16_program
#include <cuda_fp16.h>
__device__ half value;
)",
               {}, {"-I/usr/local/cuda/include"})
      .kernel("")
      .instantiate();

[qavl]

Ok, I see. Thank you) I guess, one need to make a workaround to pass the exact cuda path at runtime to make it portable, right?

[benbarsdell]

That's right. One option would be to use an environment variable like CUDA_PATH.

[maddyscientist]

We don't need another environment variable surely, we can just use the JITIFY_OPTIONS envarg to set this at run time already?