8 files changed, 268 insertions, 45 deletions
diff --git a/test/SemaCUDA/asm-constraints-device.cu b/test/SemaCUDA/asm-constraints-device.cu
new file mode 100644
index 000000000000..cdd2d32a3068
--- /dev/null
+++ b/test/SemaCUDA/asm-constraints-device.cu
@@ -0,0 +1,24 @@
+// Verify that we do check for constraints in device-side inline
+// assembly. Passing an illegal input/output constraint and look 
+// for corresponding error
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify %s
+
+__attribute__((device)) void df() {
+  short h;
+  int a;
+  // asm with PTX constraints. Some of them are PTX-specific.
+  __asm__("output constraints"
+          : "=h"(h), // .u16 reg, OK
+            "=a"(a)  // expected-error {{invalid output constraint '=a' in asm}}
+          :          // None
+          );
+  __asm__("input constraints"
+          :           // None
+          : "f"(0.0), // .f32 reg, OK
+            "d"(0.0), // .f64 reg, OK
+            "h"(0),   // .u16 reg, OK
+            "r"(0),   // .u32 reg, OK
+            "l"(0),   // .u64 reg, OK
+            "a"(0)    // expected-error {{invalid input constraint 'a' in asm}}
+          );
+}
diff --git a/test/SemaCUDA/asm-constraints-mixed.cu b/test/SemaCUDA/asm-constraints-mixed.cu
new file mode 100644
index 000000000000..a4ac9c65c99f
--- /dev/null
+++ b/test/SemaCUDA/asm-constraints-mixed.cu
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// expected-no-diagnostics
+
+__attribute__((device)) void df() {
+  short h;
+  // asm with PTX constraints. Some of them are PTX-specific.
+  __asm__("dont care" : "=h"(h): "f"(0.0), "d"(0.0), "h"(0), "r"(0), "l"(0));
+}
+
+void hf() {
+  int a;
+  // Asm with x86 constraints that are not supported by PTX.
+  __asm__("dont care" : "=a"(a): "a"(0), "b"(0), "c"(0));
+}
diff --git a/test/SemaCUDA/cuda-builtin-vars.cu b/test/SemaCUDA/cuda-builtin-vars.cu
new file mode 100644
index 000000000000..97c5111cebdc
--- /dev/null
+++ b/test/SemaCUDA/cuda-builtin-vars.cu
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s
+
+#include "cuda_builtin_vars.h"
+__attribute__((global))
+void kernel(int *out) {
+  int i = 0;
+  out[i++] = threadIdx.x;
+  threadIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = threadIdx.y;
+  threadIdx.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = threadIdx.z;
+  threadIdx.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = blockIdx.x;
+  blockIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = blockIdx.y;
+  blockIdx.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = blockIdx.z;
+  blockIdx.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = blockDim.x;
+  blockDim.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = blockDim.y;
+  blockDim.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = blockDim.z;
+  blockDim.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = gridDim.x;
+  gridDim.x = 0; // expected-error {{no setter defined for property 'x'}}
+  out[i++] = gridDim.y;
+  gridDim.y = 0; // expected-error {{no setter defined for property 'y'}}
+  out[i++] = gridDim.z;
+  gridDim.z = 0; // expected-error {{no setter defined for property 'z'}}
+
+  out[i++] = warpSize;
+  warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}}
+  // expected-note@cuda_builtin_vars.h:104 {{variable 'warpSize' declared const here}}
+
+  // Make sure we can't construct or assign to the special variables.
+  __cuda_builtin_threadIdx_t x; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda_builtin_vars.h:67 {{declared private here}}
+
+  __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda_builtin_vars.h:67 {{declared private here}}
+
+  threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda_builtin_vars.h:67 {{declared private here}}
+
+  void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_threadIdx_t'}}
+  // expected-note@cuda_builtin_vars.h:67 {{declared private here}}
+
+  // Following line should've caused an error as one is not allowed to
+  // take address of a built-in variable in CUDA. Alas there's no way
+  // to prevent getting address of a 'const int', so the line
+  // currently compiles without errors or warnings.
+  const void *wsptr = &warpSize;
+}
diff --git a/test/SemaCUDA/function-target-disabled-check.cu b/test/SemaCUDA/function-target-disabled-check.cu
new file mode 100644
index 000000000000..979d4edbf892
--- /dev/null
+++ b/test/SemaCUDA/function-target-disabled-check.cu
@@ -0,0 +1,26 @@
+// Test that we can disable cross-target call checks in Sema with the
+// -fcuda-disable-target-call-checks flag. Without this flag we'd get a bunch
+// of errors here, since there are invalid cross-target calls present.
+
+// RUN: %clang_cc1 -fsyntax-only -verify %s -fcuda-disable-target-call-checks
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -fcuda-disable-target-call-checks
+
+// expected-no-diagnostics
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+
+__attribute__((host)) void h1();
+
+__attribute__((device)) void d1() {
+  h1();
+}
+
+__attribute__((host)) void h2() {
+  d1();
+}
+
+__attribute__((global)) void g1() {
+  h2();
+}
diff --git a/test/SemaCUDA/function-target-hd.cu b/test/SemaCUDA/function-target-hd.cu
new file mode 100644
index 000000000000..25fcc6e9188f
--- /dev/null
+++ b/test/SemaCUDA/function-target-hd.cu
@@ -0,0 +1,71 @@
+// Test the Sema analysis of caller-callee relationships of host device
+// functions when compiling CUDA code. There are 4 permutations of this test as
+// host and device compilation are separate compilation passes, and clang has
+// an option to allow host calls from host device functions. __CUDA_ARCH__ is
+// defined when compiling for the device and TEST_WARN_HD when host calls are
+// allowed from host device functions. So for example, if __CUDA_ARCH__ is
+// defined and TEST_WARN_HD is not then device compilation is happening but
+// host device functions are not allowed to call device functions.
+
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD
+
+#include "Inputs/cuda.h"
+
+__host__ void hd1h(void);
+#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
+// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+__device__ void hd1d(void);
+#ifndef __CUDA_ARCH__
+// expected-note@-2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#endif
+__host__ void hd1hg(void);
+__device__ void hd1dg(void);
+#ifdef __CUDA_ARCH__
+__host__ void hd1hig(void);
+#if !defined(TEST_WARN_HD)
+// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+#else
+__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#endif
+__host__ __device__ void hd1hd(void);
+__global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
+
+__host__ __device__ void hd1(void) {
+#if defined(TEST_WARN_HD) && defined(__CUDA_ARCH__)
+// expected-warning@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}}
+// expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}}
+#endif
+  hd1d();
+#ifndef __CUDA_ARCH__
+// expected-error@-2 {{no matching function}}
+#endif
+  hd1h();
+#if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD)
+// expected-error@-2 {{no matching function}}
+#endif
+
+  // No errors as guarded
+#ifdef __CUDA_ARCH__
+  hd1d();
+#else
+  hd1h();
+#endif
+
+  // Errors as incorrectly guarded
+#ifndef __CUDA_ARCH__
+  hd1dig(); // expected-error {{no matching function}}
+#else
+  hd1hig();
+#ifndef TEST_WARN_HD
+// expected-error@-2 {{no matching function}}
+#endif
+#endif
+
+  hd1hd();
+  hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
+}
diff --git a/test/SemaCUDA/function-target.cu b/test/SemaCUDA/function-target.cu
index ca56030309de..83dce50b4af8 100644
--- a/test/SemaCUDA/function-target.cu
+++ b/test/SemaCUDA/function-target.cu
@@ -31,41 +31,3 @@ __device__ void d1(void) {
   d1hd();
   d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
 }
-
-// Expected 0-1 as in one of host/device side compilation it is an error, while
-// not in the other
-__host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-__device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
-__host__ void hd1hg(void);
-__device__ void hd1dg(void);
-#ifdef __CUDA_ARCH__
-__host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-#else
-__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
-#endif
-__host__ __device__ void hd1hd(void);
-__global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
-
-__host__ __device__ void hd1(void) {
-  // Expected 0-1 as in one of host/device side compilation it is an error,
-  // while not in the other
-  hd1d(); // expected-error 0-1 {{no matching function}}
-  hd1h(); // expected-error 0-1 {{no matching function}}
-
-  // No errors as guarded
-#ifdef __CUDA_ARCH__
-  hd1d();
-#else
-  hd1h();
-#endif
-
-  // Errors as incorrectly guarded
-#ifndef __CUDA_ARCH__
-  hd1dig(); // expected-error {{no matching function}}
-#else
-  hd1hig(); // expected-error {{no matching function}}
-#endif
-
-  hd1hd();
-  hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}
-}
diff --git a/test/SemaCUDA/launch_bounds.cu b/test/SemaCUDA/launch_bounds.cu
index 8edc41b6ce91..468954a3aab1 100644
--- a/test/SemaCUDA/launch_bounds.cu
+++ b/test/SemaCUDA/launch_bounds.cu
@@ -1,11 +1,49 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
 
 #include "Inputs/cuda.h"
 
-__launch_bounds__(128, 7) void Test1(void);
-__launch_bounds__(128) void Test2(void);
+__launch_bounds__(128, 7) void Test2Args(void);
+__launch_bounds__(128) void Test1Arg(void);
 
-__launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
-__launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
+__launch_bounds__(0xffffffff) void TestMaxArg(void);
+__launch_bounds__(0x100000000) void TestTooBigArg(void); // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__launch_bounds__(0x10000000000000000) void TestWayTooBigArg(void); // expected-error {{integer literal is too large to be represented in any integer type}}
 
-int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
+__launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
+__launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
+
+__launch_bounds__(1, 2, 3) void Test3Args(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
+__launch_bounds__() void TestNoArgs(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
+
+int TestNoFunction __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
+
+__launch_bounds__(true) void TestBool(void);
+__launch_bounds__(128.0) void TestFP(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
+__launch_bounds__((void*)0) void TestNullptr(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
+
+int nonconstint = 256;
+__launch_bounds__(nonconstint) void TestNonConstInt(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
+
+const int constint = 512;
+__launch_bounds__(constint) void TestConstInt(void);
+__launch_bounds__(constint * 2 + 3) void TestConstIntExpr(void);
+
+template <int a, int b> __launch_bounds__(a, b) void TestTemplate2Args(void) {}
+template void TestTemplate2Args<128,7>(void);
+
+template <int a> __launch_bounds__(a) void TestTemplate1Arg(void) {}
+template void TestTemplate1Arg<128>(void);
+
+template <class a>
+__launch_bounds__(a) void TestTemplate1ArgClass(void) {} // expected-error {{'a' does not refer to a value}}
+// expected-note@-2 {{declared here}}
+
+template <int a, int b, int c>
+__launch_bounds__(a + b, c + constint) void TestTemplateExpr(void) {}
+template void TestTemplateExpr<128+constint, 3, 7>(void);
+
+template <int... Args>
+__launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
+
+template <int... Args>
+__launch_bounds__(1, Args) void TestTemplateVariadicArgs2(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
diff --git a/test/SemaCUDA/qualifiers.cu b/test/SemaCUDA/qualifiers.cu
index 42a80b8b38c7..4be850586fbf 100644
--- a/test/SemaCUDA/qualifiers.cu
+++ b/test/SemaCUDA/qualifiers.cu
@@ -1,7 +1,37 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify -fcuda-is-device %s
+//
+// We run clang_cc1 with 'not' because source file contains
+// intentional errors. CC1 failure is expected and must be ignored
+// here. We're interested in what ends up in AST and that's what
+// FileCheck verifies.
+// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -ast-dump %s \
+// RUN:   | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-HOST
+// RUN: not %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -ast-dump -fcuda-is-device %s \
+// RUN:   | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-DEVICE
 
 #include "Inputs/cuda.h"
 
+// Host (x86) supports TLS and device-side compilation should ignore
+// host variables. No errors in either case.
+int __thread host_tls_var;
+// CHECK-ALL: host_tls_var 'int' tls
+
+#if defined(__CUDA_ARCH__)
+// NVPTX does not support TLS
+__device__ int __thread device_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
+// CHECK-DEVICE: device_tls_var 'int' tls
+__shared__ int __thread shared_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
+// CHECK-DEVICE: shared_tls_var 'int' tls
+#else
+// Device-side vars should not produce any errors during host-side
+// compilation.
+__device__ int __thread device_tls_var;
+// CHECK-HOST: device_tls_var 'int' tls
+__shared__ int __thread shared_tls_var;
+// CHECK-HOST: shared_tls_var 'int' tls
+#endif
+
 __global__ void g1(int x) {}
 __global__ int g2(int x) { // expected-error {{must have void return type}}
   return 1;