[Openmp-commits] [PATCH] D95013: [OpenMP][NVPTX] Replaced CUDA builtin vars with LLVM intrinsics

[Shilei Tian via Phabricator via Openmp-commits]

tianshilei1992 created this revision.
tianshilei1992 added reviewers: jdoerfert, JonChesterfield.
Herald added subscribers: guansong, yaxunl.
tianshilei1992 requested review of this revision.
Herald added subscribers: openmp-commits, sstefan1.
Herald added a project: OpenMP.

Replaced CUDA builtin vars with LLVM intrinsics such that we don't need
definitions of those intrinsics.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D95013

Files:
  openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu


Index: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu
+++ openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu
@@ -114,10 +114,10 @@
 DEVICE void __kmpc_impl_threadfence_system() { __threadfence_system(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
-DEVICE int GetThreadIdInBlock() { return threadIdx.x; }
-DEVICE int GetBlockIdInKernel() { return blockIdx.x; }
-DEVICE int GetNumberOfBlocksInKernel() { return gridDim.x; }
-DEVICE int GetNumberOfThreadsInBlock() { return blockDim.x; }
+DEVICE int GetThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
+DEVICE int GetBlockIdInKernel() { return __nvvm_read_ptx_sreg_ctaid_x(); }
+DEVICE int GetNumberOfBlocksInKernel() { return __nvvm_read_ptx_sreg_nctaid_x(); }
+DEVICE int GetNumberOfThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); }
 DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D95013.317741.patch
Type: text/x-patch
Size: 1107 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20210120/e388d9f4/attachment.bin>