Remove __host__ annotation when compiling for NVRTC

Author: stijnh

include/kernel_float/macros.h

@@ -10,11 +10,16 @@
     #define KERNEL_FLOAT_IS_CUDA (1)
     #define KERNEL_FLOAT_DEVICE    __forceinline__ __device__
 
-    #ifdef __CUDA_ARCH__
+    // NVRTC cannot deal with __host__ annotations
+    #ifdef __CUDACC_RTC__
         #define KERNEL_FLOAT_INLINE    __forceinline__ __device__
+    #else
+        #define KERNEL_FLOAT_INLINE    __forceinline__ __host__ __device__
+    #endif
+
+    #ifdef __CUDA_ARCH__
         #define KERNEL_FLOAT_IS_DEVICE (1)
     #else  // __CUDA_ARCH__
-        #define KERNEL_FLOAT_INLINE  __forceinline__ __host__
         #define KERNEL_FLOAT_IS_HOST (1)
     #endif  // __CUDA_ARCH__
 #elif defined(__HIPCC__)

single_include/kernel_float.h

@@ -16,8 +16,8 @@
 
 //================================================================================
 // this file has been auto-generated, do not modify its contents!
-// date: 2025-09-15 12:39:02.709972
-// git hash: 81efb0fbfcc587343ba798590ba1541babe378d6
+// date: 2025-09-15 12:44:05.768243
+// git hash: 9b41485a27b669ea6f4aae118b4d251947608bf6
 //================================================================================
 
 #ifndef KERNEL_FLOAT_MACROS_H
@@ -32,11 +32,16 @@
     #define KERNEL_FLOAT_IS_CUDA (1)
     #define KERNEL_FLOAT_DEVICE    __forceinline__ __device__
 
-    #ifdef __CUDA_ARCH__
+    // NVRTC cannot deal with __host__ annotations
+    #ifdef __CUDACC_RTC__
         #define KERNEL_FLOAT_INLINE    __forceinline__ __device__
+    #else
+        #define KERNEL_FLOAT_INLINE    __forceinline__ __host__ __device__
+    #endif
+
+    #ifdef __CUDA_ARCH__
         #define KERNEL_FLOAT_IS_DEVICE (1)
     #else  // __CUDA_ARCH__
-        #define KERNEL_FLOAT_INLINE  __forceinline__ __host__
         #define KERNEL_FLOAT_IS_HOST (1)
     #endif  // __CUDA_ARCH__
 #elif defined(__HIPCC__)