[clang] 75e941b - [NFC][OpenMP][CUDA] Add test for using `-x cuda -fopenmp`
Joachim Meyer via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 2 09:59:59 PDT 2021
Author: Joachim Meyer
Date: 2021-07-02T19:03:15+02:00
New Revision: 75e941b05c78e19c51e5c960c925f334b08109bb
URL: https://github.com/llvm/llvm-project/commit/75e941b05c78e19c51e5c960c925f334b08109bb
DIFF: https://github.com/llvm/llvm-project/commit/75e941b05c78e19c51e5c960c925f334b08109bb.diff
LOG: [NFC][OpenMP][CUDA] Add test for using `-x cuda -fopenmp`
This adds a very basic test in `cuda_with_openmp.cu` that just checks whether the CUDA & OpenMP integrated headers do compile, when a CUDA file is compiled with OpenMP (CPU) enabled.
Thus this basically adds the missing test for https://reviews.llvm.org/D90415.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D105322
Added:
clang/test/Headers/Inputs/include/crt/device_double_functions.hpp
clang/test/Headers/Inputs/include/crt/device_functions.hpp
clang/test/Headers/Inputs/include/crt/device_runtime.h
clang/test/Headers/Inputs/include/crt/host_runtime.h
clang/test/Headers/Inputs/include/crt/math_functions.hpp
clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp
clang/test/Headers/Inputs/include/cuda.h
clang/test/Headers/Inputs/include/cuda_runtime.h
clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h
clang/test/Headers/Inputs/include/device_atomic_functions.h
clang/test/Headers/Inputs/include/device_atomic_functions.hpp
clang/test/Headers/Inputs/include/device_double_functions.h
clang/test/Headers/Inputs/include/driver_types.h
clang/test/Headers/Inputs/include/host_config.h
clang/test/Headers/Inputs/include/host_defines.h
clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp
clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp
clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp
clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp
clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp
clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp
clang/test/Headers/Inputs/include/string.h
clang/test/Headers/Inputs/include/texture_indirect_functions.h
clang/test/Headers/cuda_with_openmp.cu
Modified:
clang/test/Headers/Inputs/include/cstdlib
clang/test/Headers/Inputs/include/new
Removed:
################################################################################
diff --git a/clang/test/Headers/Inputs/include/crt/device_double_functions.hpp b/clang/test/Headers/Inputs/include/crt/device_double_functions.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/device_double_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/crt/device_functions.hpp b/clang/test/Headers/Inputs/include/crt/device_functions.hpp
new file mode 100644
index 0000000000000..41dc8722fe643
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/device_functions.hpp
@@ -0,0 +1,3 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+__device__ void __brkpt();
diff --git a/clang/test/Headers/Inputs/include/crt/device_runtime.h b/clang/test/Headers/Inputs/include/crt/device_runtime.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/device_runtime.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/crt/host_runtime.h b/clang/test/Headers/Inputs/include/crt/host_runtime.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/host_runtime.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/crt/math_functions.hpp b/clang/test/Headers/Inputs/include/crt/math_functions.hpp
new file mode 100644
index 0000000000000..fbd2bb52fa842
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/math_functions.hpp
@@ -0,0 +1,12 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+__device__ int __isinff(float);
+__device__ int __isinf(double);
+__device__ int __finitef(float);
+__device__ int __isfinited(double);
+__device__ int __isnanf(float);
+__device__ int __isnan(double);
+__device__ int __signbitf(float);
+__device__ int __signbitd(double);
+__device__ double max(double, double);
+__device__ float max(float, float);
diff --git a/clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp b/clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/cstdlib b/clang/test/Headers/Inputs/include/cstdlib
index 1d1864a98976b..689b5e06edec9 100644
--- a/clang/test/Headers/Inputs/include/cstdlib
+++ b/clang/test/Headers/Inputs/include/cstdlib
@@ -5,11 +5,9 @@
#if __cplusplus >= 201703L
extern int abs (int __x) throw() __attribute__ ((__const__)) ;
extern long int labs (long int __x) throw() __attribute__ ((__const__)) ;
-extern float fabs (float __x) throw() __attribute__ ((__const__)) ;
#else
extern int abs (int __x) __attribute__ ((__const__)) ;
extern long int labs (long int __x) __attribute__ ((__const__)) ;
-extern float fabs (float __x) __attribute__ ((__const__)) ;
#endif
namespace std
diff --git a/clang/test/Headers/Inputs/include/cuda.h b/clang/test/Headers/Inputs/include/cuda.h
new file mode 100644
index 0000000000000..32870938a8e18
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/cuda.h
@@ -0,0 +1,127 @@
+/* Minimal declarations for CUDA support. Testing purposes only. */
+#pragma once
+
+#include <stddef.h>
+
+// Make this file work with nvcc, for testing compatibility.
+
+#ifndef __NVCC__
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __managed__ __attribute__((managed))
+#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) {}
+};
+
+// Host- and device-side placement new overloads.
+void *operator new(__SIZE_TYPE__, void *p) { return p; }
+void *operator new[](__SIZE_TYPE__, void *p) { return p; }
+__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; }
+__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; }
+
+#define CUDA_VERSION 10100
+
+struct char2 {
+ char x, y;
+ __host__ __device__ char2(char x = 0, char y = 0) : x(x), y(y) {}
+};
+struct char4 {
+ char x, y, z, w;
+ __host__ __device__ char4(char x = 0, char y = 0, char z = 0, char w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct uchar2 {
+ unsigned char x, y;
+ __host__ __device__ uchar2(unsigned char x = 0, unsigned char y = 0) : x(x), y(y) {}
+};
+struct uchar4 {
+ unsigned char x, y, z, w;
+ __host__ __device__ uchar4(unsigned char x = 0, unsigned char y = 0, unsigned char z = 0, unsigned char w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct short2 {
+ short x, y;
+ __host__ __device__ short2(short x = 0, short y = 0) : x(x), y(y) {}
+};
+struct short4 {
+ short x, y, z, w;
+ __host__ __device__ short4(short x = 0, short y = 0, short z = 0, short w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct ushort2 {
+ unsigned short x, y;
+ __host__ __device__ ushort2(unsigned short x = 0, unsigned short y = 0) : x(x), y(y) {}
+};
+struct ushort4 {
+ unsigned short x, y, z, w;
+ __host__ __device__ ushort4(unsigned short x = 0, unsigned short y = 0, unsigned short z = 0, unsigned short w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct int2 {
+ int x, y;
+ __host__ __device__ int2(int x = 0, int y = 0) : x(x), y(y) {}
+};
+struct int4 {
+ int x, y, z, w;
+ __host__ __device__ int4(int x = 0, int y = 0, int z = 0, int w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct uint2 {
+ unsigned x, y;
+ __host__ __device__ uint2(unsigned x = 0, unsigned y = 0) : x(x), y(y) {}
+};
+struct uint3 {
+ unsigned x, y, z;
+ __host__ __device__ uint3(unsigned x = 0, unsigned y = 0, unsigned z = 0) : x(x), y(y), z(z) {}
+};
+struct uint4 {
+ unsigned x, y, z, w;
+ __host__ __device__ uint4(unsigned x = 0, unsigned y = 0, unsigned z = 0, unsigned w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+
+struct longlong2 {
+ long long x, y;
+ __host__ __device__ longlong2(long long x = 0, long long y = 0) : x(x), y(y) {}
+};
+struct longlong4 {
+ long long x, y, z, w;
+ __host__ __device__ longlong4(long long x = 0, long long y = 0, long long z = 0, long long w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct ulonglong2 {
+ unsigned long long x, y;
+ __host__ __device__ ulonglong2(unsigned long long x = 0, unsigned long long y = 0) : x(x), y(y) {}
+};
+struct ulonglong4 {
+ unsigned long long x, y, z, w;
+ __host__ __device__ ulonglong4(unsigned long long x = 0, unsigned long long y = 0, unsigned long long z = 0, unsigned long long w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+
+struct float2 {
+ float x, y;
+ __host__ __device__ float2(float x = 0, float y = 0) : x(x), y(y) {}
+};
+struct float4 {
+ float x, y, z, w;
+ __host__ __device__ float4(float x = 0, float y = 0, float z = 0, float w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct double2 {
+ double x, y;
+ __host__ __device__ double2(double x = 0, double y = 0) : x(x), y(y) {}
+};
+struct double4 {
+ double x, y, z, w;
+ __host__ __device__ double4(double x = 0, double y = 0, double z = 0, double w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+
+#endif // !__NVCC__
diff --git a/clang/test/Headers/Inputs/include/cuda_runtime.h b/clang/test/Headers/Inputs/include/cuda_runtime.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/cuda_runtime.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h b/clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/device_atomic_functions.h b/clang/test/Headers/Inputs/include/device_atomic_functions.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/device_atomic_functions.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/device_atomic_functions.hpp b/clang/test/Headers/Inputs/include/device_atomic_functions.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/device_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/device_double_functions.h b/clang/test/Headers/Inputs/include/device_double_functions.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/device_double_functions.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/driver_types.h b/clang/test/Headers/Inputs/include/driver_types.h
new file mode 100644
index 0000000000000..b5d366aee29fd
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/driver_types.h
@@ -0,0 +1,4 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+
+#include <limits.h>
diff --git a/clang/test/Headers/Inputs/include/host_config.h b/clang/test/Headers/Inputs/include/host_config.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/host_config.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/host_defines.h b/clang/test/Headers/Inputs/include/host_defines.h
new file mode 100644
index 0000000000000..03fed3a2a5d46
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/host_defines.h
@@ -0,0 +1,3 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+#define __forceinline__
diff --git a/clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp b/clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/new b/clang/test/Headers/Inputs/include/new
index 8159d5527cc3a..5110731fb2f89 100644
--- a/clang/test/Headers/Inputs/include/new
+++ b/clang/test/Headers/Inputs/include/new
@@ -1,3 +1,4 @@
+#pragma once
namespace std
{
diff --git a/clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp b/clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp b/clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp b/clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp b/clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp b/clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/string.h b/clang/test/Headers/Inputs/include/string.h
new file mode 100644
index 0000000000000..98cf77fd564c0
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/string.h
@@ -0,0 +1,3 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+void* memcpy(void* dst, const void* src, size_t num);
diff --git a/clang/test/Headers/Inputs/include/texture_indirect_functions.h b/clang/test/Headers/Inputs/include/texture_indirect_functions.h
new file mode 100644
index 0000000000000..bffa775cb2822
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/texture_indirect_functions.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/cuda_with_openmp.cu b/clang/test/Headers/cuda_with_openmp.cu
new file mode 100644
index 0000000000000..efde4ecdc6626
--- /dev/null
+++ b/clang/test/Headers/cuda_with_openmp.cu
@@ -0,0 +1,8 @@
+// Test using -x cuda -fopenmp does not clash integrated headers.
+// Reported in https://bugs.llvm.org/show_bug.cgi?id=48014
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang -x cuda -fopenmp -c %s -o - --cuda-path=%S/../Driver/Inputs/CUDA/usr/local/cuda -nocudalib -isystem %S/Inputs/include -isystem %S/../../lib/Headers -fsyntax-only
+
More information about the cfe-commits
mailing list