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
+}