pdhaliwal updated this revision to Diff 354471.
pdhaliwal added a comment.
Fix format errors
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D104904/new/
https://reviews.llvm.org/D104904
Files:
clang/lib/Driver/ToolChains/Clang.cpp
clang/lib/Headers/__clang_hip_cmath.h
clang/lib/Headers/__clang_hip_math.h
clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
clang/lib/Headers/openmp_wrappers/cmath
clang/lib/Headers/openmp_wrappers/math.h
clang/test/Headers/Inputs/include/algorithm
clang/test/Headers/Inputs/include/cstdlib
clang/test/Headers/Inputs/include/utility
clang/test/Headers/amdgcn_openmp_device_math.c
clang/test/Headers/openmp_device_math_isnan.cpp
Index: clang/test/Headers/openmp_device_math_isnan.cpp
===================================================================
--- clang/test/Headers/openmp_device_math_isnan.cpp
+++ clang/test/Headers/openmp_device_math_isnan.cpp
@@ -21,14 +21,14 @@
double math(float f, double d) {
double r = 0;
// INT_RETURN: call i32 @__nv_isnanf(float
- // AMD_INT_RETURN: call i32 @_{{.*}}isnanf(float
+ // AMD_INT_RETURN: call i32 @__ocml_isnan_f32(float
// BOOL_RETURN: call i32 @__nv_isnanf(float
- // AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnanf(float
+ // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f32(float
r += std::isnan(f);
// INT_RETURN: call i32 @__nv_isnand(double
- // AMD_INT_RETURN: call i32 @_{{.*}}isnand(double
+ // AMD_INT_RETURN: call i32 @__ocml_isnan_f64(double
// BOOL_RETURN: call i32 @__nv_isnand(double
- // AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnand(double
+ // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f64(double
r += std::isnan(d);
return r;
}
Index: clang/test/Headers/amdgcn_openmp_device_math.c
===================================================================
--- /dev/null
+++ clang/test/Headers/amdgcn_openmp_device_math.c
@@ -0,0 +1,51 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK-C,CHECK
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK-CPP,CHECK
+
+#ifdef __cplusplus
+#include <cmath>
+#else
+#include <math.h>
+#endif
+
+void test_math_f64(double x) {
+// CHECK-LABEL: define {{.*}}test_math_f64
+#pragma omp target
+ {
+ // CHECK: call double @__ocml_sin_f64
+ double l1 = sin(x);
+ // CHECK: call double @__ocml_cos_f64
+ double l2 = cos(x);
+ // CHECK: call double @__ocml_fabs_f64
+ double l3 = fabs(x);
+ }
+}
+
+void test_math_f32(float x) {
+// CHECK-LABEL: define {{.*}}test_math_f32
+#pragma omp target
+ {
+ // CHECK-C: call double @__ocml_sin_f64
+ // CHECK-CPP: call float @__ocml_sin_f32
+ float l1 = sin(x);
+ // CHECK-C: call double @__ocml_cos_f64
+ // CHECK-CPP: call float @__ocml_cos_f32
+ float l2 = cos(x);
+ // CHECK-C: call double @__ocml_fabs_f64
+ // CHECK-CPP: call float @__ocml_fabs_f32
+ float l3 = fabs(x);
+ }
+}
+void test_math_f32_suffix(float x) {
+// CHECK-LABEL: define {{.*}}test_math_f32_suffix
+#pragma omp target
+ {
+ // CHECK: call float @__ocml_sin_f32
+ float l1 = sinf(x);
+ // CHECK: call float @__ocml_cos_f32
+ float l2 = cosf(x);
+ // CHECK: call float @__ocml_fabs_f32
+ float l3 = fabsf(x);
+ }
+}
Index: clang/test/Headers/Inputs/include/utility
===================================================================
--- /dev/null
+++ clang/test/Headers/Inputs/include/utility
@@ -0,0 +1,2 @@
+#pragma once
+
Index: clang/test/Headers/Inputs/include/cstdlib
===================================================================
--- clang/test/Headers/Inputs/include/cstdlib
+++ clang/test/Headers/Inputs/include/cstdlib
@@ -12,6 +12,7 @@
extern float fabs (float __x) __attribute__ ((__const__)) ;
#endif
+#ifndef __AMDGCN__
namespace std
{
@@ -29,3 +30,5 @@
double abs(double __x) { return fabs(__x); }
}
+
+#endif
Index: clang/test/Headers/Inputs/include/algorithm
===================================================================
--- /dev/null
+++ clang/test/Headers/Inputs/include/algorithm
@@ -0,0 +1,6 @@
+#pragma once
+
+namespace std {
+ template<class T> constexpr const T& min(const T& a, const T& b);
+ template<class T> constexpr const T& max(const T& a, const T& b);
+}
\ No newline at end of file
Index: clang/lib/Headers/openmp_wrappers/math.h
===================================================================
--- clang/lib/Headers/openmp_wrappers/math.h
+++ clang/lib/Headers/openmp_wrappers/math.h
@@ -48,4 +48,13 @@
#pragma omp end declare variant
+#pragma omp begin declare variant match( \
+ device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#define __OPENMP_AMDGCN__
+#include <__clang_hip_math.h>
+#undef __OPENMP_AMDGCN__
+
+#pragma omp end declare variant
+
#endif
Index: clang/lib/Headers/openmp_wrappers/cmath
===================================================================
--- clang/lib/Headers/openmp_wrappers/cmath
+++ clang/lib/Headers/openmp_wrappers/cmath
@@ -75,4 +75,14 @@
#pragma omp end declare variant
+#ifdef __AMDGCN__
+#pragma omp begin declare variant match( \
+ device = {arch(amdgcn)}, implementation = {extension(match_any, allow_templates)})
+#define __OPENMP_AMDGCN__
+#include <__clang_hip_cmath.h>
+#undef __OPENMP_AMDGCN__
+
+
+#pragma omp end declare variant
+#endif // __AMDGCN__
#endif
Index: clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
===================================================================
--- clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
+++ clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
@@ -81,4 +81,34 @@
#pragma pop_macro("OPENMP_NOEXCEPT")
#endif
+#pragma omp begin declare variant match( \
+ device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define __OPENMP_AMDGCN__
+
+#define __device__ __attribute__((device))
+
+// Import types which will be used by __clang_hip_libdevice_declares.h
+#ifndef __cplusplus
+#include <stdbool.h>
+#include <stdint.h>
+#endif
+
+/// Include declarations for libdevice functions.
+#include <__clang_hip_libdevice_declares.h>
+
+#undef __device__
+#undef __OPENMP_AMDGCN__
+
+#ifdef __cplusplus
+} // extern "C"
+
+#endif
+
+#pragma omp end declare variant
+
#endif
Index: clang/lib/Headers/__clang_hip_math.h
===================================================================
--- clang/lib/Headers/__clang_hip_math.h
+++ clang/lib/Headers/__clang_hip_math.h
@@ -9,7 +9,7 @@
#ifndef __CLANG_HIP_MATH_H__
#define __CLANG_HIP_MATH_H__
-#if !defined(__HIP__)
+#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
#error "This file is for HIP and OpenMP AMDGCN device compilation only."
#endif
@@ -19,18 +19,27 @@
#endif
#include <limits.h>
#include <stdint.h>
-#endif // __HIPCC_RTC__
+#endif // !defined(__HIPCC_RTC__)
#pragma push_macro("__DEVICE__")
+
+#ifdef __OPENMP_AMDGCN__
+#define __DEVICE__ static inline __attribute__((always_inline, nothrow))
+#else
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#endif
// A few functions return bool type starting only in C++11.
#pragma push_macro("__RETURN_TYPE")
+#ifdef __OPENMP_AMDGCN__
+#define __RETURN_TYPE int
+#else
#if defined(__cplusplus)
#define __RETURN_TYPE bool
#else
#define __RETURN_TYPE int
#endif
+#endif // __OPENMP_AMDGCN__
#if defined (__cplusplus) && __cplusplus < 201103L
// emulate static_assert on type sizes
@@ -1262,7 +1271,7 @@
__DEVICE__
double min(double __x, double __y) { return fmin(__x, __y); }
-#if !defined(__HIPCC_RTC__)
+#if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
__host__ inline static int min(int __arg1, int __arg2) {
return std::min(__arg1, __arg2);
}
@@ -1270,7 +1279,7 @@
__host__ inline static int max(int __arg1, int __arg2) {
return std::max(__arg1, __arg2);
}
-#endif // __HIPCC_RTC__
+#endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
#endif
#pragma pop_macro("__DEVICE__")
Index: clang/lib/Headers/__clang_hip_cmath.h
===================================================================
--- clang/lib/Headers/__clang_hip_cmath.h
+++ clang/lib/Headers/__clang_hip_cmath.h
@@ -10,7 +10,7 @@
#ifndef __CLANG_HIP_CMATH_H__
#define __CLANG_HIP_CMATH_H__
-#if !defined(__HIP__)
+#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
#error "This file is for HIP and OpenMP AMDGCN device compilation only."
#endif
@@ -25,7 +25,12 @@
#endif // !defined(__HIPCC_RTC__)
#pragma push_macro("__DEVICE__")
+#ifdef __OPENMP_AMDGCN__
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+#define __constant__ __attribute__((constant))
+#else
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#endif // __OPENMP_AMDGCN__
// Start with functions that cannot be defined by DEF macros below.
#if defined(__cplusplus)
Index: clang/lib/Driver/ToolChains/Clang.cpp
===================================================================
--- clang/lib/Driver/ToolChains/Clang.cpp
+++ clang/lib/Driver/ToolChains/Clang.cpp
@@ -1249,7 +1249,8 @@
// If we are offloading to a target via OpenMP we need to include the
// openmp_wrappers folder which contains alternative system headers.
if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
- getToolChain().getTriple().isNVPTX()){
+ (getToolChain().getTriple().isNVPTX() ||
+ getToolChain().getTriple().isAMDGCN())) {
if (!Args.hasArg(options::OPT_nobuiltininc)) {
// Add openmp_wrappers/* to our system include path. This lets us wrap
// standard library headers.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
