[test-suite, CUDA] Compile-time test for builtin variables.
diff --git a/External/CUDA/CMakeLists.txt b/External/CUDA/CMakeLists.txt
index 2c0b710..743f72e 100644
--- a/External/CUDA/CMakeLists.txt
+++ b/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
diff --git a/External/CUDA/builtin_var.cu b/External/CUDA/builtin_var.cu
new file mode 100644
index 0000000..78d1447
--- /dev/null
+++ b/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;
+}
