Repository: pytorch
Issue: Patch: Fix for using `clang` to compile CUDA
Labels: module: build,module: cuda,triaged

## 🐛 Patch (bug seemed the closest category I could choose)

If you write custom kernels using the libtorch API and use `clang` rather than `nvcc` to compile the CUDA code, you receive (at least) the following compilation error:

```
In file included from bazel-out/k8-fastbuild/bin/kernels/correlation/src/correlation.cu.cc
:1:
In file included from external/libtorch/torch/include/ATen/ATen.h:5:
In file included from external/libtorch/torch/include/ATen/Context.h:4:
In file included from external/libtorch/torch/include/ATen/Type.h:2:
In file included from external/libtorch/torch/include/ATen/core/Type.h:8:
In file included from external/libtorch/torch/include/c10/core/Scalar.h:9:
In file included from external/libtorch/torch/include/c10/core/ScalarType.h:4:
external/libtorch/torch/include/c10/util/Half.h:53:12: error: no matching function for call to '__uint_as_float'
    return __uint_as_float((unsigned int)w);
           ^~~~~~~~~~~~~~~
external/llvm_toolchain/lib/clang/9.0.0/include/__clang_cuda_device_functions.h:678:18: note: candidate function not viable: call to __device__ function from __host__ function
__DEVICE__ float __uint_as_float(unsigned int __a) {
                 ^
In file included from bazel-out/k8-fastbuild/bin/kernels/correlation/src/correlation.cu.cc:1:
In file included from external/libtorch/torch/include/ATen/ATen.h:5:
In file included from external/libtorch/torch/include/ATen/Context.h:4:
In file included from external/libtorch/torch/include/ATen/Type.h:2:
In file included from external/libtorch/torch/include/ATen/core/Type.h:8:
In file included from external/libtorch/torch/include/c10/core/Scalar.h:9:
In file included from external/libtorch/torch/include/c10/core/ScalarType.h:4:
external/libtorch/torch/include/c10/util/Half.h:69:22: error: no matching function for call to '__float_as_uint'
    return (uint32_t)__float_as_uint(f);
                     ^~~~~~~~~~~~~~~
external/llvm_toolchain/lib/clang/9.0.0/include/__clang_cuda_device_functions.h:282:25: note: candidate function not viable: call to __device__ function from __host__ function
__DEVICE__ unsigned int __float_as_uint(float __a) {
```

AFAICT, this is probably due to [clang's CUDA compilation model](http://llvm.org/docs/CompileCudaWithLLVM.html#compilation-models) possibly combined with different inlining behaviour between `nvcc` and `clang`. Note in particular that clang requires "__device__ code is parsed and must be semantically correct, even though we’re not generating code for the device at this time" during the `__host__` compilation, when the `__CUDA_ARCH__` macro will not be defined.

## Proposed (possibly incomplete) solution

For my immediate purposes, this patch file sufficed:

```
--- torch/include/c10/util/Half.h	2019-07-16 13:12:24.000000000 +0100
+++ torch/include/c10/util/Half.h	2019-10-18 17:26:58.703472487 +0100
@@ -44,0 +44,7 @@
+
+#if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__)
+#define __DEVICE_HOST_FUNCTION__ __device__ __host__
+#else
+#define __DEVICE_HOST_FUNCTION__
+#endif
@@ -49,7 +49,7 @@
-  inline float fp32_from_bits(uint32_t w) {
+  __DEVICE_HOST_FUNCTION__ inline float fp32_from_bits(uint32_t w) {
   #if defined(__OPENCL_VERSION__)
     return as_float(w);
   #elif defined(__CUDA_ARCH__)
     return __uint_as_float((unsigned int)w);
   #elif defined(__INTEL_COMPILER)
     return _castu32_f32(w);
@@ -65,7 +65,7 @@
-  inline uint32_t fp32_to_bits(float f) {
+  __DEVICE_HOST_FUNCTION__ inline uint32_t fp32_to_bits(float f) {
   #if defined(__OPENCL_VERSION__)
     return as_uint(f);
   #elif defined(__CUDA_ARCH__)
     return (uint32_t)__float_as_uint(f);
   #elif defined(__INTEL_COMPILER)
     return _castf32_u32(f);
```

## To Reproduce

Steps to reproduce the behavior:

1. Write a non-trivial CUDA kernel using libtorch API
1. Compile with `clang -x cuda ...`(http://llvm.org/docs/CompileCudaWithLLVM.html)

## Expected behavior

Everything works as if you'd compiled with `nvcc`.

## Additional context

I tried this with clang 8.0.0 and 9.0.0 with libtorch at master. This isn't a full PR because a) I haven't found time to run a full test suite to check that I've found all of the bugs and b) I'm not confident enough that my proposed solution is acceptable to the pytorch devs. Hence this Issue instead. If anybody would like to outline for me the steps I should take to create a properly tested PR for this, I'd be happy to (assuming this is something you want to support). Ideally, there should be CI compiles against clang to check for compatibility.