python-3.xcudacupyhalf-precision-float

Using half precision with CuPy

I am trying to compile a simple CUDA kernel with CuPy using the half precision format provided by the cuda_fp16 header file.

My kernel looks like this:

code = r'''
extern "C" {

#include <cuda_fp16.h>

__global__ void kernel(half * const f1, half * const f2)
{
   if (blockDim.x*blockIdx.x + threadIdx.x < 12 && blockDim.y*blockIdx.y + threadIdx.y < 12)
   {
      const int ctr_0 = blockDim.x*blockIdx.x + threadIdx.x;
      const int ctr_1 = blockDim.y*blockIdx.y + threadIdx.y;
      f1[12*ctr_1 + ctr_0] = f2[12*ctr_1 + ctr_0];
   } 
}

}

I try to compile like this:

options = ('-I/path/to/cuda/include/', )

mod = cp.RawModule(code=code, options=options, backend="nvrtc", jitify=True)
func = mod.get_function("kernel")

However, this results in several compiler errors:

---------------------------------------------------
--- JIT compile log for /tmp/tmpdiqnmomv/25f306a7612419fcd799b8c90718648c2c1313ca.cubin.cu ---
---------------------------------------------------
cuda_fp16.hpp(266): error: more than one instance of overloaded function "operator++" has "C" linkage

cuda_fp16.hpp(267): error: more than one instance of overloaded function "operator--" has "C" linkage

cuda_fp16.hpp(270): error: more than one instance of overloaded function "operator+" has "C" linkage

cuda_fp16.hpp(271): error: more than one instance of overloaded function "operator-" has "C" linkage

cuda_fp16.hpp(314): error: more than one instance of overloaded function "operator+" has "C" linkage

cuda_fp16.hpp(315): error: more than one instance of overloaded function "operator-" has "C" linkage

cuda_fp16.hpp(316): error: more than one instance of overloaded function "operator*" has "C" linkage

cuda_fp16.hpp(317): error: more than one instance of overloaded function "operator/" has "C" linkage

cuda_fp16.hpp(319): error: more than one instance of overloaded function "operator+=" has "C" linkage

cuda_fp16.hpp(320): error: more than one instance of overloaded function "operator-=" has "C" linkage

cuda_fp16.hpp(321): error: more than one instance of overloaded function "operator*=" has "C" linkage

cuda_fp16.hpp(322): error: more than one instance of overloaded function "operator/=" has "C" linkage

cuda_fp16.hpp(324): error: more than one instance of overloaded function "operator++" has "C" linkage

cuda_fp16.hpp(325): error: more than one instance of overloaded function "operator--" has "C" linkage

cuda_fp16.hpp(326): error: more than one instance of overloaded function "operator++" has "C" linkage

cuda_fp16.hpp(327): error: more than one instance of overloaded function "operator--" has "C" linkage

cuda_fp16.hpp(329): error: more than one instance of overloaded function "operator+" has "C" linkage

cuda_fp16.hpp(330): error: more than one instance of overloaded function "operator-" has "C" linkage

cuda_fp16.hpp(332): error: more than one instance of overloaded function "operator==" has "C" linkage

cuda_fp16.hpp(333): error: more than one instance of overloaded function "operator!=" has "C" linkage

cuda_fp16.hpp(334): error: more than one instance of overloaded function "operator>" has "C" linkage

cuda_fp16.hpp(335): error: more than one instance of overloaded function "operator<" has "C" linkage

cuda_fp16.hpp(336): error: more than one instance of overloaded function "operator>=" has "C" linkage

cuda_fp16.hpp(337): error: more than one instance of overloaded function "operator<=" has "C" linkage

24 errors detected in the compilation of "/tmp/tmpdiqnmomv/25f306a7612419fcd799b8c90718648c2c1313ca.cubin.cu".

Is there anything obvious I am missing?

I am using cupy-cuda11x and cuda 11.2

Solution

  • The error message is pretty clear -- the compiler is telling you that cuda_fp16.hpp contains features which aren't supported by C linkage. In this case function overloading, it appears.

    I would expect something like this should compile correctly:

    #include <cuda_fp16.h>

extern "C" 
__global__ void kernel(half * const f1, half * const f2)
{
   if (blockDim.x*blockIdx.x + threadIdx.x < 12 && blockDim.y*blockIdx.y + threadIdx.y < 12)
   {
      const int ctr_0 = blockDim.x*blockIdx.x + threadIdx.x;
      const int ctr_1 = blockDim.y*blockIdx.y + threadIdx.y;
      f1[12*ctr_1 + ctr_0] = f2[12*ctr_1 + ctr_0];
   } 
}

    i.e. your function has C linkage, but the included header does not.