GPU RTC: Add RTCoverrideArchitecture option to override RTC architecture compile command line part

Author: davidrohr

GPU/GPUTracking/Base/cuda/CMakeLists.txt

@@ -39,9 +39,10 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT")
 
   # build flags to use for RTC
   set(GPU_RTC_FLAGS "${CMAKE_CUDA_FLAGS} ${CMAKE_CUDA_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -std=c++${CMAKE_CUDA_STANDARD}")
+  set(GPU_RTC_FLAGS_ARCH "")
   if(CUDA_COMPUTETARGET)
       foreach(CUDA_ARCH ${CUDA_COMPUTETARGET})
-          set(GPU_RTC_FLAGS "${GPU_RTC_FLAGS} -gencode arch=compute_${CUDA_ARCH},code=sm_${CUDA_ARCH}")
+          set(GPU_RTC_FLAGS_ARCH "${GPU_RTC_FLAGS_ARCH} -gencode arch=compute_${CUDA_ARCH},code=sm_${CUDA_ARCH}")
       endforeach()
       list (GET CUDA_COMPUTETARGET 0 RTC_CUDA_ARCH)
       set(RTC_CUDA_ARCH "${RTC_CUDA_ARCH}0")
@@ -85,7 +86,16 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT")
   )
   create_binary_resource(${GPU_RTC_BIN}.command ${GPU_RTC_BIN}.command.o)
 
-  set(SRCS ${SRCS} ${GPU_RTC_BIN}.src.o ${GPU_RTC_BIN}.command.o)
+  add_custom_command(
+      OUTPUT ${GPU_RTC_BIN}.command.arch
+      COMMAND echo -n "${GPU_RTC_FLAGS_ARCH}" > ${GPU_RTC_BIN}.command.arch
+      COMMAND_EXPAND_LISTS
+      VERBATIM
+      COMMENT "Preparing CUDA RTC ARCH file ${GPU_RTC_BIN}.command.arch"
+  )
+  create_binary_resource(${GPU_RTC_BIN}.command.arch ${GPU_RTC_BIN}.command.arch.o)
+
+  set(SRCS ${SRCS} ${GPU_RTC_BIN}.src.o ${GPU_RTC_BIN}.command.o ${GPU_RTC_BIN}.command.arch.o)
 endif()
 # -------------------------------- End RTC -------------------------------------------------------
 

GPU/GPUTracking/Base/cuda/GPUReconstructionCUDAGenRTC.cxx

@@ -33,6 +33,7 @@ using namespace GPUCA_NAMESPACE::gpu;
 #include "utils/qGetLdBinarySymbols.h"
 QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_src);
 QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_command);
+QGET_LD_BINARY_SYMBOLS(GPUReconstructionCUDArtc_command_arch);
 #endif
 
 int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
@@ -53,12 +54,16 @@ int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
     kernelsall += kernels[i] + "\n";
   }
 
+  std::string baseCommand = (mProcessingSettings.RTCprependCommand != "" ? (mProcessingSettings.RTCprependCommand + " ") : "");
+  baseCommand += (getenv("O2_GPU_RTC_OVERRIDE_CMD") ? std::string(getenv("O2_GPU_RTC_OVERRIDE_CMD")) : std::string(_binary_GPUReconstructionCUDArtc_command_start, _binary_GPUReconstructionCUDArtc_command_len));
+  baseCommand += std::string(" ") + (mProcessingSettings.RTCoverrideArchitecture != "" ? mProcessingSettings.RTCoverrideArchitecture : std::string(_binary_GPUReconstructionCUDArtc_command_arch_start, _binary_GPUReconstructionCUDArtc_command_arch_len));
+
 #ifdef GPUCA_HAVE_O2HEADERS
   char shasource[21], shaparam[21], shacmd[21], shakernels[21];
   if (mProcessingSettings.rtc.cacheOutput) {
     o2::framework::internal::SHA1(shasource, _binary_GPUReconstructionCUDArtc_src_start, _binary_GPUReconstructionCUDArtc_src_len);
     o2::framework::internal::SHA1(shaparam, rtcparam.c_str(), rtcparam.size());
-    o2::framework::internal::SHA1(shacmd, _binary_GPUReconstructionCUDArtc_command_start, _binary_GPUReconstructionCUDArtc_command_len);
+    o2::framework::internal::SHA1(shacmd, baseCommand.c_str(), baseCommand.size());
     o2::framework::internal::SHA1(shakernels, kernelsall.c_str(), kernelsall.size());
   }
 #endif
@@ -159,8 +164,6 @@ int GPUReconstructionCUDA::genRTC(std::string& filename, unsigned int& nCompile)
     }
     HighResTimer rtcTimer;
     rtcTimer.ResetStart();
-    std::string baseCommand = (mProcessingSettings.RTCprependCommand != "" ? (mProcessingSettings.RTCprependCommand + " ") : "");
-    baseCommand += (getenv("O2_GPU_RTC_OVERRIDE_CMD") ? std::string(getenv("O2_GPU_RTC_OVERRIDE_CMD")) : std::string(_binary_GPUReconstructionCUDArtc_command_start, _binary_GPUReconstructionCUDArtc_command_len));
 #ifdef WITH_OPENMP
 #pragma omp parallel for schedule(dynamic, 1)
 #endif

GPU/GPUTracking/Base/hip/CMakeLists.txt

@@ -80,8 +80,9 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT")
 
   # build flags to use for RTC
   set(GPU_RTC_FLAGS "${CMAKE_HIP_FLAGS} ${CMAKE_HIP_FLAGS_${CMAKE_BUILD_TYPE_UPPER}} -std=c++${CMAKE_HIP_STANDARD}")
+  set(GPU_RTC_FLAGS_ARCH "")
   foreach(HIP_ARCH ${CMAKE_HIP_ARCHITECTURES})
-    set(GPU_RTC_FLAGS "${GPU_RTC_FLAGS} --offload-arch=${HIP_ARCH}")
+    set(GPU_RTC_FLAGS_ARCH "${GPU_RTC_FLAGS_ARCH} --offload-arch=${HIP_ARCH}")
   endforeach()
 
   set(GPU_RTC_FLAGS_SEPARATED "${GPU_RTC_FLAGS}")
@@ -118,7 +119,16 @@ if(NOT ALIGPU_BUILD_TYPE STREQUAL "ALIROOT")
   )
   create_binary_resource(${GPU_RTC_BIN}.command ${GPU_RTC_BIN}.command.o)
 
-  set(SRCS ${SRCS} ${GPU_RTC_BIN}.src.o ${GPU_RTC_BIN}.command.o)
+  add_custom_command(
+      OUTPUT ${GPU_RTC_BIN}.command.arch
+      COMMAND echo -n "${GPU_RTC_FLAGS_ARCH}" > ${GPU_RTC_BIN}.command.arch
+      COMMAND_EXPAND_LISTS
+      VERBATIM
+      COMMENT "Preparing HIP RTC ARCH file ${GPU_RTC_BIN}.command.arch"
+  )
+  create_binary_resource(${GPU_RTC_BIN}.command.arch ${GPU_RTC_BIN}.command.arch.o)
+
+  set(SRCS ${SRCS} ${GPU_RTC_BIN}.src.o ${GPU_RTC_BIN}.command.o ${GPU_RTC_BIN}.command.arch.o)
 endif()
 # -------------------------------- End RTC -------------------------------------------------------
 

GPU/GPUTracking/Base/hip/per_kernel/CMakeLists.txt

@@ -11,5 +11,5 @@
 
 add_library(GPUTrackingHIPKernels OBJECT $<JOIN:$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<LIST:TRANSFORM,$<TARGET_PROPERTY:O2_GPU_KERNELS,O2_GPU_KERNEL_NAMES>,REPLACE,[^A-Za-z0-9]+,_>,PREPEND,${O2_GPU_KERNEL_WRAPPER_FOLDER}/krnl_>,APPEND,.hip.cxx>, >)
 set(CMAKE_CXX_COMPILER ${hip_HIPCC_EXECUTABLE})
-set(CMAKE_CXX_FLAGS "${GPU_RTC_FLAGS} --genco")
+set(CMAKE_CXX_FLAGS "${GPU_RTC_FLAGS} ${GPU_RTC_FLAGS_ARCH} --genco")
 unset(CMAKE_CXX_FLAGS_${CMAKE_BUILD_TYPE_UPPER})

GPU/GPUTracking/Definitions/GPUSettingsList.h

@@ -283,6 +283,7 @@ AddOption(tpcMaxAttachedClustersPerSectorRow, unsigned int, 51000, "", 0, "Maxim
 AddOption(tpcUseOldCPUDecoding, bool, false, "", 0, "Enable old CPU-based TPC decoding")
 AddOption(RTCcacheFolder, std::string, "./rtccache/", "", 0, "Folder in which the cache file is stored")
 AddOption(RTCprependCommand, std::string, "", "", 0, "Prepend RTC compilation commands by this string")
+AddOption(RTCoverrideArchitecture, std::string, "", "", 0, "Override arhcitecture part of RTC compilation command line")
 AddOption(printSettings, bool, false, "", 0, "Print all settings when initializing")
 AddVariable(eventDisplay, GPUCA_NAMESPACE::gpu::GPUDisplayFrontendInterface*, nullptr)
 AddSubConfig(GPUSettingsProcessingRTC, rtc)