[libomptarget] Build a minimal deviceRTL for amdgcn

[JonChesterfield]

Summary

[libomptarget] Build a minimal deviceRTL for amdgcn

The CMakeLists.txt file is functionally identical to the one used in the aomp fork.
Whitespace changes were made based on nvptx/CMakeLists.txt, plus the
copyright notice updated to match (Greg was the original author so would
like his sign off on that here).

This change will build a small subset of the deviceRTL if an appropriate toolchain is
available, e.g. a local install of rocm. Support.h is moved from nvptx as a dependency
of debug.h.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

57

Glob has the advantage that this file will keep working as other source is added to amdgcn or common.

This will break as soon as common contains a file that amdgcn doesn't want to compile. At that point, this can become a whitelist instead.

I'm happy to list out the sources explicitly if that's strongly preferred.

74

Builds a specialised version of the library for various amdgcn architectures

[Hahnfeld]

Random drive-by comments to parts of the patch, I won't review in full.

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

57

As far as I know globs have the disadvantage that they need an explicit invocation of CMake to get an updated list of source files. That's why LLVM usually avoids them.

71–72

I don't think it's a good idea to honor random environment variables. I'd suggest to make this a proper configuration variable.

133

OPENMP_INSTALL_LIBDIR

[JonChesterfield]

• address review comments

[JonChesterfield]

This is actually quite close to compiling under Release, at least for cancel.cu and critical.cu which don't include omptarget-nvptx.h. Debug.h unconditionally includes support.h (looks unintentional) and amdgcn's target_impl currently doesn't work without the cuda shim.

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

57

Fair enough. Explicit lists it is.

71–72

Agreed. Added a variable loosely based on the equivalent one for nvptx and checked I can still plumb it into the out of tree build scripts.

133

Nice, thanks.

[JonChesterfield]

• Drop device_environment from this patch

[JonChesterfield]

• Import support.h, update target_impl such that the deviceRTL builds successfully

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

17

This file is correct for the not yet upstream hsa plugin. For the purposes of this change the num_devices and device_num fields could be dropped.

[jdoerfert]

Thanks for pushing this. Two comments below, nothing major.

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

137

Nit: I'm not a cmake person but I am surprised to see "cuda" above so many times, e.g. in file names?

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

114

Where did these implementations go?

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

137

The source is (currently still) cuda so it's not too bad that variables refer to cuda. I didn't rename anything from Greg's original script. Happy to rename the variables before or after commit.

I'd like to put the remaining cuda intrinsics behind an API and rename the files to .cpp, but the renaming caused problems with nvcc that I am not yet sure how to solve.

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

114

The previous implementation called through to functions in a support library with fairly ugly prototypes.

I'm planning to implement these in target_impl.hip, which doesn't exist yet. Until then they're just another couple of missing symbols.

[jdoerfert]

LGTM.

[ABataev]

Seems to me, it causes some problems at the compile-time:

/build_llvm/llvm-project/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h:27:21: fatal error: support.h: No such file or directory
 #include "support.h"
                     ^
compilation terminated.

[ABataev]

Reverted in 02b9c5d963c6c87a5dba46642c63bdb6b35901f1

[JonChesterfield]

In D70414#1767277, @ABataev wrote:

Seems to me, it causes some problems at the compile-time:

Ah. Sorry about this, a clean build would show the error locally for me. The fix is trivial, will do it now (plus the 20mins or so for a clean build to sanity check)

[JonChesterfield]

In D70414#1767357, @ABataev wrote:

Reverted in 02b9c5d963c6c87a5dba46642c63bdb6b35901f1

Fair enough. The patch is literally s/"support.h"/"common/support.h" in two, optionally three files.

[JonChesterfield]

@ABataev I reproduced locally, fixed, posted D70971. Would you mind checking it builds for you as a redundant check? Thanks

openmp/libomptarget/deviceRTLs/CMakeLists.txt

 ##===----------------------------------------------------------------------===##
 #
 # Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 # See https://llvm.org/LICENSE.txt for license information.
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 #
 # ##===----------------------------------------------------------------------===##
 #
-# Build a device RTL for each available machine available.
+# Build a device RTL for each available machine.
 #
 ##===----------------------------------------------------------------------===##
 
+add_subdirectory(amdgcn)
 add_subdirectory(nvptx)

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

+##===----------------------------------------------------------------------===##
+#
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+##===----------------------------------------------------------------------===##
+#
+# Build the AMDGCN Device RTL if the ROCM tools are available
+#
+##===----------------------------------------------------------------------===##
+
+find_package(LLVM QUIET CONFIG
+  PATHS
+  $ENV{AOMP}
+  $ENV{HOME}/rocm/aomp
+  /opt/rocm/aomp
+  /usr/lib/rocm/aomp
+  ${LIBOMPTARGET_NVPTX_CUDA_COMPILER_DIR}
+  ${LIBOMPTARGET_NVPTX_CUDA_LINKER_DIR}
+  ${CMAKE_CXX_COMPILER_DIR}
+  NO_DEFAULT_PATH)
+
+if (LLVM_DIR)
+  libomptarget_say("Found LLVM ${LLVM_PACKAGE_VERSION}. Configure: ${LLVM_DIR}/LLVMConfig.cmake")
+else()
+  libomptarget_say("Not building AMDGCN device RTL: AOMP not found")
+  return()
+endif()
+
+set(AOMP_INSTALL_PREFIX ${LLVM_INSTALL_PREFIX})
+
+if (AOMP_INSTALL_PREFIX)
+  set(AOMP_BINDIR ${AOMP_INSTALL_PREFIX}/bin)
+else()
+  set(AOMP_BINDIR ${LLVM_BUILD_BINARY_DIR}/bin)
+endif()
+
+libomptarget_say("Building AMDGCN device RTL. LLVM_COMPILER_PATH=${AOMP_BINDIR}")
+
+project(omptarget-amdgcn)
+
+add_custom_target(omptarget-amdgcn ALL)
+
+#optimization level
+set(optimization_level 2)
+
+# Activate RTL message dumps if requested by the user.
+if(LIBOMPTARGET_NVPTX_DEBUG)
+  set(CUDA_DEBUG -DOMPTARGET_NVPTX_DEBUG=-1)
+endif()
+
+get_filename_component(devicertl_base_directory
+  ${CMAKE_CURRENT_SOURCE_DIR}
+  DIRECTORY)
+
+set(cuda_sources
+  ${devicertl_base_directory}/common/src/cancel.cu
+  ${devicertl_base_directory}/common/src/critical.cu)
+
+set(h_files
+  ${CMAKE_CURRENT_SOURCE_DIR}/src/amdgcn_interface.h
+  ${CMAKE_CURRENT_SOURCE_DIR}/src/device_environment.h
+  ${CMAKE_CURRENT_SOURCE_DIR}/src/target_impl.h
+  ${devicertl_base_directory}/common/debug.h
+  ${devicertl_base_directory}/common/state-queue.h
+  ${devicertl_base_directory}/common/state-queuei.h
+  ${devicertl_base_directory}/common/support.h)
+
+# for both in-tree and out-of-tree build
+if (NOT CMAKE_ARCHIVE_OUTPUT_DIRECTORY)
+  set(OUTPUTDIR ${CMAKE_CURRENT_BINARY_DIR})
+else()
+  set(OUTPUTDIR ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
+endif()
+
+# create libraries
+set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900)
+if (DEFINED LIBOMPTARGET_AMDGCN_GFXLIST)
+  set(mcpus ${LIBOMPTARGET_AMDGCN_GFXLIST})
+endif()
+
+macro(add_cuda_bc_library)
+  set(cu_cmd ${AOMP_BINDIR}/clang++
+    -std=c++11
+    -fcuda-rdc
+    -fvisibility=default
+    --cuda-device-only
+    -Wno-unused-value
+    -x hip
+    -O${optimization_level}
+    --cuda-gpu-arch=${mcpu}
+    ${CUDA_DEBUG}
+    -I${CMAKE_CURRENT_SOURCE_DIR}/src
+    -I${devicertl_base_directory})
+
+  set(bc1_files)
+
+  foreach(file ${ARGN})
+    get_filename_component(fname ${file} NAME_WE)
+    set(bc1_filename ${fname}.${mcpu}.bc)
+
+    add_custom_command(
+      OUTPUT ${bc1_filename}
+      COMMAND ${cu_cmd} ${file} -o ${bc1_filename}
+      DEPENDS ${file} ${h_files})
+
+    list(APPEND bc1_files ${bc1_filename})
+  endforeach()
+
+  add_custom_command(
+    OUTPUT linkout.cuda.${mcpu}.bc
+    COMMAND ${AOMP_BINDIR}/llvm-link ${bc1_files} -o linkout.cuda.${mcpu}.bc
+    DEPENDS ${bc1_files})
+
+  list(APPEND bc_files linkout.cuda.${mcpu}.bc)
+endmacro()
+
+set(libname "omptarget-amdgcn")
+
+foreach(mcpu ${mcpus})
+  set(bc_files)
+  add_cuda_bc_library(${cuda_sources})
+
+  set(bc_libname lib${libname}-${mcpu}.bc)
+  add_custom_command(
+    OUTPUT ${bc_libname}
+    COMMAND ${AOMP_BINDIR}/llvm-link ${bc_files} | ${AOMP_BINDIR}/opt --always-inline -o ${OUTPUTDIR}/${bc_libname}
+    DEPENDS ${bc_files})
+
+  add_custom_target(lib${libname}-${mcpu} ALL DEPENDS ${bc_libname})
+
+  install(FILES ${OUTPUTDIR}/${bc_libname}
+     DESTINATION "${OPENMP_INSTALL_LIBDIR}/libdevice"
+  )
+endforeach()

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

+//===---- device_environment.h - OpenMP GPU device environment --- CUDA -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Global device environment
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _OMPTARGET_DEVICE_ENVIRONMENT_H_
+#define _OMPTARGET_DEVICE_ENVIRONMENT_H_
+
+#include "target_impl.h"
+
+struct omptarget_device_environmentTy {
+  int32_t debug_level;   // gets value of envvar LIBOMPTARGET_DEVICE_RTL_DEBUG
+                         // only useful for Debug build of deviceRTLs 
+  int32_t num_devices;   // gets number of active offload devices 
+  int32_t device_num;    // gets a value 0 to num_devices-1
+};
+
+extern DEVICE omptarget_device_environmentTy omptarget_device_environment;
+
+#endif

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

 EXTERN uint64_t __ballot64(int predicate);
 // initialized with a 64-bit mask with bits set in positions less than the
 // thread's lane number in the warp
 EXTERN uint64_t __lanemask_lt();
 // initialized with a 64-bit mask with bits set in positions greater than the
 // thread's lane number in the warp
 EXTERN uint64_t __lanemask_gt();
 
-EXTERN void llvm_amdgcn_s_barrier();
-
 // CU id
 EXTERN unsigned __smid();
 
 INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   lo = (uint32_t)(val & UINT64_C(0x00000000FFFFFFFF));
   hi = (uint32_t)((val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
 }
 

 INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
   return __lanemask_gt();
 }
 
 INLINE uint32_t __kmpc_impl_smid() {
   return __smid();
 }
 
-INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __ffsll(x); }
+INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __builtin_ffsl(x); }
 
-INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __popcll(x); }
+INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __builtin_popcountl(x); }
 
 INLINE __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   return __ballot64(1);
 }
 
-INLINE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
-                                     int32_t SrcLane) {
-  return __shfl(Var, SrcLane, WARPSIZE);
-}
+EXTERN int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
+                                     int32_t SrcLane);
 
-INLINE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
-                                          uint32_t Delta, int32_t Width) {
-  return __shfl_down(Var, Delta, Width);
-}
+EXTERN int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
+                                          uint32_t Delta, int32_t Width);
 
-INLINE void __kmpc_impl_syncthreads() { llvm_amdgcn_s_barrier(); }
+INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
 
 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.
   // So it is ok to issue a full workgroup barrier here.
   __builtin_amdgcn_s_barrier();
 }
 
+// DEVICE versions of part of libc
+extern "C" {
+DEVICE __attribute__((noreturn)) void
+__assertfail(const char *, const char *, unsigned, const char *, size_t);
+INLINE static void __assert_fail(const char *__message, const char *__file,
+                                 unsigned int __line, const char *__function) {
+  __assertfail(__message, __file, __line, __function, sizeof(char));
+}
+DEVICE int printf(const char *, ...);
+}
+
 #endif