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