[CUDA] Support for built-in cuda variables (threadIdx and its friends).

[tra]

Summary

The patch implements support for built-in CUDA variables using __declattr((property)) attribute to convert all built-in variable reads into appropriate llvm builtin calls.

This makes it possible to compile typical CUDA source code that relies on built-in variables.

[eliben]

lib/Headers/cuda/cuda_builtin_vars.h

3 ↗

(On Diff #23887)

A top-level comment on the file that explains what this does with a usage example?

12 ↗

(On Diff #23887)

Arent blockIdx and threadIdx unsigned?

[tra]

Made fields unsigned. Added copyright header and comments.

[tra]

lib/Headers/cuda/cuda_builtin_vars.h

3 ↗

(On Diff #23887)

Done.

12 ↗

(On Diff #23887)

Yes. So are blockDim and gridDim as dim3 is based on uint3. Fixed.

[tra]

Really changed the fields to be unsigned.

[rsmith]

lib/Headers/cuda/cuda_builtin_vars.h

89 ↗

(On Diff #23976)

This is a (strong) definition due to the extern in the __CUDA_BUILTIN_VAR macro, and will cause a link error if this file is #included into multiple objects in the same binary. You could solve this by removing the extern (that'd give the variable internal linkage) or by using enum { warpSize = 32 }; or similar. I'm not sure exactly what the CUDA spec requires here.

11 ↗

(On Diff #23887)

Should these types also have deleted copy constructors and copy assignment operators? I'd imagine these should both be ill-formed:

auto x = threadIdx;
threadIdx = threadIdx;

[tra]

• Restricted access to built-in variables.
• Implemented warpSize as class subject to the same restrictions as other built-in variables.
• Made all built-in variables extern weak in case we ever emit any references to them.
• Added a test case to check diagnostics on attempt to construct or improperly access a built-in variable.

lib/Headers/cuda/cuda_builtin_vars.h

89 ↗

(On Diff #23976)

Those have to be extern because no instances of those special variables can ever exist, yet they are referred to as if they do. 'extern' matches this functionality best. I've added a weak attribute to make linking work. That is, if/when we'll get to do any linking on the GPU code. Currently we produce PTX assembly and we don't ever link one PTX file with another.

I've reworked warpSize into a class with the same restrictions on construction and access as the other built-in variables.

As for CUDA spec, it does not say much -- 'They [built-in variables] are only valid within functions that are executed on the device. [warpSize] is of type int and contains the warp size in threads.'

11 ↗

(On Diff #23887)

I've disabled constructors for the special variables and also disabled address-of operator as nvcc does not allow taking their address.

[rsmith]

lib/Headers/cuda/cuda_builtin_vars.h

89 ↗

(On Diff #23976)

I was only talking about warpSize, not the others. The others seem OK as (non-weak) extern declarations, on the assumption that programs aren't actually allowed to use the address of these.

Your new approach doesn't satisfy the requirement that "warpSize is of type int": this is detectable through template argument deduction, for instance:

template<typename T> T max(T, T)
max(warpSize, 10); // error, T is ambiguous, could be int or __cuda_builtin_warpSize_t

Using

__attribute__((device)) const int warpSize = 32;

would probably work (that is, drop the extern from what you had before).

[tra]

Reverted warpSize back to plain old int. It's 'static const int' now.

lib/Headers/cuda/cuda_builtin_vars.h

89 ↗

(On Diff #23976)

Another downside of warpSize-in-a-class was that a reference to an instance of that class was required for the type conversion operator as it can't be a static member function. Scratch that.

Dropping extern is not sufficient. 'const int' suffers from the same problem as 'extern' -- we end up with warpSize visible externally which causes linking error.

I could use 'static const int' -- it keeps the symbol local and does not create linking issues. Unfortunately it allows one to take address of the variable. It's not perfect, but we can live with that for now.

[rsmith]

LGTM

lib/Headers/cuda/cuda_builtin_vars.h

89 ↗

(On Diff #23976)

In C++, const on a global variable definition (that is not explicitly marked extern) implies internal linkage just like static does -- the static is harmless but redundant.

[tra]

lib/Headers/cuda/cuda_builtin_vars.h

89 ↗

(On Diff #23976)

Thanks. I was unaware of that. Will fix.

cfe/trunk/lib/Headers/CMakeLists.txt

 set(files
   adxintrin.h
   altivec.h
   ammintrin.h
   arm_acle.h
   avx2intrin.h
   avx512bwintrin.h
   avx512erintrin.h
   avx512fintrin.h
   avx512vlbwintrin.h
   avx512vlintrin.h
   avxintrin.h
   bmi2intrin.h
   bmiintrin.h
   cpuid.h
+  cuda/cuda_builtin_vars.h
   emmintrin.h
   f16cintrin.h
   float.h
   fma4intrin.h
   fmaintrin.h
   htmintrin.h
   htmxlintrin.h
   ia32intrin.h

cfe/trunk/lib/Headers/cuda/cuda_builtin_vars.h

+/*===---- cuda_builtin_vars.h - CUDA built-in variables ---------------------===
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __CUDA_BUILTIN_VARS_H
+#define __CUDA_BUILTIN_VARS_H
+
+// The file implements built-in CUDA variables using __declspec(property).
+// https://msdn.microsoft.com/en-us/library/yhfk0thd.aspx
+// All read accesses of built-in variable fields get converted into calls to a
+// getter function which in turn would call appropriate builtin to fetch the
+// value.
+//
+// Example:
+//    int x = threadIdx.x;
+// IR output:
+//  %0 = call i32 @llvm.ptx.read.tid.x() #3
+// PTX output:
+//  mov.u32     %r2, %tid.x;
+
+#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC)                                \
+  __declspec(property(get = __fetch_builtin_##FIELD)) unsigned int FIELD;      \
+  static inline __attribute__((always_inline))                                 \
+      __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) {     \
+    return INTRINSIC;                                                          \
+  }
+
+#if __cplusplus >= 201103L
+#define __DELETE =delete
+#else
+#define __DELETE
+#endif
+
+// Make sure nobody can create instances of the special varible types.  nvcc
+// also disallows taking address of special variables, so we disable address-of
+// operator as well.
+#define __CUDA_DISALLOW_BUILTINVAR_ACCESS(TypeName)                            \
+  __attribute__((device)) TypeName() __DELETE;                                 \
+  __attribute__((device)) TypeName(const TypeName &) __DELETE;                 \
+  __attribute__((device)) void operator=(const TypeName &) const __DELETE;     \
+  __attribute__((device)) TypeName *operator&() const __DELETE
+
+struct __cuda_builtin_threadIdx_t {
+  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_tid_x());
+  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_tid_y());
+  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_tid_z());
+private:
+  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_threadIdx_t);
+};
+
+struct __cuda_builtin_blockIdx_t {
+  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_ctaid_x());
+  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_ctaid_y());
+  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_ctaid_z());
+private:
+  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockIdx_t);
+};
+
+struct __cuda_builtin_blockDim_t {
+  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_ntid_x());
+  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_ntid_y());
+  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_ntid_z());
+private:
+  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockDim_t);
+};
+
+struct __cuda_builtin_gridDim_t {
+  __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_nctaid_x());
+  __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_nctaid_y());
+  __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_nctaid_z());
+private:
+  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_gridDim_t);
+};
+
+#define __CUDA_BUILTIN_VAR                                                     \
+  extern const __attribute__((device)) __attribute__((weak))
+__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx;
+__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx;
+__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim;
+__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim;
+
+// warpSize should translate to read of %WARP_SZ but there's currently no
+// builtin to do so. According to PTX v4.2 docs 'to date, all target
+// architectures have a WARP_SZ value of 32'.
+__attribute__((device)) const int warpSize = 32;
+
+#undef __CUDA_DEVICE_BUILTIN
+#undef __CUDA_BUILTIN_VAR
+#undef __CUDA_DISALLOW_BUILTINVAR_ACCESS
+
+#endif /* __CUDA_BUILTIN_VARS_H */

cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu

+// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+
+#include "cuda/cuda_builtin_vars.h"
+
+// CHECK: define void @_Z6kernelPi(i32* %out)
+__attribute__((global))
+void kernel(int *out) {
+  int i = 0;
+  out[i++] = threadIdx.x; // CHECK: call i32 @llvm.ptx.read.tid.x()
+  out[i++] = threadIdx.y; // CHECK: call i32 @llvm.ptx.read.tid.y()
+  out[i++] = threadIdx.z; // CHECK: call i32 @llvm.ptx.read.tid.z()
+
+  out[i++] = blockIdx.x; // CHECK: call i32 @llvm.ptx.read.ctaid.x()
+  out[i++] = blockIdx.y; // CHECK: call i32 @llvm.ptx.read.ctaid.y()
+  out[i++] = blockIdx.z; // CHECK: call i32 @llvm.ptx.read.ctaid.z()
+
+  out[i++] = blockDim.x; // CHECK: call i32 @llvm.ptx.read.ntid.x()
+  out[i++] = blockDim.y; // CHECK: call i32 @llvm.ptx.read.ntid.y()
+  out[i++] = blockDim.z; // CHECK: call i32 @llvm.ptx.read.ntid.z()
+
+  out[i++] = gridDim.x; // CHECK: call i32 @llvm.ptx.read.nctaid.x()
+  out[i++] = gridDim.y; // CHECK: call i32 @llvm.ptx.read.nctaid.y()
+  out[i++] = gridDim.z; // CHECK: call i32 @llvm.ptx.read.nctaid.z()
+
+  out[i++] = warpSize; // CHECK: store i32 32,
+
+  // CHECK: ret void
+}

cfe/trunk/test/SemaCUDA/cuda-builtin-vars.cu

+// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s
+
+#include "cuda/cuda_builtin_vars.h"
+__attribute__((global))
+void kernel(int *out) {
+  int i = 0;
+  out[i++] = threadIdx.x;
+  threadIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = threadIdx.y;
+  threadIdx.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = threadIdx.z;
+  threadIdx.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = blockIdx.x;
+  blockIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = blockIdx.y;
+  blockIdx.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = blockIdx.z;
+  blockIdx.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = blockDim.x;
+  blockDim.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = blockDim.y;
+  blockDim.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = blockDim.z;
+  blockDim.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = gridDim.x;
+  gridDim.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = gridDim.y;
+  gridDim.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = gridDim.z;
+  gridDim.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = warpSize;
+  warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}}
+  // expected-note@cuda/cuda_builtin_vars.h:104 {{variable 'warpSize' declared const here}}
+
+  // Make sure we can't construct or assign to the special variables.
+  __cuda_builtin_threadIdx_t x; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda/cuda_builtin_vars.h:67 {{declared private here}}
+
+  __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda/cuda_builtin_vars.h:67 {{declared private here}}
+
+  threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda/cuda_builtin_vars.h:67 {{declared private here}}
+
+  void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda/cuda_builtin_vars.h:67 {{declared private here}}
+
+  // Following line should've caused an error as one is not allowed to
+  // take address of a built-in variable in CUDA. Alas there's no way
+  // to prevent getting address of a 'const int', so the line
+  // currently compiles without errors or warnings.
+  const void *wsptr = &warpSize;
+}