From: Eli Bendersky Date: Mon, 28 Apr 2014 22:21:28 +0000 (+0000) Subject: Move all CUDA testing inputs to Inputs/ subdirectory inside the tests. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=495bb269c08f513c0561f3bc5b5666fde2785526;p=clang Move all CUDA testing inputs to Inputs/ subdirectory inside the tests. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@207453 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/test/SemaCUDA/cuda.h b/test/CodeGenCUDA/Inputs/cuda.h similarity index 100% rename from test/SemaCUDA/cuda.h rename to test/CodeGenCUDA/Inputs/cuda.h diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index a28886f4b1..b80820683f 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/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; diff --git a/test/CodeGenCUDA/device-stub.cu b/test/CodeGenCUDA/device-stub.cu index af73ea993f..ed94d10848 100644 --- a/test/CodeGenCUDA/device-stub.cu +++ b/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. diff --git a/test/CodeGenCUDA/filter-decl.cu b/test/CodeGenCUDA/filter-decl.cu index 008eaaed88..faaeb69fe1 100644 --- a/test/CodeGenCUDA/filter-decl.cu +++ b/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 diff --git a/test/CodeGenCUDA/kernel-call.cu b/test/CodeGenCUDA/kernel-call.cu index f134624eec..9b849db908 100644 --- a/test/CodeGenCUDA/kernel-call.cu +++ b/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) {} diff --git a/test/CodeGenCUDA/launch-bounds.cu b/test/CodeGenCUDA/launch-bounds.cu index 3f88aeb7e7..ed4c2bfc88 100644 --- a/test/CodeGenCUDA/launch-bounds.cu +++ b/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 diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu index 211692fcc7..11b92b5876 100644 --- a/test/CodeGenCUDA/ptx-kernels.cu +++ b/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" diff --git a/test/PCH/Inputs/cuda.h b/test/PCH/Inputs/cuda.h new file mode 100644 index 0000000000..a9a4595a14 --- /dev/null +++ b/test/PCH/Inputs/cuda.h @@ -0,0 +1,20 @@ +/* Minimal declarations for CUDA support. Testing purposes only. */ + +#include + +#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); diff --git a/test/PCH/cuda-kernel-call.cu b/test/PCH/cuda-kernel-call.cu index ef12c59207..ffb0c1444f 100644 --- a/test/PCH/cuda-kernel-call.cu +++ b/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>>>(); diff --git a/test/SemaCUDA/Inputs/cuda.h b/test/SemaCUDA/Inputs/cuda.h new file mode 100644 index 0000000000..a9a4595a14 --- /dev/null +++ b/test/SemaCUDA/Inputs/cuda.h @@ -0,0 +1,20 @@ +/* Minimal declarations for CUDA support. Testing purposes only. */ + +#include + +#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); diff --git a/test/SemaCUDA/function-target.cu b/test/SemaCUDA/function-target.cu index c7a55e2fad..51bc8c9f55 100644 --- a/test/SemaCUDA/function-target.cu +++ b/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}} diff --git a/test/SemaCUDA/kernel-call.cu b/test/SemaCUDA/kernel-call.cu index 91b1d49e2d..9a3d86c47f 100644 --- a/test/SemaCUDA/kernel-call.cu +++ b/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) {} diff --git a/test/SemaCUDA/launch_bounds.cu b/test/SemaCUDA/launch_bounds.cu index a1a4e64f1c..bed7658aed 100644 --- a/test/SemaCUDA/launch_bounds.cu +++ b/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}} diff --git a/test/SemaCUDA/qualifiers.cu b/test/SemaCUDA/qualifiers.cu index 1346d654b8..42a80b8b38 100644 --- a/test/SemaCUDA/qualifiers.cu +++ b/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}}