Move all CUDA testing inputs to Inputs/ subdirectory inside the tests.

llvm-svn: 207453

clang/test/CodeGenCUDA/address-spaces.cu

@@ -3,7 +3,7 @@
 // Verifies Clang emits correct address spaces and addrspacecast instructions
 // for CUDA code.
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // CHECK: @i = addrspace(1) global
 __device__ int i;

clang/test/CodeGenCUDA/device-stub.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // Test that we build the correct number of calls to cudaSetupArgument followed
 // by a call to cudaLaunch.

clang/test/CodeGenCUDA/filter-decl.cu

@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // CHECK-HOST-NOT: constantdata = global
 // CHECK-DEVICE: constantdata = global

clang/test/CodeGenCUDA/kernel-call.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -emit-llvm %s -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 __global__ void g1(int x) {}
 

clang/test/CodeGenCUDA/launch-bounds.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 #define MAX_THREADS_PER_BLOCK 256
 #define MIN_BLOCKS_PER_MP     2

clang/test/CodeGenCUDA/ptx-kernels.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 // CHECK-LABEL: define void @device_function
 extern "C"

clang/test/PCH/Inputs/cuda.h

@@ -0,0 +1,20 @@
+/* Minimal declarations for CUDA support.  Testing purposes only. */
+
+#include <stddef.h>
+
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+typedef struct cudaStream *cudaStream_t;
+
+int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                      cudaStream_t stream = 0);

clang/test/PCH/cuda-kernel-call.cu

@@ -5,7 +5,7 @@
 #define HEADER
 // Header.
 
-#include "../SemaCUDA/cuda.h"
+#include "Inputs/cuda.h"
 
 void kcall(void (*kp)()) {
   kp<<<1, 1>>>();

clang/test/SemaCUDA/Inputs/cuda.h

@@ -0,0 +1,20 @@
+/* Minimal declarations for CUDA support.  Testing purposes only. */
+
+#include <stddef.h>
+
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+typedef struct cudaStream *cudaStream_t;
+
+int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
+                      cudaStream_t stream = 0);

clang/test/SemaCUDA/function-target.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 
-#include "cuda.h"
+#include "Inputs/cuda.h"
 
 __host__ void h1h(void);
 __device__ void h1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ function}}

clang/test/SemaCUDA/kernel-call.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 
-#include "cuda.h"
+#include "Inputs/cuda.h"
 
 __global__ void g1(int x) {}
 

clang/test/SemaCUDA/launch_bounds.cu

@@ -1,14 +1,14 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-
-#include "cuda.h"
-
-__launch_bounds__(128, 7) void Test1(void);
-__launch_bounds__(128) void Test2(void);
-
-__launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
-
-// FIXME: the error should read that the attribute takes exactly one or two arguments, but there
-// is no support for such a diagnostic currently.
-__launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
-
-int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+__launch_bounds__(128, 7) void Test1(void);
+__launch_bounds__(128) void Test2(void);
+
+__launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
+
+// FIXME: the error should read that the attribute takes exactly one or two arguments, but there
+// is no support for such a diagnostic currently.
+__launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
+
+int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}

clang/test/SemaCUDA/qualifiers.cu

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 
-#include "cuda.h"
+#include "Inputs/cuda.h"
 
 __global__ void g1(int x) {}
 __global__ int g2(int x) { // expected-error {{must have void return type}}