Hi eliben, echristo, rnk,
The patch implements support for [[
http://docs.nvidia.com/cuda/cuda-c-programming-guide/#built-in-variables|
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.
http://reviews.llvm.org/D9064
Files:
lib/Headers/CMakeLists.txt
lib/Headers/cuda/cuda_builtin_vars.h
test/CodeGenCUDA/cuda-builtin-vars.cu
EMAIL PREFERENCES
http://reviews.llvm.org/settings/panel/emailpreferences/
Index: lib/Headers/CMakeLists.txt
===================================================================
--- lib/Headers/CMakeLists.txt
+++ lib/Headers/CMakeLists.txt
@@ -13,6 +13,7 @@
bmi2intrin.h
bmiintrin.h
cpuid.h
+ cuda/cuda_builtin_vars.h
emmintrin.h
f16cintrin.h
float.h
Index: lib/Headers/cuda/cuda_builtin_vars.h
===================================================================
--- /dev/null
+++ lib/Headers/cuda/cuda_builtin_vars.h
@@ -0,0 +1,53 @@
+#ifndef __CUDA_BUILTIN_VARS_H__
+#define __CUDA_BUILTIN_VARS_H__
+
+#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC) \
+ __declspec(property(get = __fetch_builtin_##FIELD)) int FIELD; \
+ static inline __attribute__((always_inline)) \
+ __attribute__((device)) int __fetch_builtin_##FIELD(void) { \
+ return INTRINSIC; \
+ }
+
+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_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_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_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_builtin_gridDim_t() {}
+};
+
+#define __CUDA_BUILTIN_VAR extern const __attribute__((device))
+
+__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;
+__CUDA_BUILTIN_VAR int warpSize = 32;
+
+#undef __CUDA_DEVICE_BUILTIN
+#undef __CUDA_BUILTIN_VAR
+#endif
Index: test/CodeGenCUDA/cuda-builtin-vars.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/cuda-builtin-vars.cu
@@ -0,0 +1,28 @@
+// 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-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
