[PATCH] D139730: [OpenMP][DeviceRTL][AMDGPU] Support code object version 5

[Saiyedul Islam via Phabricator via cfe-commits]

saiislam added inline comments.


================
Comment at: openmp/libomptarget/DeviceRTL/src/Mapping.cpp:50
 
-uint32_t getNumHardwareThreadsInBlock() {
-  return __builtin_amdgcn_workgroup_size_x();
-}
+uint32_t getNumHardwareThreadsInBlock() { return external_get_local_size(0); }
 
----------------
If we still don't want to depend on rocm-device-libs then we will have to do something like (haven't tried this code yet):

```
uint32_t getNumHardwareThreadsInBlock() {
   if (__oclc_ABI_version < 500) {
      return __builtin_amdgcn_workgroup_size_x();
   } else {
      void *implicitArgPtr = __builtin_amdgcn_implicitarg_ptr();
      return (ushort)implicitArgPtr[6];
}
```


================
Comment at: openmp/libomptarget/DeviceRTL/src/Mapping.cpp:80
 
-uint32_t getNumberOfBlocks() {
-  return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
-}
+uint32_t getNumberOfBlocks() { return external_get_num_groups(0); }
 
----------------
```
uint32_t getNumberOfBlocks() {
   if (__oclc_ABI_version < 500) {
      return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
   } else {
      void *implicitArgPtr = __builtin_amdgcn_implicitarg_ptr();
      return (uint)implicitArgPtr[0];
}
```


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D139730