[maleadt]

Each KernelAbstractions launch queried HIP for the device's maximum workgroup size (hipDeviceGetAttribute, which formats a string), twice for the kernel's maximum workgroup size, and for the kernel's occupancy. In a profiled launch loop, these queries took about 330 of the 390 ns that KernelAbstractions adds to a launch.

Store the device limit in HIPDevice (HIP.max_workgroup_size), alongside max_workgroup_dims, and cache the kernel limit and occupancy without dynamic shared memory in HIPFunction, as CUDA.jl does in CuFunction. Both are fixed when the function is loaded and need no invalidation: HIP's cache-configuration and attribute setters do not change the inputs to the occupancy calculation.

HIP picks the largest workgroup size occupancy allows, or the limit rounded up to the wavefront size if smaller (ihipOccupancyMaxActiveBlocksPerMultiprocessor in clr), so limits that admit the cached result reuse it. Smaller limits, e.g. an ndrange below the cached workgroup size, still query HIP. The test compares cached and queried configurations across a range of limits.

Host time per launch, 64 launches per batch without synchronization, median over 6 interleaved runs (9950X iGPU, gfx1036; Julia 1.12.7, KernelAbstractions 0.10.3, KernelInterface 0.4.3; base 3e70165):

launch	before	after
@roc, for reference	816 ns	827 ns
KA, ndrange=1000	1472 ns	1187 ns
KA, ndrange=10_000	1451 ns	1100 ns
KA, static workgroup size 256	1283 ns	1079 ns
KA, 3-d ndrange=(30, 30, 30)	1543 ns	1121 ns

The "after" runs also include #1197, which only changes compile-time latency.

Tests: the full suite on gfx1036 (30553 passed, 14 broken), including kernelabstractions_tests and kernelinterface_tests.

[luraess]

cscs-ci run MI250,MI300

[github-actions]

Documentation preview: https://amdgpu.juliagpu.org/previews/PR1196/

[github-actions]

Benchmark results

Compared with 3e7016582, the merge base.
The environment differs from the baseline's: AMDGPUExt 0.5.9 → 0.10.3, AcceleratedKernels 0.5.9 → 0.5.10, Atomix 1.6.0 → 1.7.0, BFloat16sExt 10.4.1 → 10.4.2, JSON 1.12.2 → 1.13.0, LLVM 10.4.1 → 10.4.2, Parsers 3.0.1 → 3.0.2, StructUtils 2.11.1 → 2.11.2, StructUtilsStaticArraysCoreExt 2.11.1 → 2.11.2, StructUtilsTablesExt 2.11.1 → 2.11.2, pocl_standalone_jll 7.2.3+1 → 7.2.3+2.

Times are the minimum time of every benchmark. A change is unusual when it exceeds both three times the benchmark's typical change between consecutive runs on the default branch and a minimum per category (performance 3%, latency 5%, applications 3%, other 3%). A benchmark that starts or stops allocating is listed however few the allocations. A single run can still be a fluke, so rerun before drawing conclusions.

	Geometric mean	Unusual: slower	Unusual: faster	Compared
Performance	−4.9%	4	19	69
Latency	−3.4%	0	0	3
Applications	−3.4%	0	1	14

Benchmark	Baseline	This PR	Change
Performance
array/random/rand!/Int64	59.8 µs	38.8 µs	−35.1%	🟢 faster
kernel/rand	48.3 µs	71.7 µs	+48.4%	🔴 slower
array/broadcast	50.6 µs	38.1 µs	−24.8%	🟢 faster
array/reductions/reduce/Float32/dims=2L	72.5 µs	57.3 µs	−21.0%	🟢 faster
array/reductions/mapreduce/Float32/dims=1	55.5 µs	44.7 µs	−19.4%	🟢 faster
array/iteration/findmin/2d	79.3 µs	64.6 µs	−18.5%	🟢 faster
array/random/rand/Int64	46.4 µs	38.0 µs	−18.1%	🟢 faster
array/accumulate/Int64/1d	104 µs	85.4 µs	−18.0%	🟢 faster
array/accumulate/Int64/dims=1	75.1 µs	62.6 µs	−16.7%	🟢 faster
array/sorting/1d	343 µs	292 µs	−14.9%	🟢 faster
array/reductions/mapreduce/Float32/dims=2L	60.7 µs	69.5 µs	+14.6%	🔴 slower
array/copyto!/gpu_to_gpu	35.1 µs	40.2 µs	+14.4%	🔴 slower
array/reductions/reduce/Float32/dims=1L	50.4 µs	57.0 µs	+13.0%	🔴 slower
array/iteration/logical	91.0 µs	83.0 µs	−8.8%	🟢 faster
array/reverse/1dL_inplace	78.8 µs	72.0 µs	−8.6%	🟢 faster
array/iteration/findfirst/bool	54.9 µs	50.3 µs	−8.3%	🟢 faster
array/reverse/2dL	80.5 µs	74.1 µs	−8.0%	🟢 faster
array/iteration/findall/bool	94.7 µs	87.4 µs	−7.7%	🟢 faster
array/accumulate/Float32/dims=2L	124 µs	116 µs	−6.5%	🟢 faster
array/reductions/reduce/Float32/dims=2	59.4 µs	55.8 µs	−6.0%	🟢 faster
array/reductions/reduce/Int64/dims=2	60.6 µs	57.0 µs	−5.9%	🟢 faster
array/reductions/mapreduce/Int64/dims=2	60.7 µs	57.2 µs	−5.8%	🟢 faster
array/reductions/mapreduce/Float32/dims=2	59.6 µs	56.2 µs	−5.6%	🟢 faster
Applications
gemm/tiled	2.75 ms	1.91 ms	−30.6%	🟢 faster

Benchmark	Baseline allocations	This PR
Applications
prefix_sum	none	2 (84 B)	🔴 allocates

All benchmarks

Benchmark	Baseline	This PR	Change
Performance
amdgpu/synchronization/context/device	1.79 µs	1.77 µs	−1.3%
amdgpu/synchronization/stream/blocking	228 ns	235 ns	+3.3%
amdgpu/synchronization/stream/nonblocking	350 ns	340 ns	−2.9%
array/accumulate/Float32/1d	82.6 µs	80.2 µs	−2.9%
array/accumulate/Float32/dims=1	64.2 µs	64.4 µs	+0.3%
array/accumulate/Float32/dims=1L	105 µs	106 µs	+0.2%
array/accumulate/Float32/dims=2	105 µs	103 µs	−1.5%
array/accumulate/Float32/dims=2L	124 µs	116 µs	−6.5%	🟢 faster
array/accumulate/Int64/1d	104 µs	85.4 µs	−18.0%	🟢 faster
array/accumulate/Int64/dims=1	75.1 µs	62.6 µs	−16.7%	🟢 faster
array/accumulate/Int64/dims=1L	111 µs	115 µs	+3.6%
array/accumulate/Int64/dims=2	91.0 µs	88.7 µs	−2.6%
array/accumulate/Int64/dims=2L	157 µs	162 µs	+2.8%
array/broadcast	50.6 µs	38.1 µs	−24.8%	🟢 faster
array/construct	2.29 µs	2.27 µs	−1.2%
array/copy	40.3 µs	38.8 µs	−3.7%
array/copyto!/cpu_to_gpu	85.9 µs	86.0 µs	+0.1%
array/copyto!/gpu_to_cpu	85.8 µs	86.0 µs	+0.2%
array/copyto!/gpu_to_gpu	35.1 µs	40.2 µs	+14.4%	🔴 slower
array/iteration/findall/bool	94.7 µs	87.4 µs	−7.7%	🟢 faster
array/iteration/findall/int	107 µs	102 µs	−4.8%
array/iteration/findfirst/bool	54.9 µs	50.3 µs	−8.3%	🟢 faster
array/iteration/findfirst/int	56.6 µs	56.7 µs	+0.3%
array/iteration/findmin/1d	95.0 µs	94.5 µs	−0.5%
array/iteration/findmin/2d	79.3 µs	64.6 µs	−18.5%	🟢 faster
array/iteration/logical	91.0 µs	83.0 µs	−8.8%	🟢 faster
array/iteration/scalar	307 µs	297 µs	−3.4%
array/permutedims/2d	62.1 µs	60.1 µs	−3.2%
array/permutedims/3d	49.4 µs	57.5 µs	+16.4%
array/permutedims/4d	54.1 µs	51.0 µs	−5.8%
array/random/rand!/Float32	50.9 µs	42.7 µs	−16.2%
array/random/rand!/Int64	59.8 µs	38.8 µs	−35.1%	🟢 faster
array/random/rand/Float32	46.9 µs	46.0 µs	−1.9%
array/random/rand/Int64	46.4 µs	38.0 µs	−18.1%	🟢 faster
array/random/randn!/Float32	82.6 µs	69.3 µs	−16.1%
array/random/randn/Float32	70.5 µs	70.6 µs	+0.2%
array/reductions/mapreduce/Float32/1d	71.0 µs	69.1 µs	−2.6%
array/reductions/mapreduce/Float32/dims=1	55.5 µs	44.7 µs	−19.4%	🟢 faster
array/reductions/mapreduce/Float32/dims=1L	59.7 µs	57.3 µs	−4.2%
array/reductions/mapreduce/Float32/dims=2	59.6 µs	56.2 µs	−5.6%	🟢 faster
array/reductions/mapreduce/Float32/dims=2L	60.7 µs	69.5 µs	+14.6%	🔴 slower
array/reductions/mapreduce/Int64/1d	71.7 µs	69.9 µs	−2.5%
array/reductions/mapreduce/Int64/dims=1	51.6 µs	49.3 µs	−4.5%
array/reductions/mapreduce/Int64/dims=1L	70.9 µs	68.9 µs	−2.8%
array/reductions/mapreduce/Int64/dims=2	60.7 µs	57.2 µs	−5.8%	🟢 faster
array/reductions/mapreduce/Int64/dims=2L	78.4 µs	75.0 µs	−4.3%
array/reductions/reduce/Float32/1d	70.9 µs	69.1 µs	−2.5%
array/reductions/reduce/Float32/dims=1	55.7 µs	53.6 µs	−3.8%
array/reductions/reduce/Float32/dims=1L	50.4 µs	57.0 µs	+13.0%	🔴 slower
array/reductions/reduce/Float32/dims=2	59.4 µs	55.8 µs	−6.0%	🟢 faster
array/reductions/reduce/Float32/dims=2L	72.5 µs	57.3 µs	−21.0%	🟢 faster
array/reductions/reduce/Int64/1d	71.4 µs	69.4 µs	−2.9%
array/reductions/reduce/Int64/dims=1	51.6 µs	49.4 µs	−4.3%
array/reductions/reduce/Int64/dims=1L	70.9 µs	68.6 µs	−3.3%
array/reductions/reduce/Int64/dims=2	60.6 µs	57.0 µs	−5.9%	🟢 faster
array/reductions/reduce/Int64/dims=2L	78.2 µs	74.6 µs	−4.6%
array/reverse/1d	37.3 µs	41.6 µs	+11.4%
array/reverse/1d_inplace	40.3 µs	41.0 µs	+1.8%
array/reverse/1dL	60.5 µs	60.3 µs	−0.3%
array/reverse/1dL_inplace	78.8 µs	72.0 µs	−8.6%	🟢 faster
array/reverse/2d	47.8 µs	46.5 µs	−2.7%
array/reverse/2d_inplace	56.1 µs	43.4 µs	−22.6%
array/reverse/2dL	80.5 µs	74.1 µs	−8.0%	🟢 faster
array/reverse/2dL_inplace	90.5 µs	67.8 µs	−25.1%
array/sorting/1d	343 µs	292 µs	−14.9%	🟢 faster
kernel/indexing	32.1 µs	31.5 µs	−1.7%
kernel/indexing_checked	37.0 µs	36.2 µs	−2.3%
kernel/launch	900 ns	895 ns	−0.6%
kernel/rand	48.3 µs	71.7 µs	+48.4%	🔴 slower
Latency
import	1.37 s	1.31 s	−4.6%
precompile	29.0 s	27.8 s	−4.3%
ttfp	1.77 s	1.75 s	−1.1%
Applications
bitonic_sort	1.11 ms	1.12 ms	+0.4%
byval/reference	39.2 µs	40.0 µs	+1.9%
byval/slices=1	41.1 µs	40.6 µs	−1.4%
byval/slices=2	158 µs	136 µs	−14.2%
byval/slices=3	233 µs	234 µs	+0.6%
convolution	76.0 µs	75.4 µs	−0.8%
floyd_warshall	9.06 ms	9.08 ms	+0.2%
gemm/tiled	2.75 ms	1.91 ms	−30.6%	🟢 faster
gemm/tiled_unbounded	2.79 ms	2.77 ms	−0.4%
histogram	782 µs	800 µs	+2.4%
prefix_sum	194 µs	195 µs	+0.5%
stencil/diffusion3d	1.57 ms	1.57 ms	+0.2%
stencil/diffusion3d_checked	1.61 ms	1.62 ms	+0.5%
volumerhs	4.86 ms	4.84 ms	−0.4%

History: https://juliagpu.org/perf/?pkg=amdgpu