Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Improve support for HIP/ROCm and Alpaka #40678

Merged
merged 3 commits into from
Feb 5, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions DataFormats/Portable/interface/PortableCollection.h
Original file line number Diff line number Diff line change
Expand Up @@ -6,13 +6,13 @@
namespace traits {

// trait for a generic SoA-based product
template <typename T, typename TDev, typename = std::enable_if_t<cms::alpakatools::is_device_v<TDev>>>
template <typename T, typename TDev, typename = std::enable_if_t<alpaka::isDevice<TDev>>>
class PortableCollectionTrait;

} // namespace traits

// type alias for a generic SoA-based product
template <typename T, typename TDev, typename = std::enable_if_t<cms::alpakatools::is_device_v<TDev>>>
template <typename T, typename TDev, typename = std::enable_if_t<alpaka::isDevice<TDev>>>
using PortableCollection = typename traits::PortableCollectionTrait<T, TDev>::CollectionType;

#endif // DataFormats_Portable_interface_PortableCollection_h
7 changes: 4 additions & 3 deletions DataFormats/Portable/interface/PortableDeviceCollection.h
Original file line number Diff line number Diff line change
Expand Up @@ -5,12 +5,13 @@
#include <optional>
#include <type_traits>

#include <alpaka/alpaka.hpp>

#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"
#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"

// generic SoA-based product in device memory
template <typename T, typename TDev, typename = std::enable_if_t<cms::alpakatools::is_device_v<TDev>>>
template <typename T, typename TDev, typename = std::enable_if_t<alpaka::isDevice<TDev>>>
class PortableDeviceCollection {
static_assert(not std::is_same_v<TDev, alpaka_common::DevHost>,
"Use PortableHostCollection<T> instead of PortableDeviceCollection<T, DevHost>");
Expand All @@ -32,7 +33,7 @@ class PortableDeviceCollection {
assert(reinterpret_cast<uintptr_t>(buffer_->data()) % Layout::alignment == 0);
}

template <typename TQueue, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
template <typename TQueue, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
PortableDeviceCollection(int32_t elements, TQueue const& queue)
: buffer_{cms::alpakatools::make_device_buffer<std::byte[]>(queue, Layout::computeDataSize(elements))},
layout_{buffer_->data(), elements},
Expand Down
5 changes: 3 additions & 2 deletions DataFormats/Portable/interface/PortableHostCollection.h
Original file line number Diff line number Diff line change
Expand Up @@ -4,10 +4,11 @@
#include <cassert>
#include <optional>

#include <alpaka/alpaka.hpp>

#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/host.h"
#include "HeterogeneousCore/AlpakaInterface/interface/memory.h"
#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"

// generic SoA-based product in host memory
template <typename T>
Expand All @@ -30,7 +31,7 @@ class PortableHostCollection {
assert(reinterpret_cast<uintptr_t>(buffer_->data()) % Layout::alignment == 0);
}

template <typename TQueue, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
template <typename TQueue, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
PortableHostCollection(int32_t elements, TQueue const& queue)
// allocate pinned host memory associated to the given work queue, accessible by the queue's device
: buffer_{cms::alpakatools::make_host_buffer<std::byte[]>(queue, Layout::computeDataSize(elements))},
Expand Down
5 changes: 3 additions & 2 deletions DataFormats/Portable/interface/Product.h
Original file line number Diff line number Diff line change
Expand Up @@ -4,10 +4,11 @@
#include <memory>
#include <utility>

#include <alpaka/alpaka.hpp>

#include "DataFormats/Portable/interface/ProductBase.h"
#include "HeterogeneousCore/AlpakaInterface/interface/ScopedContextFwd.h"
#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"

namespace edm {
template <typename T>
Expand All @@ -30,7 +31,7 @@ namespace cms::alpakatools {
* it. Here is a somewhat natural place. If the overhead is too much, we
* can use them only where synchronization between queues is needed.
*/
template <typename TQueue, typename T, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
template <typename TQueue, typename T, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
class Product : public ProductBase<TQueue> {
public:
using Queue = TQueue;
Expand Down
3 changes: 1 addition & 2 deletions DataFormats/Portable/interface/ProductBase.h
Original file line number Diff line number Diff line change
Expand Up @@ -8,15 +8,14 @@
#include <alpaka/alpaka.hpp>

#include "HeterogeneousCore/AlpakaInterface/interface/ScopedContextFwd.h"
#include "HeterogeneousCore/AlpakaInterface/interface/traits.h"

namespace cms::alpakatools {

/**
* Base class for all instantiations of Product<TQueue, T> to hold the
* non-T-dependent members.
*/
template <typename TQueue, typename = std::enable_if_t<cms::alpakatools::is_queue_v<TQueue>>>
template <typename TQueue, typename = std::enable_if_t<alpaka::isQueue<TQueue>>>
class ProductBase {
public:
using Queue = TQueue;
Expand Down
8 changes: 7 additions & 1 deletion DataFormats/SoATemplate/interface/SoACommon.h
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@
#include "FWCore/Utilities/interface/typedefs.h"

// CUDA attributes
#ifdef __CUDACC__
#if defined(__CUDACC__) || defined(__HIPCC__)
#define SOA_HOST_ONLY __host__
#define SOA_DEVICE_ONLY __device__
#define SOA_HOST_DEVICE __host__ __device__
Expand All @@ -35,6 +35,12 @@
printf("%s\n", (A)); \
__trap(); \
}
#elif defined(__HIPCC__) && defined(__HIP_DEVICE_COMPILE__)
#define SOA_THROW_OUT_OF_RANGE(A) \
{ \
printf("%s\n", (A)); \
abort(); \
}
#else
#define SOA_THROW_OUT_OF_RANGE(A) \
{ throw std::out_of_range(A); }
Expand Down
25 changes: 17 additions & 8 deletions DataFormats/SoATemplate/test/BuildFile.xml
Original file line number Diff line number Diff line change
@@ -1,17 +1,26 @@
<iftool name="cuda-gcc-support">
<bin file="SoALayoutAndView_t.cu" name="SoALayoutAndView_t">
<bin file="SoALayoutAndView_t.cu" name="testCudaSoALayoutAndView_t">
<use name="boost"/>
<use name="cuda"/>
<use name="eigen"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
</bin>
</iftool>

<!-- dictionaries for FakeSoA -->
<library file="" name="FakeSoADict"/>

<!-- This test triggers a bug in ROOT, so it's kept disabled
<bin file="SoAStreamer_t.cpp" name="SoAStreamer_t">
<use name="root"/>
<iftool name="rocm">
<bin file="SoALayoutAndView_t.hip.cc" name="testRocmSoALayoutAndView_t">
<use name="boost"/>
<use name="eigen"/>
<use name="rocm"/>
<use name="HeterogeneousCore/ROCmUtilities"/>
</bin>
-->
</iftool>

<!-- This test triggers a bug in ROOT, so it's kept disabled
<bin file="SoAStreamer_t.cpp" name="SoAStreamer_t">
<use name="root"/>
</bin>
-->

<!-- dictionaries for FakeSoA -->
<library file="" name="FakeSoADict"/>
Loading