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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
