sdfg::gpu::rocm Namespace Reference

Classes
struct	RocmMmaSupport
class	RocmArch
struct	RocmDeviceInfo
	A ROCm/HIP GPU device as reported by rocminfo: its LLVM/ROCm ISA (gfx) name together with the marketing name(s) of the device(s) that report it and the wavefront size. More...
class	RocmMmaBareStrategy
	Raw v_mfma_f32_32x32x8f16 emit strategy (see RocmMfma32Support). More...
class	RocmMmaFragmentStrategy
	Classic rocwmma fragment-based emit strategy (load/store/mma_sync/fill_fragment). More...
class	RocmWmmaLibDependency
class	GpuMmaTransform

Enumerations
enum class	RocmMmaVariant { Fragment = 1 , Bare = 2 }

Functions
const std::vector< RocmDeviceInfo > &	query_rocm_devices ()
	Queries the ROCm GPUs available on this machine via rocminfo.
const RocmArch *	rocm_arch_for_device (const RocmDeviceInfo &device)
	Resolves a queried device to a RocmArch.
const RocmArch *	rocm_arch_from_schedule_type (const structured_control_flow::ScheduleType &schedule)
	Wrapper, that allows us to store this inside the ScheduleType in the future.
const RocmArch *	rocm_arch_from_env ()
	Use the DOCC_ROCM_ARCH env var or if unset, query the available cards.
const RocmArch *	rocm_arch_from_available_hardware ()
const RocmArch *	rocm_arch_parse (const std::string &name)
data_flow::ImplementationType	ImplementationType_ROCM_MMA_BARE ("ROCM_MMA_BARE")
data_flow::ImplementationType	ImplementationType_ROCM_MMA_FRAGMENT ("ROCM_MMA")

Variables
RocmArch	ROCM_ARCH_GFX1201
RocmArch	ROCM_ARCH_GFX90A
RocmArch	ROCM_ARCH_GFX942

Enumeration Type Documentation

RocmMmaVariant

enum class sdfg::gpu::rocm::RocmMmaVariant

strong

Enumerator
Fragment	rocwmma fragment path (RocmMmaFragmentStrategy)
Bare	raw MFMA builtin path (RocmMmaBareStrategy)

Function Documentation

ImplementationType_ROCM_MMA_BARE()

data_flow::ImplementationType sdfg::gpu::rocm::ImplementationType_ROCM_MMA_BARE ( "ROCM_MMA_BARE" )

inline

ImplementationType_ROCM_MMA_FRAGMENT()

data_flow::ImplementationType sdfg::gpu::rocm::ImplementationType_ROCM_MMA_FRAGMENT ( "ROCM_MMA" )

inline

query_rocm_devices()

const std::vector< RocmDeviceInfo > & sdfg::gpu::rocm::query_rocm_devices

(

)

Queries the ROCm GPUs available on this machine via rocminfo.

Internally this invokes rocminfo and parses its per-agent report, keeping only GPU agents and matching their gfx… ISA name. The list is uniqued by gfx name (marketing names sharing a target are collected together).

The result is cached on first call, so rocminfo is spawned at most once per process. Returns an empty list if no ROCm device could be queried (e.g. rocminfo is not available).

rocm_arch_for_device()

const RocmArch * sdfg::gpu::rocm::rocm_arch_for_device

(

const RocmDeviceInfo &

device

)

Resolves a queried device to a RocmArch.

Prefers a curated arch (with correct MMA capabilities) when the gfx target is recognised; otherwise synthesises a custom arch from the fields rocminfo reports (name + wavefront, MMA disabled). The returned pointer has static lifetime.

rocm_arch_from_available_hardware()

const RocmArch * sdfg::gpu::rocm::rocm_arch_from_available_hardware

(

)

rocm_arch_from_env()

const RocmArch * sdfg::gpu::rocm::rocm_arch_from_env

(

)

Use the DOCC_ROCM_ARCH env var or if unset, query the available cards.

rocm_arch_from_schedule_type()

const RocmArch * sdfg::gpu::rocm::rocm_arch_from_schedule_type

(

const structured_control_flow::ScheduleType &

schedule

)

Wrapper, that allows us to store this inside the ScheduleType in the future.

For now, we look at the env-vars or query the local system

rocm_arch_parse()

const RocmArch * sdfg::gpu::rocm::rocm_arch_parse

(

const std::string &

name

)

Variable Documentation

ROCM_ARCH_GFX1201

RocmArch sdfg::gpu::rocm::ROCM_ARCH_GFX1201

extern

ROCM_ARCH_GFX90A

RocmArch sdfg::gpu::rocm::ROCM_ARCH_GFX90A

extern

ROCM_ARCH_GFX942

RocmArch sdfg::gpu::rocm::ROCM_ARCH_GFX942

extern