[Openmp-commits] [PATCH] D71446: [libomptarget] Build most of common/src for amdgcn

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: jdoerfert, ABataev, grokos.
Herald added subscribers: openmp-commits, mgorny, jvesely.
Herald added a project: OpenMP.

[libomptarget] Build most of common/src for amdgcn

Excluding parallel.cu, which uses an integer min() from cuda,
Excluding support.cu, which calls malloc that is not yet available for amdgcn


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D71446

Files:
  openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
  openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h


Index: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
===================================================================
--- openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
+++ openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
@@ -121,6 +121,10 @@
 
 INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
 
+INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t) {
+  // AMDGCN doesn't need to sync threads in a warp
+}
+
 INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
   // we have protected the master warp from releasing from its barrier
   // due to a full workgroup barrier in the middle of a work function.
Index: openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
===================================================================
--- openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
+++ openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
@@ -56,7 +56,11 @@
 
 set(cuda_sources
   ${devicertl_base_directory}/common/src/cancel.cu
-  ${devicertl_base_directory}/common/src/critical.cu)
+  ${devicertl_base_directory}/common/src/critical.cu
+  ${devicertl_base_directory}/common/src/loop.cu
+  ${devicertl_base_directory}/common/src/omptarget.cu
+  ${devicertl_base_directory}/common/src/sync.cu
+  ${devicertl_base_directory}/common/src/task.cu)
 
 set(h_files
   ${CMAKE_CURRENT_SOURCE_DIR}/src/amdgcn_interface.h
@@ -64,6 +68,8 @@
   ${CMAKE_CURRENT_SOURCE_DIR}/src/target_impl.h
   ${devicertl_base_directory}/common/debug.h
   ${devicertl_base_directory}/common/device_environment.h
+  ${devicertl_base_directory}/common/omptarget.h
+  ${devicertl_base_directory}/common/omptargeti.h
   ${devicertl_base_directory}/common/state-queue.h
   ${devicertl_base_directory}/common/state-queuei.h
   ${devicertl_base_directory}/common/support.h)


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D71446.233720.patch
Type: text/x-patch
Size: 1827 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20191213/c8da72f4/attachment.bin>