Namespace Acts::Cuda

namespace Cuda

Typedefs

using device_array = std::unique_ptr<T, Details::DeviceArrayDeleter>

Convenience type for using primitive variable arrays on a CUDA device.

using host_array = std::unique_ptr<T, Details::HostArrayDeleter>

Convenience type for using primitive variable arrays on the host.

Functions

template<typename T>
void copyToDevice(device_array<T> &dev, const host_array<T> &host, std::size_t arraySize)

Copy one array from the host to the device.

template<typename T>
void copyToDevice(device_array<T> &dev, const host_array<T> &host, std::size_t arraySize, const StreamWrapper &stream)

Copy one array from the host to the device asynchronously.

template<typename T>
void copyToHost(host_array<T> &host, const device_array<T> &dev, std::size_t arraySize)

Copy one array from the device to the host.

template<typename T>
void copyToHost(host_array<T> &host, const device_array<T> &dev, std::size_t arraySize, const StreamWrapper &stream)

Copy one array from the device to the host asynchronously.

StreamWrapper createStreamFor(const Acts::Cuda::Info::Device &device)

Create a stream for a particular CUDA device.

template<typename T>
device_array<T> make_device_array(std::size_t size)

Function creating a primitive array in CUDA device memory.

template<typename T>
host_array<T> make_host_array(std::size_t size)

Function creating a primitive array in the host’s memory.

std::ostream &operator<<(std::ostream &out, const Info::Device &device)

Print operator for Acts::Cuda::Info::Device.

class Info
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Utilities/Info.hpp>

Class providing information about the CUDA devices at runtime.

Without exposing any CUDA dependencies publicly to the clients.

Declarations preventing any copies of the singleton object

Info(const Info&) = delete

Explicitly delete the copy constructor.

Info(Info&&) = delete

Explicitly delete the move constructor.

Info &operator=(const Info&) = delete

Explicitly delete the copy assignment operator.

Info &operator=(Info&&) = delete

Explicitly delete the move assignment operator.

Public Functions

const std::vector<Device> &devices() const

Get all the available CUDA devices.

Public Static Functions

static Info &instance()

Singleton accessor function.

struct Device
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Utilities/Info.hpp>

Helper struct describing one available CUDA device.

Public Members

bool concurrentKernels = false

Whether the device supports multiple kernel executions in parallel.

int id = -1

Identifier that CUDA knows this device by.

int maxThreadsPerBlock = -1

The maximal number of threads per block for this device.

std::string name

The name of this device.

std::size_t totalMemory = 0

The total amount of (global) memory on the device.

class MemoryManager
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Utilities/MemoryManager.hpp>

Singleton class used for allocating memory on CUDA device(s)

In order to avoid calling cudaMalloc(…) and cudaFree(…) too many times in the code (which can turn out to be pretty slow), device memory is allocated using this singleton memory manager for the Acts::Cuda::device_array arrays.

It is implemented in a very simple way. It allocates a big blob of memory, and then hands out pointers from this blob to anyone that asks for device memory.

The class doesn’t handle memory returns in any sophisticated way. It assumes that any calculation will need all allocated memory until the end of that calculation. At which point all of that memory gets re-purpused in one call.

The code is not thread safe currently in any shape or form. But there should be ways of making it at least “thread friendly” later on.

Declarations preventing any copies of the singleton object

MemoryManager(const MemoryManager&) = delete

Disallow copy construction.

MemoryManager(MemoryManager&&) = delete

Disallow move construction.

MemoryManager &operator=(const MemoryManager&) = delete

Disallow copy assignment.

MemoryManager &operator=(MemoryManager&&) = delete

Disallow move assignment.

Functions that the users of Acts may be interacting with

void setMemorySize(std::size_t sizeInBytes, int device = -1)

Set the amount of memory to use on a particular device.

static MemoryManager &instance()

Singleton object accessor.

Functions used internally by the Acts code

std::size_t availableMemory(int device = -1) const

Get the amount of memory still available on a specific device.

void *allocate(std::size_t sizeInBytes, int device = -1)

Get a pointer to an available memory block on the device.

void reset(int device = -1)

Reset all allocations.

Public Functions

~MemoryManager()

Destructor, freeing up all allocated memory.

template<typename external_spacepoint_t>
class SeedFinder
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Seeding2/SeedFinder.hpp>

Public Functions

SeedFinder(SeedfinderConfig<external_spacepoint_t> commonConfig, const SeedFilterConfig &seedFilterConfig, const TripletFilterConfig &tripletFilterConfig, int device = 0, std::unique_ptr<const Logger> logger = getDefaultLogger("Cuda::SeedFinder", Logging::INFO))

Create a CUDA backed seed finder object.

Parameters

  • commonConfig – Configuration shared with Acts::Seedfinder

  • seedFilterConfig – Configuration shared with Acts::SeedFilter

  • tripletFilterConfig – Configuration for the GPU based triplet filtering

  • device – The identifier of the CUDA device to run on

  • logger – A Logger instance

template<typename sp_range_t>
std::vector<Seed<external_spacepoint_t>> createSeedsForGroup(sp_range_t bottomSPs, sp_range_t middleSPs, sp_range_t topSPs) const

Create all seeds from the space points in the three iterators.

Can be used to parallelize the seed creation

Parameters

  • bottomSPs – group of space points to be used as innermost SP in a seed.

  • middleSPs – group of space points to be used as middle SP in a seed.

  • topSPs – group of space points to be used as outermost SP in a seed. Ranges must return pointers. Ranges must be separate objects for each parallel call.

Returns

vector in which all found seeds for this group are stored.

void setLogger(std::unique_ptr<const Logger> newLogger)

set logging instance

Parameters

newLogger – [in] is the logging istance to be set

class StreamWrapper
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Utilities/StreamWrapper.hpp>

Helper class for passing around cudaStream_t objects (pointers)

In order to be able to create user interfaces that return/receive CUDA streams, while not exposing the users of those interfaces to the CUDA Runtime API, this class helps us hiding the concrete CUDA types from our interfaces.

Public Functions

StreamWrapper(void *stream, bool ownsStream = true)

Declare the Acts::Cuda::getStreamFrom function a frient of the class.

Note that it’s not practical to put that function into the Acts::Cuda::details namespace, because then we would be forced to forward declare it in this header. Constructor with the stream to be wrapped

StreamWrapper(StreamWrapper &&parent)

Move constructor.

StreamWrapper(const StreamWrapper&) = delete

Disabled copy constructor.

~StreamWrapper()

Destructor.

StreamWrapper &operator=(StreamWrapper &&rhs)

Move assignment operator.

StreamWrapper &operator=(const StreamWrapper&) = delete

Disabled copy assignment operator.

void synchronize() const

Wait for all scheduled operations to finish in the stream.

struct TripletFilterConfig
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Seeding2/TripletFilterConfig.hpp>

Structure holding pointers to the user defined filter functions.

Public Types

typedef float (*seedWeightFunc_t)(const Details::SpacePoint&, const Details::SpacePoint&, const Details::SpacePoint&)

Type for the seed weighting functions.

typedef bool (*singleSeedCutFunc_t)(float, const Details::SpacePoint&, const Details::SpacePoint&, const Details::SpacePoint&)

Type for the seed filtering functions.

Public Members

seedWeightFunc_t seedWeight = nullptr

Pointer to a function assigning weights to seed candidates.

The function receives the bottom, middle and top spacepoints (in this order), and needs to return a float weight for the combination.

Note that you can not set this pointer directly. You must use cudaMemcpyFromSymbol to set it from a global function pointer.

singleSeedCutFunc_t singleSeedCut = nullptr

Pointer to a function filtering seed candidates.

The function receives a previously assigned “seed weight”, and references to the bottom, middle and top spacepoints (in this order). It needs to return an accept/reject decision for the combination.

Note that you can not set this pointer directly. You must use cudaMemcpyFromSymbol to set it from a global function pointer.

namespace Details

Namespace holding some implementation detail types that should not be used directly in client code.

Functions

DubletCounts countDublets(std::size_t maxBlockSize, std::size_t nMiddleSP, const device_array<unsigned int> &middleBottomCounts, const device_array<unsigned int> &middleTopCounts)

Calculate summary values for the dublet search.

After the dublet search is done, we need to know some information about how many duplets were found exactly. As this information is necessary for the scheduing of the subsequent steps of the execution on the GPU. This function is used to collect this information

Parameters

  • maxBlockSize – The maximum block size to use on the GPU

  • nMiddleSP – The number of middle spacepoints for which the dublet reconstruction was run

  • middleBottomCounts – The output from the Acts::Cuda::Details::findDublets(…) function with the same name

  • middleTopCounts – The output from the Acts::Cuda::Details::findDublets(…) function with the same name

Returns

An object holding all the summary statistics necessary for the subsequent steps of GPU execution

void findDublets(std::size_t maxBlockSize, std::size_t nBottomSPs, const device_array<SpacePoint> &bottomSPs, std::size_t nMiddleSPs, const device_array<SpacePoint> &middleSPs, std::size_t nTopSPs, const device_array<SpacePoint> &topSPs, float deltaRMin, float deltaRMax, float cotThetaMax, float collisionRegionMin, float collisionRegionMax, device_array<unsigned int> &middleBottomCounts, device_array<std::size_t> &middleBottomDublets, device_array<unsigned int> &middleTopCounts, device_array<std::size_t> &middleTopDublets)

Find all viable middle-bottom and middle-top dublets.

This function is run as the first step in the seed finding, looking for viable middle-bottom and middle-top spacepoint pairs for the subsequent steps of the code.

Note that middleBottomCounts and middleTopCounts have type “unsigned int” instead of “std::size_t”, because the GPU code needs to execute atomic operations on these arrays. And CUDA does not define such operations on std::size_t (i.e. unsigned long).

Parameters

  • maxBlockSize – [in] The maximum block size to use on the GPU

  • nBottomSPs – [in] The number of bottom spacepoints in bottomSPs

  • bottomSPs – [in] Properties of all of the bottom spacepoints

  • nMiddleSPs – [in] The number of middle spacepoints in middleSPs

  • middleSPs – [in] Properties of all of the middle spacepoints

  • nTopSPs – [in] The number of top spacepoints in topSPs

  • topSPs – [in] Properties of all of the top spacepoints

  • deltaRMin – [in] Configuration parameter from Acts::SeedfinderConfig

  • deltaRMax – [in] Configuration parameter from Acts::SeedfinderConfig

  • cotThetaMax – [in] Configuration parameter from Acts::SeedfinderConfig

  • collisionRegionMin – [in] Configuration parameter from Acts::SeedfinderConfig

  • collisionRegionMax – [in] Configuration parameter from Acts::SeedfinderConfig

  • middleBottomCounts – [out] 1-D array of the number of middle-bottom dublets found for each middle spacepoint

  • middleBottomDublets – [out] 2-D matrix of size nMiddleSPs x nBottomSPs, holding the bottom spacepoint indices for the identified middle-bottom dublets

  • middleTopCounts – [out] 1-D array of the number of middle-top dublets found for each middle spacepoint

  • middleTopDublets – [out] 2-D matrix of size nMiddleSPs x nTopSPs, holding the top spacepoint indices for the identified middle-top dublets

std::vector<std::vector<Triplet>> findTriplets(const Info::Device &device, std::size_t maxBlockSize, const DubletCounts &dubletCounts, const SeedFilterConfig &seedConfig, const TripletFilterConfig &filterConfig, std::size_t nBottomSPs, const device_array<SpacePoint> &bottomSPs, std::size_t nMiddleSPs, const device_array<SpacePoint> &middleSPs, std::size_t nTopSPs, const device_array<SpacePoint> &topSPs, const device_array<unsigned int> &middleBottomCounts, const device_array<std::size_t> &middleBottomDublets, const device_array<unsigned int> &middleTopCounts, const device_array<std::size_t> &middleTopDublets, float maxScatteringAngle2, float sigmaScattering, float minHelixDiameter2, float pT2perRadius, float impactMax)

Find all viable triplets from the provided spacepoint dublets.

This function is used to find a “loosely selected” set of seed candidates that still need to be filtered through Acts::SeedFilter::filterSeeds_1SpFixed before returning it to the user.

Parameters

  • device – [in] Properties of the device that the code will be running on

  • maxBlockSize – [in] The maximum block size to use on the GPU

  • dubletCounts – [in] The output object from Acts::Cuda::Details::countDublets

  • seedConfig – [in] Configuration parameters for the triplet finding/filtering

  • filterConfig – [in] User provided settings (code…) for the triplet filtering

  • nBottomSPs – [in] The number of bottom spacepoints in bottomSPs

  • bottomSPs – [in] Properties of all of the bottom spacepoints

  • nMiddleSPs – [in] The number of middle spacepoints in middleSPs

  • middleSPs – [in] Properties of all of the middle spacepoints

  • nTopSPs – [in] The number of top spacepoints in topSPs

  • topSPs – [in] Properties of all of the top spacepoints

  • middleBottomCounts – [in] 1-D array of the number of middle-bottom dublets found for each middle spacepoint

  • middleBottomDublets – [in] 2-D matrix of size nMiddleSPs x nBottomSPs, holding the bottom spacepoint indices for the identified middle-bottom dublets

  • middleTopCounts – [in] 1-D array of the number of middle-top dublets found for each middle spacepoint

  • middleTopDublets – [in] 2-D matrix of size nMiddleSPs x nTopSPs, holding the top spacepoint indices for the identified middle-top dublets

  • maxScatteringAngle2 – [in] Configuration parameter from Acts::SeedfinderConfig

  • sigmaScattering – [in] Configuration parameter from Acts::SeedfinderConfig

  • minHelixDiameter2 – [in] Configuration parameter from Acts::SeedfinderConfig

  • pT2perRadius – [in] Configuration parameter from Acts::SeedfinderConfig

  • impactMax – [in] Configuration parameter from Acts::SeedfinderConfig

Returns

A 2-D structure holding the parameters of the identified triplets for each middle spacepoint

class DeviceArrayDeleter
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Utilities/Arrays.hpp>

Class performing the deletion of a CUDA device memory array.

Public Functions

void operator()(void *ptr)

Operator performing the deletion of the memory.

struct DubletCounts
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Seeding2/Details/Types.hpp>

Helper struct summarising the results of the dublet search.

Public Members

unsigned int maxMBDublets = 0

The maximal number of middle-bottom dublets.

unsigned int maxMTDublets = 0

The maximal number of middle-top dublets.

unsigned int maxTriplets = 0

The maximal number of triplets for any middle SP.

unsigned int nDublets = 0

The total number of dublets (M-B and M-T) found.

unsigned int nTriplets = 0

The total number of triplet candidates found.

class HostArrayDeleter
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Utilities/Arrays.hpp>

Class performing the deletion of pinned host memory.

Public Functions

void operator()(void *ptr)

Operator performing the deletion of the memory.

struct LinCircle
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Seeding2/Details/Types.hpp>

Helper struct holding the linearly transformed coordinates of spacepoints.

Public Members

float cotTheta = 0.0f
float Er = 0.0f
float iDeltaR = 0.0f
float U = 0.0f
float V = 0.0f
float Zo = 0.0f
struct SpacePoint
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Seeding2/Details/Types.hpp>

Helper struct describing a spacepoint on the device.

Public Members

float radius = 0.0f

radius in beam system coordinates

float varianceR = 0.0f
float varianceZ = 0.0f
float x = 0.0f

x-coordinate in beam system coordinates

float y = 0.0f

y-coordinate in beam system coordinates

float z = 0.0f

z-coordinate in beam system coordinates

struct Triplet
#include </home/docs/checkouts/readthedocs.org/user_builds/acts/checkouts/v19.7.0/Plugins/Cuda/include/Acts/Plugins/Cuda/Seeding2/Details/Types.hpp>

Structure used in the CUDA-based triplet finding.

Public Members

unsigned int bottomIndex = static_cast<unsigned int>(-1)
float impactParameter = 0.0f
float invHelixDiameter = 0.0f
unsigned int middleIndex = static_cast<unsigned int>(-1)
unsigned int topIndex = static_cast<unsigned int>(-1)
float weight = 0.0f