[PATCH] D104904: [OpenMP][AMDGCN] Initial math headers support

[Pushpinder Singh via Phabricator via cfe-commits]

pdhaliwal added inline comments.


================
Comment at: clang/lib/Headers/__clang_hip_math.h:29
+#else
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#endif
----------------
JonChesterfield wrote:
> wonder if HIP would benefit from nothrow here
Would like to keep hip changes minimal in this patch.


================
Comment at: clang/lib/Headers/__clang_hip_math.h:35
+#ifdef __OPENMP_AMDGCN__
+#define __RETURN_TYPE int
+#else
----------------
jdoerfert wrote:
> JonChesterfield wrote:
> > I'd expect openmp to match the cplusplus/c distinction here, as openmp works on C source
> ^ Agreed. Though, we use a different trick because it's unfortunately not as straight forward always and can be decided based on the C vs C++.
This is somewhat tricky. Since declaration of `__finite/__isnan /__isinff` is with int return type in standard library (and the corresponding methods in C++ seems to be isfinite, isnan and isinf with bool return type), the compiler fails to resolve these functions when using bool. I don't know how HIP is working.

__RETURN_TYPE macro is only being used with the following methods:
1. __finite
2. __isnan
3. __isinf
4. __signbit

and with the corresponding float versions.


================
Comment at: clang/lib/Headers/openmp_wrappers/cmath:83
+#include <__clang_hip_cmath.h>
+#undef __OPENMP_AMDGCN__
+
----------------
jdoerfert wrote:
> No match_any needed (here and elsewhere).
> 
> Also, don't we want all but the includes to be the same for both GPUs. Maybe we have a device(kind(gpu)) variant and inside the nvptx and amdgpu just for the respective include?
device(kind(gpu)) breaks nvptx and hip with lots of errors like below,

```
...
__clang_cuda_device_functions.h:29:40: error: use of undeclared identifier '__nvvm_vote_all'
...
```

Maybe I am doing something wrong.


================
Comment at: clang/test/Headers/Inputs/include/cstdlib:15
 
+#ifndef __AMDGCN__
 namespace std
----------------
jdoerfert wrote:
> JonChesterfield wrote:
> > I think I'd expect builtin_labs et al to work on amdgcn, are we missing lowering for them?
> Yeah, looks weird that we cannot compile this mock-up header.
>From what I understand, hip is defining fabs to use ocml's version into the std namespace, which was already defined in this header. So that's causing multiple declaration error. I will wrap only fabs in the ifdef's


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D104904/new/

https://reviews.llvm.org/D104904