Merge pull request llvm#490 from AMD-Lightning-Internal/amd/dev/rlieberm/jhuber-preland-123673

preland  123673 [OpenMP] Port the OpenMP device runtime to direct C++ compila…

Author: ronlieb

clang/lib/Headers/__clang_hip_math.h

@@ -21,8 +21,11 @@
 #include <limits.h>
 #include <stdint.h>
 #ifdef __OPENMP_AMDGCN__
+// FIXME: A hack for the OpenMP DeviceRTL's `LibM.h` that should be removed.
+#ifndef __OPENMP_SKIP_INCLUDE__
 #include <omp.h>
 #endif
+#endif
 #endif // !defined(__HIPCC_RTC__)
 
 // __DEVICE__ is a helper macro with common set of attributes for the wrappers

offload/DeviceRTL/CMakeLists.txt

@@ -90,7 +90,7 @@ set(src_files
 # propagation. That said, we will run the vectorizer again after the runtime
 # has been linked into the user program.
 set(clang_opt_flags -O3 -mllvm -openmp-opt-disable -DSHARED_SCRATCHPAD_SIZE=512 -mllvm -vectorize-slp=false )
-set(link_opt_flags  -O3        -openmp-opt-disable -attributor-enable=module -vectorize-slp=false )
+set(link_opt_flags  -O3 -openmp-opt-disable -attributor-enable=module -vectorize-slp=false )
 set(link_export_flag -passes=internalize -internalize-public-api-file=${source_directory}/exports)
 
 # If the user built with the GPU C library enabled we will use that instead.
@@ -103,11 +103,10 @@ set (LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL "${LIBOMPTARGET_LLVM_INCLUDE_DIRS}
 list(TRANSFORM LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL PREPEND "-I")
 
 # Set flags for LLVM Bitcode compilation.
-set(bc_flags -c -foffload-lto -std=c++17 -fvisibility=hidden
-              ${clang_opt_flags} --offload-device-only
-             -nocudalib -nogpulib -nogpuinc
-             -fopenmp -fopenmp-cuda-mode
-             -Wno-unknown-cuda-version -Wno-openmp-target
+set(bc_flags -c -flto -std=c++17 -fvisibility=hidden
+             ${clang_opt_flags} -nocudalib -nogpulib
+             -nogpuinc -nostdlibinc -Wno-unknown-cuda-version
+             -fno-rtti -fno-exceptions -fconvergent-functions
              -I${CMAKE_BINARY_DIR}/openmp/runtime/src # Need omp.h for LibM.
              -I${CMAKE_BINARY_DIR}/projects/openmp/runtime/src # Need omp.h for LibM.
              -I${CMAKE_BINARY_DIR}/runtime/src
@@ -139,9 +138,8 @@ function(compileDeviceRTLLibrary target_name target_triple)
     add_custom_target(${outfile}
       COMMAND ${CLANG_TOOL}
       ${bc_flags}
-      -fopenmp-targets=${target_triple}
-      -Xopenmp-target=${target_triple} -march=
       ${target_bc_flags}
+      --target=${target_triple}
       -MD -MF ${depfile}
       ${infile} -o ${outfile}
       DEPENDS ${infile} ${include_files}
@@ -310,10 +308,7 @@ function(compileDeviceRTLLibrary target_name target_triple)
     set(ide_target_name omptarget-ide-${target_name})
     add_library(${ide_target_name} STATIC EXCLUDE_FROM_ALL ${src_files})
     target_compile_options(${ide_target_name} PRIVATE
-      -fopenmp -fopenmp-cuda-mode
-      -fopenmp-targets=${target_triple} -Xopenmp-target=${target_triple} -march=
-      -mllvm -openmp-opt-disable
-      -foffload-lto -fvisibility=hidden --offload-device-only
+      --target=${target_triple} -flto -fvisibility=hidden
       -nocudalib -nogpulib -nogpuinc -nostdlibinc -Wno-unknown-cuda-version
     )
     target_compile_definitions(${ide_target_name} PRIVATE SHARED_SCRATCHPAD_SIZE=512)

offload/DeviceRTL/include/Allocator.h

@@ -17,8 +17,6 @@
 // Forward declaration.
 struct KernelEnvironmentTy;
 
-#pragma omp begin declare target device_type(nohost)
-
 namespace ompx {
 
 namespace allocator {
@@ -44,6 +42,4 @@ extern "C" {
 [[gnu::weak]] void free(void *Ptr);
 }
 
-#pragma omp end declare target
-
 #endif

offload/DeviceRTL/include/DevRTLExtras.h

@@ -13,13 +13,7 @@
 #ifndef OPENMP_LIBOMPTARGET_DEVICERTL_INCLUDE_DEVRTLEXTRAS_H
 #define OPENMP_LIBOMPTARGET_DEVICERTL_INCLUDE_DEVRTLEXTRAS_H
 
-/// Base type declarations for freestanding mode
-///
-///{
-using uint64_t = unsigned long;
-// TODO: Properly implement this
-using uintptr_t = uint64_t;
-///}
+#include <stdint.h>
 
 /// Macros for allocating variables in different address spaces.
 ///{

offload/DeviceRTL/include/DeviceTypes.h

@@ -100,14 +100,7 @@ struct TaskDescriptorTy {
   TaskFnTy TaskFn;
 };
 
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
 using LaneMaskTy = uint64_t;
-#pragma omp end declare variant
-
-#pragma omp begin declare variant match(                                       \
-        device = {arch(amdgcn)}, implementation = {extension(match_none)})
-using LaneMaskTy = uint64_t;
-#pragma omp end declare variant
 
 namespace lanes {
 enum : LaneMaskTy { All = ~(LaneMaskTy)0 };
@@ -164,8 +157,7 @@ typedef enum omp_allocator_handle_t {
 #define OMP_PRAGMA(STR) __PRAGMA(omp STR)
 
 #define SHARED(NAME)                                                           \
-  NAME [[clang::loader_uninitialized]];                                        \
-  OMP_PRAGMA(allocate(NAME) allocator(omp_pteam_mem_alloc))
+  [[clang::address_space(3)]] NAME [[clang::loader_uninitialized]];
 
 // TODO: clang should use address space 5 for omp_thread_mem_alloc, but right
 //       now that's not the case.

offload/DeviceRTL/include/DeviceUtils.h

@@ -15,8 +15,6 @@
 #include "DeviceTypes.h"
 #include "Shared/Utils.h"
 
-#pragma omp begin declare target device_type(nohost)
-
 namespace utils {
 
 template <typename T> struct type_identity {
@@ -95,6 +93,4 @@ bool isThreadLocalMemPtr(void *Ptr);
 
 } // namespace utils
 
-#pragma omp end declare target
-
 #endif

offload/DeviceRTL/include/Mapping.h

@@ -24,12 +24,8 @@ enum {
   DIM_Z = 2,
 };
 
-#pragma omp begin declare target device_type(nohost)
-
 inline constexpr uint32_t MaxThreadsPerTeam = 1024;
 
-#pragma omp end declare target
-
 /// Initialize the mapping machinery.
 void init(bool IsSPMD);
 

offload/DeviceRTL/include/State.h

@@ -23,8 +23,6 @@
 // Forward declaration.
 struct KernelEnvironmentTy;
 
-#pragma omp begin declare target device_type(nohost)
-
 namespace ompx {
 
 namespace memory {
@@ -89,8 +87,7 @@ struct TeamStateTy {
   ParallelRegionFnTy ParallelRegionFnVar;
 };
 
-extern TeamStateTy TeamState;
-#pragma omp allocate(TeamState) allocator(omp_pteam_mem_alloc)
+extern TeamStateTy [[clang::address_space(3)]] TeamState;
 
 struct ThreadStateTy {
 
@@ -116,8 +113,7 @@ struct ThreadStateTy {
   }
 };
 
-extern ThreadStateTy **ThreadStates;
-#pragma omp allocate(ThreadStates) allocator(omp_pteam_mem_alloc)
+extern ThreadStateTy **[[clang::address_space(3)]] ThreadStates;
 
 /// Initialize the state machinery. Must be called by all threads.
 void init(bool IsSPMD, KernelEnvironmentTy &KernelEnvironment,
@@ -379,6 +375,4 @@ inline state::Value<uint32_t, state::VK_RunSched> RunSched;
 
 } // namespace ompx
 
-#pragma omp end declare target
-
 #endif

offload/DeviceRTL/include/Synchronization.h

@@ -15,8 +15,6 @@
 #include "DeviceTypes.h"
 #include "DeviceUtils.h"
 
-#pragma omp begin declare target device_type(nohost)
-
 namespace ompx {
 namespace atomic {
 
@@ -221,6 +219,4 @@ void system(atomic::OrderingTy Ordering);
 
 } // namespace ompx
 
-#pragma omp end declare target
-
 #endif

offload/DeviceRTL/include/Workshare.h

@@ -12,8 +12,6 @@
 #ifndef OMPTARGET_WORKSHARE_H
 #define OMPTARGET_WORKSHARE_H
 
-#pragma omp begin declare target device_type(nohost)
-
 namespace ompx {
 
 namespace workshare {
@@ -25,6 +23,4 @@ void init(bool IsSPMD);
 
 } // namespace ompx
 
-#pragma omp end declare target
-
 #endif

offload/DeviceRTL/include/extra_allocators.h

@@ -61,7 +61,7 @@ void omp_free(void *ptr, omp_allocator_handle_t allocator = omp_null_allocator);
 //// Any re-definitions of malloc/free delete the native CUDA
 //// but they are necessary
 #ifdef __AMDGCN__
-void *malloc(uint64_t Size);
+void *malloc(size_t Size);
 void free(void *Ptr);
 #endif
 } // extern "C"

offload/DeviceRTL/src/Allocator.cpp

@@ -19,8 +19,6 @@
 
 using namespace ompx;
 
-#pragma omp begin declare target device_type(nohost)
-
 [[gnu::used, gnu::retain, gnu::weak,
   gnu::visibility(
       "protected")]] DeviceMemoryPoolTy __omp_rtl_device_memory_pool;
@@ -77,5 +75,3 @@ void *allocator::alloc(uint64_t Size) { return BumpAllocator.alloc(Size); }
 void allocator::free(void *Ptr) { BumpAllocator.free(Ptr); }
 
 ///}
-
-#pragma omp end declare target

offload/DeviceRTL/src/Configuration.cpp

@@ -17,8 +17,6 @@
 
 using namespace ompx;
 
-#pragma omp begin declare target device_type(nohost)
-
 // Weak definitions will be overridden by CGOpenmpRuntimeGPU if enabled.
 [[gnu::weak]] extern const uint32_t __omp_rtl_debug_kind = 0;
 [[gnu::weak]] extern const uint32_t __omp_rtl_assume_no_thread_state = 0;
@@ -85,5 +83,3 @@ bool config::mayUseNestedParallelism() {
     return false;
   return state::getKernelEnvironment().Configuration.MayUseNestedParallelism;
 }
-
-#pragma omp end declare target

offload/DeviceRTL/src/Debug.cpp

@@ -21,8 +21,6 @@
 
 using namespace ompx;
 
-#pragma omp begin declare target device_type(nohost)
-
 extern "C" {
 void __assert_assume(bool condition) { __builtin_assume(condition); }
 
@@ -44,5 +42,3 @@ void __assert_fail_internal(const char *expr, const char *msg, const char *file,
   __builtin_trap();
 }
 }
-
-#pragma omp end declare target

offload/DeviceRTL/src/DeviceUtils.cpp

@@ -15,8 +15,6 @@
 #include "Interface.h"
 #include "Mapping.h"
 
-#pragma omp begin declare target device_type(nohost)
-
 using namespace ompx;
 
 extern "C" [[gnu::weak]] int IsSPMDMode;
@@ -33,8 +31,6 @@ __keep_alive() {
 
 namespace impl {
 
-bool isSharedMemPtr(const void *Ptr) { return false; }
-
 void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits) {
   static_assert(sizeof(unsigned long) == 8, "");
   *LowBits = static_cast<uint32_t>(Val & 0x00000000FFFFFFFFUL);
@@ -54,7 +50,7 @@ uint64_t ballotSync(uint64_t Mask, int32_t Pred);
 /// AMDGCN Implementation
 ///
 ///{
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
+#ifdef __AMDGPU__
 
 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane, int32_t Width) {
   int Self = mapping::getThreadIdInWarp();
@@ -78,15 +74,13 @@ bool isSharedMemPtr(const void *Ptr) {
   return __builtin_amdgcn_is_shared(
       (const __attribute__((address_space(0))) void *)Ptr);
 }
-#pragma omp end declare variant
+#endif
 ///}
 
 /// NVPTX Implementation
 ///
 ///{
-#pragma omp begin declare variant match(                                       \
-        device = {arch(nvptx, nvptx64)},                                       \
-            implementation = {extension(match_any)})
+#ifdef __NVPTX__
 
 int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane, int32_t Width) {
   return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, Width - 1);
@@ -103,7 +97,7 @@ uint64_t ballotSync(uint64_t Mask, int32_t Pred) {
 
 bool isSharedMemPtr(const void *Ptr) { return __nvvm_isspacep_shared(Ptr); }
 
-#pragma omp end declare variant
+#endif
 ///}
 } // namespace impl
 
@@ -149,5 +143,3 @@ int64_t __kmpc_shuffle_int64(int64_t Val, int16_t Delta, int16_t Width) {
   return utils::shuffleDown(lanes::All, Val, Delta, Width);
 }
 }
-
-#pragma omp end declare target

offload/DeviceRTL/src/ExtraMapping.cpp

@@ -1,8 +1,6 @@
 #include "Interface.h"
 #include "Mapping.h"
 
-#pragma omp declare target
-
 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
 using namespace ompx::mapping;
@@ -13,10 +11,8 @@ namespace impl {
 /// AMDGCN Implementation
 ///
 ///{
-uint32_t __kmpc_impl_smid();
-uint32_t getGenericModeMainThreadId();
 
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
+#ifdef __AMDGPU__
 
 // Partially derived fom hcc_detail/device_functions.h
 
@@ -75,15 +71,14 @@ static uint32_t getGenericModeMainThreadId() {
   return (__kmpc_get_hardware_num_threads_in_block() - 1) & (~Mask);
 }
 
-#pragma omp end declare variant
+#endif
 ///}
 
 /// NVPTX Implementation
 ///
 ///{
 
-#pragma omp begin declare variant match(                                       \
-    device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
+#ifdef __NVPTX__
 
 static uint32_t __kmpc_impl_smid() { return 0; }
 
@@ -92,7 +87,7 @@ static uint32_t getGenericModeMainThreadId() {
   return (__kmpc_get_hardware_num_threads_in_block() - 1) & (~Mask);
 }
 
-#pragma omp end declare variant
+#endif
 ///}
 
 } // namespace impl
@@ -113,7 +108,7 @@ int omp_ext_get_lane_id() {
 }
 
 int omp_ext_get_smid() {
-  int rc = impl::__kmpc_impl_smid();
+  int rc = ompx::mapping::impl::__kmpc_impl_smid();
   return rc;
 }
 
@@ -137,5 +132,3 @@ unsigned long long omp_ext_get_active_threads_mask() {
 }
 
 } // end extern "C"
-
-#pragma omp end declare target