tilelang.contrib.hip_resource_info

Parse AMD GPU per-kernel resource usage out of clang’s -Rpass-analysis=kernel-resource-usage remarks and expose them on JITKernel.

clang emits a block like::

remark: src.cc:9:0: Function Name: main_kernel [-Rpass-analysis=kernel-resource-usage] remark: src.cc:9:0: TotalSGPRs: 16 [-Rpass-analysis=kernel-resource-usage] remark: src.cc:9:0: VGPRs: 5 [-Rpass-analysis=kernel-resource-usage] remark: src.cc:9:0: ScratchSize [bytes/lane]: 0 [-Rpass-analysis=kernel-resource-usage] remark: src.cc:9:0: SGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage] remark: src.cc:9:0: VGPRs Spill: 0 [-Rpass-analysis=kernel-resource-usage] …

right alongside any real warnings/errors. We parse and strip those lines before printing or raising, so autotune logs don’t drown in hundreds of remark blocks while real diagnostics still surface.

Classes

KernelResourceUsage

Resource counts as reported by clang's kernel-resource-usage pass.

Functions

hipcc_remark_flag()	The clang flag callers should pass to hipcc to enable the remark
reset_recorder()	Begin a fresh recording window on this thread.
pop_recorded()	Return everything recorded since the last reset_recorder and
filter_and_record(output)	Strip kernel-resource-usage remarks from output, parse them,
dump_to_file(usage, path)	Persist parsed resource usage so it survives kernel-cache hits.
load_from_file(path)	Inverse of dump_to_file. Tolerant of missing / unknown fields

Module Contents

class tilelang.contrib.hip_resource_info.KernelResourceUsage

Resource counts as reported by clang’s kernel-resource-usage pass.

Field names mirror the remark labels (lower-cased, normalized) so we can extend without breaking callers.

n_regs: int = 0

n_spills: int = 0

scratch_bytes: int = 0

n_max_threads: int | None = None

extra: dict[str, str]

tilelang.contrib.hip_resource_info.hipcc_remark_flag()

The clang flag callers should pass to hipcc to enable the remark output we parse here.

Return type:

str

tilelang.contrib.hip_resource_info.reset_recorder()

Begin a fresh recording window on this thread.

Return type:

None

tilelang.contrib.hip_resource_info.pop_recorded()

Return everything recorded since the last reset_recorder and clear the buffer.

Return type:

dict[str, KernelResourceUsage]

tilelang.contrib.hip_resource_info.filter_and_record(output)

Strip kernel-resource-usage remarks from output, parse them, and append the parsed entries to the active recorder (if any). Returns the filtered output with the remark lines removed.

Parameters:

output (str)

Return type:

str

tilelang.contrib.hip_resource_info.dump_to_file(usage, path)

Persist parsed resource usage so it survives kernel-cache hits.

Parameters:

• usage (dict[str, KernelResourceUsage])

• path (str)

Return type:

None

tilelang.contrib.hip_resource_info.load_from_file(path)

Inverse of dump_to_file. Tolerant of missing / unknown fields so older cache entries keep working when the dataclass evolves.

Parameters:

path (str)

Return type:

dict[str, KernelResourceUsage]