Merge pull request cms-sw#40832 from fwyzard/more_Alpaka_updates_131x

CUDA, ROCm and Alpaka-related updates

Author: cmsbuild

Alignment/CommonAlignment/python/tools/trackselectionRefitting.py

@@ -363,10 +363,6 @@ def getSequence(process, collection,
     ## put the sequence together ##
     ###############################
 
-    if "Fast" in TTRHBuilder:
-        print("PixelCPEFast has been chosen, here we must include CUDAService first")
-        process.load('HeterogeneousCore.CUDAServices.CUDAService_cfi')
-
     modules = []
     src = collection
     prevsrc = None

Configuration/StandardSequences/python/Accelerators_cff.py

@@ -4,3 +4,4 @@
 # used in production
 
 from HeterogeneousCore.CUDACore.ProcessAcceleratorCUDA_cfi import ProcessAcceleratorCUDA
+from HeterogeneousCore.ROCmCore.ProcessAcceleratorROCm_cfi import ProcessAcceleratorROCm

DataFormats/PortableTestObjects/src/alpaka/classes_rocm.h

@@ -0,0 +1,5 @@
+#include "DataFormats/Common/interface/DeviceProduct.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+#include "DataFormats/Portable/interface/Product.h"
+#include "DataFormats/PortableTestObjects/interface/TestSoA.h"
+#include "DataFormats/PortableTestObjects/interface/alpaka/TestDeviceCollection.h"

DataFormats/PortableTestObjects/src/alpaka/classes_rocm_def.xml

@@ -0,0 +1,5 @@
+<lcgdict>
+  <class name="alpaka_rocm_async::portabletest::TestDeviceCollection" persistent="false"/>
+  <class name="edm::DeviceProduct<alpaka_rocm_async::portabletest::TestDeviceCollection>" persistent="false"/>
+  <class name="edm::Wrapper<edm::DeviceProduct<alpaka_rocm_async::portabletest::TestDeviceCollection>>" persistent="false"/>
+</lcgdict>

EventFilter/HcalRawToDigi/plugins/HcalDigisProducerGPU.cc

@@ -9,7 +9,7 @@
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
 #include "FWCore/ServiceRegistry/interface/Service.h"
 #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
-#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAInterface.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
 
 class HcalDigisProducerGPU : public edm::stream::EDProducer<edm::ExternalWork> {
@@ -97,8 +97,8 @@ HcalDigisProducerGPU::HcalDigisProducerGPU(const edm::ParameterSet& ps)
   hf3_.stride = hcal::compute_stride<hcal::Flavor3>(QIE11DigiCollection::MAXSAMPLES);
 
   // preallocate pinned host memory only if CUDA is available
-  edm::Service<CUDAService> cs;
-  if (cs and cs->enabled()) {
+  edm::Service<CUDAInterface> cuda;
+  if (cuda and cuda->enabled()) {
     hf01_.reserve(config_.maxChannelsF01HE);
     hf5_.reserve(config_.maxChannelsF5HB);
     hf3_.reserve(config_.maxChannelsF3HB);

EventFilter/HcalRawToDigi/plugins/HcalRawToDigiGPU.cc

@@ -1,5 +1,4 @@
-#include <iostream>
-
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CondFormats/DataRecord/interface/HcalElectronicsMapRcd.h"
 #include "DataFormats/FEDRawData/interface/FEDNumbering.h"
 #include "DataFormats/FEDRawData/interface/FEDRawDataCollection.h"
@@ -8,11 +7,7 @@
 #include "FWCore/Framework/interface/MakerMacros.h"
 #include "FWCore/Framework/interface/stream/EDProducer.h"
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
-#include "FWCore/ServiceRegistry/interface/Service.h"
 #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
-#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 
 #include "DeclsForKernels.h"
 #include "DecodeGPU.h"

HeterogeneousCore/AlpakaCore/python/ProcessAcceleratorAlpaka.py

@@ -1,11 +1,17 @@
 import FWCore.ParameterSet.Config as cms
 
+import os
+
+from HeterogeneousCore.Common.PlatformStatus import PlatformStatus
+
 class ModuleTypeResolverAlpaka:
     def __init__(self, accelerators, backend):
-        # first element is used as the default is nothing is set
+        # first element is used as the default if nothing is set
         self._valid_backends = []
         if "gpu-nvidia" in accelerators:
             self._valid_backends.append("cuda_async")
+        if "gpu-amd" in accelerators:
+            self._valid_backends.append("rocm_async")
         if "cpu" in accelerators:
             self._valid_backends.append("serial_sync")
         if len(self._valid_backends) == 0:
@@ -45,26 +51,64 @@ class ProcessAcceleratorAlpaka(cms.ProcessAccelerator):
     ProcessAcceleratorCUDA) define.
     """
     def __init__(self):
-        super(ProcessAcceleratorAlpaka,self).__init__()
+        super(ProcessAcceleratorAlpaka, self).__init__()
         self._backend = None
+
     # User-facing interface
     def setBackend(self, backend):
         self._backend = backend
+
     # Framework-facing interface
     def moduleTypeResolver(self, accelerators):
         return ModuleTypeResolverAlpaka(accelerators, self._backend)
+
     def apply(self, process, accelerators):
-        if not hasattr(process, "AlpakaServiceSerialSync"):
+        # Propagate the AlpakaService messages through the MessageLogger
+        if not hasattr(process.MessageLogger, "AlpakaService"):
+            process.MessageLogger.AlpakaService = cms.untracked.PSet()
+
+        # Check if the CPU backend is available
+        try:
+            if not "cpu" in accelerators:
+                raise False
             from HeterogeneousCore.AlpakaServices.AlpakaServiceSerialSync_cfi import AlpakaServiceSerialSync
-            process.add_(AlpakaServiceSerialSync)
-        if not hasattr(process, "AlpakaServiceCudaAsync"):
+        except:
+            # the CPU backend is not available, do not load the AlpakaServiceSerialSync
+            if hasattr(process, "AlpakaServiceSerialSync"):
+                del process.AlpakaServiceSerialSync
+        else:
+            # the CPU backend is available, ensure the AlpakaServiceSerialSync is loaded
+            if not hasattr(process, "AlpakaServiceSerialSync"):
+                process.add_(AlpakaServiceSerialSync)
+
+        # Check if CUDA is available, and if the system has at least one usable NVIDIA GPU
+        try:
+            if not "gpu-nvidia" in accelerators:
+                raise False
             from HeterogeneousCore.AlpakaServices.AlpakaServiceCudaAsync_cfi import AlpakaServiceCudaAsync
-            process.add_(AlpakaServiceCudaAsync)
+        except:
+            # CUDA is not available, do not load the AlpakaServiceCudaAsync
+            if hasattr(process, "AlpakaServiceCudaAsync"):
+                del process.AlpakaServiceCudaAsync
+        else:
+            # CUDA is available, ensure the AlpakaServiceCudaAsync is loaded
+            if not hasattr(process, "AlpakaServiceCudaAsync"):
+                process.add_(AlpakaServiceCudaAsync)
 
-        if not hasattr(process.MessageLogger, "AlpakaService"):
-            process.MessageLogger.AlpakaService = cms.untracked.PSet()
+        # Check if ROCm is available, and if the system has at least one usable AMD GPU
+        try:
+            if not "gpu-amd" in accelerators:
+                raise False
+            from HeterogeneousCore.AlpakaServices.AlpakaServiceROCmAsync_cfi import AlpakaServiceROCmAsync
+        except:
+            # ROCm is not available, do not load the AlpakaServiceROCmAsync
+            if hasattr(process, "AlpakaServiceROCmAsync"):
+                del process.AlpakaServiceROCmAsync
+        else:
+            # ROCm is available, ensure the AlpakaServiceROCmAsync is loaded
+            if not hasattr(process, "AlpakaServiceROCmAsync"):
+                process.add_(AlpakaServiceROCmAsync)
 
-        process.AlpakaServiceSerialSync.enabled = "cpu" in accelerators
-        process.AlpakaServiceCudaAsync.enabled = "gpu-nvidia" in accelerators
 
+# Ensure this module is kept in the configuration when dumping it
 cms.specialImportRegistry.registerSpecialImportForType(ProcessAcceleratorAlpaka, "from HeterogeneousCore.AlpakaCore.ProcessAcceleratorAlpaka import ProcessAcceleratorAlpaka")

HeterogeneousCore/AlpakaCore/src/module_backend_config.cc

@@ -17,7 +17,7 @@ namespace cms::alpakatools {
     descAlpaka.addUntracked<std::string>("backend", "")
         ->setComment(
             "Alpaka backend for this module. Can be empty string (for the global default), 'serial_sync', or "
-            "'cuda_async'");
+            " - depending on the architecture and available hardware - 'cuda_async', 'rocm_async'");
 
     if (iDesc.defaultDescription()) {
       if (iDesc.defaultDescription()->isLabelUnused(kPSetName)) {

HeterogeneousCore/AlpakaServices/src/alpaka/AlpakaService.cc

@@ -16,12 +16,12 @@
 
 #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
 #include "FWCore/ServiceRegistry/interface/Service.h"
-#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAInterface.h"
 #endif  // ALPAKA_ACC_GPU_CUDA_ENABLED
 
 #ifdef ALPAKA_ACC_GPU_HIP_ENABLED
 #include "FWCore/ServiceRegistry/interface/Service.h"
-#include "HeterogeneousCore/ROCmServices/interface/ROCmService.h"
+#include "HeterogeneousCore/ROCmServices/interface/ROCmInterface.h"
 #endif  // ALPAKA_ACC_GPU_HIP_ENABLED
 
 namespace ALPAKA_ACCELERATOR_NAMESPACE {
@@ -31,11 +31,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
         verbose_(config.getUntrackedParameter<bool>("verbose")) {
 #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
     // rely on the CUDAService to initialise the CUDA devices
-    edm::Service<CUDAService> cudaService;
+    edm::Service<CUDAInterface> cuda;
 #endif  // ALPAKA_ACC_GPU_CUDA_ENABLED
 #ifdef ALPAKA_ACC_GPU_HIP_ENABLED
     // rely on the ROCmService to initialise the ROCm devices
-    edm::Service<ROCmService> rocmService;
+    edm::Service<ROCmInterface> rocm;
 #endif  // ALPAKA_ACC_GPU_HIP_ENABLED
 
     // TODO from Andrea Bocci:
@@ -48,14 +48,14 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
     }
 
 #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
-    if (not cudaService->enabled()) {
+    if (not cuda or not cuda->enabled()) {
       enabled_ = false;
       edm::LogInfo("AlpakaService") << ALPAKA_TYPE_ALIAS_NAME(AlpakaService) << " disabled by CUDAService";
       return;
     }
 #endif  // ALPAKA_ACC_GPU_CUDA_ENABLED
 #ifdef ALPAKA_ACC_GPU_HIP_ENABLED
-    if (not rocmService->enabled()) {
+    if (not rocm or not rocm->enabled()) {
       enabled_ = false;
       edm::LogInfo("AlpakaService") << ALPAKA_TYPE_ALIAS_NAME(AlpakaService) << " disabled by ROCmService";
       return;

HeterogeneousCore/CUDACore/README.md

@@ -83,7 +83,14 @@ This page documents the CUDA integration within CMSSW
        stream must synchronize with the work queued on other CUDA
        streams (with CUDA events and `cudaStreamWaitEvent()`)
 4. Outside of `acquire()`/`produce()`, CUDA API functions may be
-   called only if `CUDAService::enabled()` returns `true`.
+   called only if the `CUDAService` implementation of the `CUDAInterface`
+   is available and `CUDAService::enabled()` returns `true`:
+     ```c++
+     edm::Service<CUDAInterface> cuda;
+     if (cuda and cuda->enabled()) {
+       // CUDA calls ca be made here
+     }
+     ```
    * With point 3 it follows that in these cases multiple devices have
      to be dealt with explicitly, as well as CUDA streams