DeviceRuntimeApi.hh File Reference

Include CUDA or HIP runtime APIs for compiling with host/cc compiler. More...

#include "corecel/Config.hh"
#include "corecel/Assert.hh"

Macros
#define	CELER_DEVICE_PLATFORM   none
	API prefix token for the device offloading type.
#define	CELER_DEVICE_PLATFORM_UPPER_STR   ""
#define	CELER_DEVICE_API_SYMBOL(TOK)   void
	Add a prefix "hip" or "cuda" to a code token.
#define	CELER_EU_PER_CU   0
	Execution units per compute unit.
#define	CELER_DEVICE_RUNTIME_INCLUDED   1
	This macro informs downstream Celeritas code (namely, Stream and DeviceEvent) that it's safe to use types from the device APIs.
#define	CELER_DEVICE_API_CALL(STMT)
	Safely and portably dispatch a CUDA/HIP API call.
#define	CELER_DEVICE_PREFIX(STMT)   CELER_DEVICE_API_SYMBOL(STMT)
#define	CELER_DEVICE_CALL_PREFIX(STMT)   CELER_DEVICE_API_CALL(STMT)
#define	CELER_DEVICE_CHECK_ERROR()   CELER_DEVICE_API_CALL(PeekAtLastError())
#define	CELER_USE_THRUST   0
#define	CELER_CUB_HAS_TRANSFORM   0
#define	CELER_CUB_HAS_FLAGGEDIF   0
#define	CELER_HIPCUB_HAS_TRANSFORM   0
	Determine if CUB or hipCUB is available, and if so, check the version.

Detailed Description

Include CUDA or HIP runtime APIs for compiling with host/cc compiler.

Macro Definition Documentation

CELER_DEVICE_API_CALL

#define CELER_DEVICE_API_CALL

(

STMT

)

Value:

        do \
        { \
            CELER_NOT_CONFIGURED("CUDA or HIP"); \
        } while (0)

Safely and portably dispatch a CUDA/HIP API call.

When CUDA or HIP support is enabled, execute the wrapped statement prepend the argument with "cuda" or "hip" and throw a RuntimeError if it fails. If no device platform is enabled, throw an unconfigured assertion.

Example:

CELER_DEVICE_API_CALL(Malloc(&ptr_gpu, 100 * sizeof(float)));
CELER_DEVICE_API_CALL(DeviceSynchronize());

CELER_EU_PER_CU

#define CELER_EU_PER_CU   0

Execution units per compute unit.

AMD multiprocessors each have 4 SIMD units per compute unit, but there is no device attribute or compiler definition that provides this information. For CUDA, each streaming multiprocessor (MP) is a single "execution unit" and a "compute unit".

CELER_HIPCUB_HAS_TRANSFORM

#define CELER_HIPCUB_HAS_TRANSFORM   0

Determine if CUB or hipCUB is available, and if so, check the version.

CUDA has included CUB since CUDA 11, but ROCm does not include hipCUB by default, so test for the availability of hipCUB and use thrust instead if it's unavailable.

DeviceTransform::Transform is unavailable in earlier versions of CUB/hipCUB, so the code uses thrust::transform in that case.

DeviceSelect::FlaggedIf is unavailable in earlier versions of CUB and doesn't work with hipCUB versions 3.4.0 through 4.1.0 when using celeritas OpaqueId data types, so use a transform and DeviceSelect::Flagged instead.

Note

These are unneeded and undefined if this isn't a HIP or CUDA source file. This check is necessary because the version header file needs to be included.