[test-suite, CUDA] Compile-time test for builtin variables.

[tra]

Landed in 93b9e85695e583d1c577b40

External/CUDA/CMakeLists.txt

macro(create_local_cuda_tests VariantSuffix)
   create_one_local_test(algorithm algorithm.cu)
   create_one_local_test(cmath cmath.cu)
   create_one_local_test(complex complex.cu)
   create_one_local_test(math_h math_h.cu)
   create_one_local_test(new new.cu)
   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
   # longest one, so by not running unnecessary instances we speed up cuda
   # buildbot a lot.
   create_one_local_test_f(simd simd.cu
                           "cuda-(8[.]0|9[.]2)-c[+][+]11-libc[+][+]")
   create_one_local_test(test_round test_round.cu)

External/CUDA/builtin_var.cu

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