[Openmp-commits] [PATCH] D96533: [libomptarget][amdgcn] Build amdgcn devicertl as openmp

Jon Chesterfield via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Thu Feb 11 14:04:02 PST 2021


JonChesterfield added inline comments.


================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_intrinsics.h:46
+
+#endif
----------------
jdoerfert wrote:
> JonChesterfield wrote:
> > jdoerfert wrote:
> > > Do we need to declare builtins? I guess we can but we should not have to, right?
> > Today, yes. I think the 'right' behaviour is to allow these within an appropriate declare variant, and an adequate approximation would be to allow them everywhere.
> > 
> > As raw openmp, the assumption is that any function can be called from the host. That doesn't mean anything for __builtin_amdgcn_dispatch_ptr, so it's reasonable that these aren't automatically available.
> > 
> > At present though, I'm just trying to get this library to compile with trunk clang as-is. Curiously, if one of these is declared with the wrong type, clang emits an error on the declaration.
> I can just call these like other builtins (when I target amdgcn):
> https://godbolt.org/z/nabGT7
That is interesting.

I wrote this patch against a version of --cuda-device-only. These may have only raised an error on the host pass, which is not invoked by the present -Xclang themed cmake.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D96533



More information about the Openmp-commits mailing list