Index: External/CUDA/CMakeLists.txt
===================================================================
--- External/CUDA/CMakeLists.txt
+++ External/CUDA/CMakeLists.txt
@@ -98,6 +98,7 @@
   create_one_local_test(empty empty.cu)
   create_one_local_test(printf printf.cu)
   create_one_local_test(future future.cu)
+  create_one_local_test(builtin_var builtin_var.cu)
   # We only need SIMD tests on CUDA-8.0 to verivy that our reference is correct
   # and matches NVIDIA-provided one.  and on CUDA-9.2 to verify that clang's
   # implementation matches the reference.  This test also happens to be the
Index: External/CUDA/builtin_var.cu
===================================================================
--- /dev/null
+++ External/CUDA/builtin_var.cu
@@ -0,0 +1,20 @@
+__global__ void kernel() {
+  // Verify that *Idx/*Dim can be assigned to uint3/dim3.
+  uint3 thread_idx = threadIdx;
+  uint3 block_idx = blockIdx;
+  dim3 block_dim = blockDim;
+  dim3 grid_dim = gridDim;
+
+  // And that they can be converted to uint3/dim3
+  dim3 thread_idx_dim = threadIdx;
+  dim3 block_idx_dim = blockIdx;
+  uint3 block_dim_uint = blockDim;
+  uint3 grid_dim_uint = gridDim;
+}
+
+int main(int argc, char* argv[]) {
+  kernel<<<2, 2>>>();
+  cudaDeviceSynchronize();
+  cudaDeviceReset();
+  return 0;
+}