Skip to content

PoCL: compute the typed index queries with % T - #807

Merged
maleadt merged 1 commit into
mainfrom
tb/pocl-spirvintrinsics-1.1.3
Sep 30, 2026
Merged

PoCL: compute the typed index queries with % T#807
maleadt merged 1 commit into
mainfrom
tb/pocl-spirvintrinsics-1.1.3

Conversation

@maleadt

@maleadt maleadt commented Sep 29, 2026

Copy link
Copy Markdown
Member

Split off from #800, as requested in review.

PoCL's typed index queries used T(x), a checked conversion, which leaves a throw_inexacterror branch in every kernel that asks for e.g. KI.get_global_id(Int32). This computes them with x % T instead.

That needs SPIRVIntrinsics 1.1.3 (JuliaGPU/OpenCL.jl#511): with older versions, InstCombine folds trunc(extractelement(load <3 x i64> @__spirv_BuiltIn…)) into a load of a vector type SPIR-V can't represent (e.g. <24 x i8>), and spirv-val rejects the kernel. So the compat bound goes up to 1.1.3.

The PoCL KernelInterface testsuite and the codegen checks pass locally.

Stacked on #803 (only for a linear stack; it doesn't depend on it). #800 is stacked on this.

@maleadt maleadt mentioned this pull request Sep 29, 2026
Merged
@maleadt
maleadt added this pull request to stack #808 September 29, 2026 20:22
@github-actions

github-actions Bot commented Sep 29, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Benchmark Results

Show table
main ce7bd83... main / ce7bd83...
const/@Const/Float32/262144 0.34 ± 0.031 ms 0.341 ± 0.031 ms 0.998 ± 0.13
const/@Const/Float32/65536 0.131 ± 0.029 ms 0.14 ± 0.028 ms 0.936 ± 0.28
const/@Const/Float64/262144 0.433 ± 0.031 ms 0.43 ± 0.03 ms 1.01 ± 0.1
const/@Const/Float64/65536 0.165 ± 0.014 ms 0.166 ± 0.014 ms 0.991 ± 0.12
const/unmarked/Float32/262144 1.15 ± 0.041 ms 1.16 ± 0.038 ms 0.997 ± 0.048
const/unmarked/Float32/65536 0.326 ± 0.032 ms 0.325 ± 0.034 ms 1 ± 0.14
const/unmarked/Float64/262144 1.46 ± 0.076 ms 1.36 ± 0.037 ms 1.07 ± 0.063
const/unmarked/Float64/65536 0.371 ± 0.037 ms 0.37 ± 0.035 ms 1 ± 0.14
launch/3D static workgroup, dynamic ndrange 0.0751 ± 0.013 ms 0.0748 ± 0.015 ms 1 ± 0.27
launch/3D static workgroup, static ndrange 0.0747 ± 0.013 ms 0.0747 ± 0.012 ms 1 ± 0.24
launch/dynamic workgroup, dynamic ndrange 0.0758 ± 0.022 ms 0.076 ± 0.017 ms 0.998 ± 0.37
launch/dynamic workgroup, dynamic ndrange, workgroupsize given 0.0706 ± 0.033 ms 0.069 ± 0.033 ms 1.02 ± 0.69
launch/static workgroup, dynamic ndrange 0.0738 ± 0.025 ms 0.0743 ± 0.013 ms 0.993 ± 0.37
launch/static workgroup, static ndrange 0.0745 ± 0.014 ms 0.0739 ± 0.018 ms 1.01 ± 0.31
partition/dynamic workgroup, dynamic ndrange 0.0583 ± 0.0012 μs 0.0581 ± 0.001 μs 1 ± 0.028
partition/static workgroup, dynamic ndrange 0.0648 ± 0.012 μs 0.0655 ± 0.012 μs 0.991 ± 0.26
partition/static workgroup, static ndrange 1.55 ± 0.01 ns 1.55 ± 0.01 ns 1 ± 0.0091
saxpy/default/Float16/1024 0.0769 ± 0.015 ms 0.0761 ± 0.021 ms 1.01 ± 0.34
saxpy/default/Float16/1048576 0.797 ± 0.041 ms 0.801 ± 0.041 ms 0.996 ± 0.072
saxpy/default/Float16/16384 0.0658 ± 0.034 ms 0.0661 ± 0.035 ms 0.995 ± 0.73
saxpy/default/Float16/2048 0.0765 ± 0.025 ms 0.0774 ± 0.018 ms 0.988 ± 0.4
saxpy/default/Float16/256 0.0776 ± 0.015 ms 0.0771 ± 0.015 ms 1.01 ± 0.27
saxpy/default/Float16/262144 0.251 ± 0.036 ms 0.249 ± 0.036 ms 1 ± 0.2
saxpy/default/Float16/32768 0.0741 ± 0.032 ms 0.0734 ± 0.032 ms 1.01 ± 0.63
saxpy/default/Float16/4096 0.0783 ± 0.031 ms 0.0794 ± 0.03 ms 0.987 ± 0.54
saxpy/default/Float16/512 0.0776 ± 0.011 ms 0.0773 ± 0.014 ms 1 ± 0.23
saxpy/default/Float16/64 0.0779 ± 0.013 ms 0.0768 ± 0.022 ms 1.01 ± 0.33
saxpy/default/Float16/65536 0.0954 ± 0.032 ms 0.0944 ± 0.03 ms 1.01 ± 0.46
saxpy/default/Float32/1024 0.0765 ± 0.012 ms 0.076 ± 0.01 ms 1.01 ± 0.21
saxpy/default/Float32/1048576 0.509 ± 0.041 ms 0.492 ± 0.034 ms 1.03 ± 0.11
saxpy/default/Float32/16384 0.0572 ± 0.031 ms 0.0557 ± 0.03 ms 1.03 ± 0.79
saxpy/default/Float32/2048 0.076 ± 0.029 ms 0.0741 ± 0.029 ms 1.03 ± 0.56
saxpy/default/Float32/256 0.078 ± 0.0097 ms 0.0775 ± 0.012 ms 1.01 ± 0.2
saxpy/default/Float32/262144 0.178 ± 0.032 ms 0.176 ± 0.032 ms 1.01 ± 0.26
saxpy/default/Float32/32768 0.0619 ± 0.028 ms 0.0606 ± 0.028 ms 1.02 ± 0.66
saxpy/default/Float32/4096 0.0728 ± 0.028 ms 0.0735 ± 0.027 ms 0.991 ± 0.52
saxpy/default/Float32/512 0.0768 ± 0.0099 ms 0.0761 ± 0.01 ms 1.01 ± 0.19
saxpy/default/Float32/64 0.0781 ± 0.013 ms 0.0768 ± 0.019 ms 1.02 ± 0.3
saxpy/default/Float32/65536 0.0831 ± 0.033 ms 0.0779 ± 0.031 ms 1.07 ± 0.6
saxpy/default/Float64/1024 0.0763 ± 0.0099 ms 0.0752 ± 0.015 ms 1.02 ± 0.25
saxpy/default/Float64/1048576 0.618 ± 0.058 ms 0.597 ± 0.063 ms 1.03 ± 0.15
saxpy/default/Float64/16384 0.0558 ± 0.029 ms 0.056 ± 0.031 ms 0.997 ± 0.76
saxpy/default/Float64/2048 0.0716 ± 0.031 ms 0.0725 ± 0.029 ms 0.987 ± 0.58
saxpy/default/Float64/256 0.0769 ± 0.0084 ms 0.0764 ± 0.015 ms 1.01 ± 0.22
saxpy/default/Float64/262144 0.199 ± 0.036 ms 0.2 ± 0.037 ms 0.997 ± 0.26
saxpy/default/Float64/32768 0.068 ± 0.031 ms 0.066 ± 0.031 ms 1.03 ± 0.68
saxpy/default/Float64/4096 0.0682 ± 0.027 ms 0.0696 ± 0.026 ms 0.98 ± 0.53
saxpy/default/Float64/512 0.076 ± 0.013 ms 0.0757 ± 0.013 ms 1 ± 0.24
saxpy/default/Float64/64 0.078 ± 0.013 ms 0.0773 ± 0.011 ms 1.01 ± 0.23
saxpy/default/Float64/65536 0.0884 ± 0.032 ms 0.0891 ± 0.032 ms 0.992 ± 0.51
saxpy/static workgroup=(1024,)/Float16/1024 0.0746 ± 0.013 ms 0.0743 ± 0.019 ms 1 ± 0.31
saxpy/static workgroup=(1024,)/Float16/1048576 0.795 ± 0.045 ms 0.798 ± 0.041 ms 0.997 ± 0.076
saxpy/static workgroup=(1024,)/Float16/16384 0.0609 ± 0.032 ms 0.0614 ± 0.032 ms 0.993 ± 0.73
saxpy/static workgroup=(1024,)/Float16/2048 0.0744 ± 0.027 ms 0.0754 ± 0.022 ms 0.987 ± 0.46
saxpy/static workgroup=(1024,)/Float16/256 0.075 ± 0.017 ms 0.0752 ± 0.012 ms 0.997 ± 0.28
saxpy/static workgroup=(1024,)/Float16/262144 0.244 ± 0.036 ms 0.248 ± 0.035 ms 0.985 ± 0.2
saxpy/static workgroup=(1024,)/Float16/32768 0.0706 ± 0.03 ms 0.0721 ± 0.03 ms 0.98 ± 0.59
saxpy/static workgroup=(1024,)/Float16/4096 0.0728 ± 0.033 ms 0.0734 ± 0.033 ms 0.992 ± 0.63
saxpy/static workgroup=(1024,)/Float16/512 0.0752 ± 0.019 ms 0.0749 ± 0.019 ms 1 ± 0.37
saxpy/static workgroup=(1024,)/Float16/64 0.0754 ± 0.008 ms 0.0752 ± 0.011 ms 1 ± 0.18
saxpy/static workgroup=(1024,)/Float16/65536 0.0917 ± 0.029 ms 0.0929 ± 0.029 ms 0.987 ± 0.44
saxpy/static workgroup=(1024,)/Float32/1024 0.0741 ± 0.011 ms 0.0735 ± 0.014 ms 1.01 ± 0.24
saxpy/static workgroup=(1024,)/Float32/1048576 0.495 ± 0.038 ms 0.498 ± 0.039 ms 0.994 ± 0.11
saxpy/static workgroup=(1024,)/Float32/16384 0.0564 ± 0.03 ms 0.058 ± 0.03 ms 0.972 ± 0.73
saxpy/static workgroup=(1024,)/Float32/2048 0.0702 ± 0.03 ms 0.0701 ± 0.03 ms 1 ± 0.61
saxpy/static workgroup=(1024,)/Float32/256 0.0752 ± 0.015 ms 0.075 ± 0.014 ms 1 ± 0.27
saxpy/static workgroup=(1024,)/Float32/262144 0.172 ± 0.031 ms 0.173 ± 0.03 ms 0.996 ± 0.25
saxpy/static workgroup=(1024,)/Float32/32768 0.0616 ± 0.029 ms 0.0613 ± 0.029 ms 1 ± 0.66
saxpy/static workgroup=(1024,)/Float32/4096 0.0658 ± 0.029 ms 0.0653 ± 0.029 ms 1.01 ± 0.63
saxpy/static workgroup=(1024,)/Float32/512 0.0741 ± 0.015 ms 0.0737 ± 0.018 ms 1.01 ± 0.32
saxpy/static workgroup=(1024,)/Float32/64 0.075 ± 0.014 ms 0.0748 ± 0.02 ms 1 ± 0.33
saxpy/static workgroup=(1024,)/Float32/65536 0.0774 ± 0.029 ms 0.0777 ± 0.029 ms 0.996 ± 0.53
saxpy/static workgroup=(1024,)/Float64/1024 0.0737 ± 0.012 ms 0.0734 ± 0.0093 ms 1 ± 0.21
saxpy/static workgroup=(1024,)/Float64/1048576 0.639 ± 0.083 ms 0.581 ± 0.044 ms 1.1 ± 0.17
saxpy/static workgroup=(1024,)/Float64/16384 0.058 ± 0.029 ms 0.0573 ± 0.029 ms 1.01 ± 0.72
saxpy/static workgroup=(1024,)/Float64/2048 0.0657 ± 0.031 ms 0.0683 ± 0.03 ms 0.962 ± 0.61
saxpy/static workgroup=(1024,)/Float64/256 0.0741 ± 0.022 ms 0.0737 ± 0.019 ms 1.01 ± 0.39
saxpy/static workgroup=(1024,)/Float64/262144 0.203 ± 0.032 ms 0.203 ± 0.034 ms 0.997 ± 0.23
saxpy/static workgroup=(1024,)/Float64/32768 0.0689 ± 0.029 ms 0.0676 ± 0.03 ms 1.02 ± 0.62
saxpy/static workgroup=(1024,)/Float64/4096 0.0616 ± 0.028 ms 0.0625 ± 0.027 ms 0.985 ± 0.62
saxpy/static workgroup=(1024,)/Float64/512 0.074 ± 0.012 ms 0.0734 ± 0.023 ms 1.01 ± 0.35
saxpy/static workgroup=(1024,)/Float64/64 0.0756 ± 0.013 ms 0.0746 ± 0.018 ms 1.01 ± 0.3
saxpy/static workgroup=(1024,)/Float64/65536 0.0903 ± 0.032 ms 0.0917 ± 0.032 ms 0.984 ± 0.49
time_to_load 0.758 ± 0.014 s 0.771 ± 0.012 s 0.983 ± 0.024
main ce7bd83... main / ce7bd83...
const/@Const/Float32/262144 1 allocs: 16 B 1 allocs: 16 B 1
const/@Const/Float32/65536 1 allocs: 16 B 1 allocs: 16 B 1
const/@Const/Float64/262144 1 allocs: 16 B 1 allocs: 16 B 1
const/@Const/Float64/65536 1 allocs: 16 B 1 allocs: 16 B 1
const/unmarked/Float32/262144 1 allocs: 16 B 1 allocs: 16 B 1
const/unmarked/Float32/65536 1 allocs: 16 B 1 allocs: 16 B 1
const/unmarked/Float64/262144 1 allocs: 16 B 1 allocs: 16 B 1
const/unmarked/Float64/65536 1 allocs: 16 B 1 allocs: 16 B 1
launch/3D static workgroup, dynamic ndrange 5 allocs: 0.141 kB 5 allocs: 0.141 kB 1
launch/3D static workgroup, static ndrange 6 allocs: 0.172 kB 6 allocs: 0.172 kB 1
launch/dynamic workgroup, dynamic ndrange 15 allocs: 0.312 kB 15 allocs: 0.312 kB 1
launch/dynamic workgroup, dynamic ndrange, workgroupsize given 1 allocs: 16 B 1 allocs: 16 B 1
launch/static workgroup, dynamic ndrange 1 allocs: 16 B 1 allocs: 16 B 1
launch/static workgroup, static ndrange 2 allocs: 0.0469 kB 2 allocs: 0.0469 kB 1
partition/dynamic workgroup, dynamic ndrange 2 allocs: 0.0625 kB 2 allocs: 0.0625 kB 1
partition/static workgroup, dynamic ndrange 2 allocs: 32 B 2 allocs: 32 B 1
partition/static workgroup, static ndrange 0 allocs: 0 B 0 allocs: 0 B
saxpy/default/Float16/1024 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/1048576 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/16384 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/2048 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/256 16 allocs: 0.328 kB 16 allocs: 0.328 kB 1
saxpy/default/Float16/262144 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/32768 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/4096 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/512 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float16/64 16 allocs: 0.328 kB 16 allocs: 0.328 kB 1
saxpy/default/Float16/65536 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/1024 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/1048576 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/16384 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/2048 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/256 16 allocs: 0.328 kB 16 allocs: 0.328 kB 1
saxpy/default/Float32/262144 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/32768 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/4096 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/512 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float32/64 16 allocs: 0.328 kB 16 allocs: 0.328 kB 1
saxpy/default/Float32/65536 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/1024 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/1048576 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/16384 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/2048 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/256 16 allocs: 0.328 kB 16 allocs: 0.328 kB 1
saxpy/default/Float64/262144 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/32768 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/4096 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/512 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/default/Float64/64 16 allocs: 0.328 kB 16 allocs: 0.328 kB 1
saxpy/default/Float64/65536 24 allocs: 0.453 kB 24 allocs: 0.453 kB 1
saxpy/static workgroup=(1024,)/Float16/1024 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/1048576 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/16384 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/2048 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/256 1 allocs: 16 B 1 allocs: 16 B 1
saxpy/static workgroup=(1024,)/Float16/262144 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/32768 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/4096 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/512 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float16/64 1 allocs: 16 B 1 allocs: 16 B 1
saxpy/static workgroup=(1024,)/Float16/65536 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/1024 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/1048576 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/16384 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/2048 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/256 1 allocs: 16 B 1 allocs: 16 B 1
saxpy/static workgroup=(1024,)/Float32/262144 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/32768 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/4096 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/512 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float32/64 1 allocs: 16 B 1 allocs: 16 B 1
saxpy/static workgroup=(1024,)/Float32/65536 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/1024 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/1048576 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/16384 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/2048 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/256 1 allocs: 16 B 1 allocs: 16 B 1
saxpy/static workgroup=(1024,)/Float64/262144 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/32768 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/4096 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/512 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
saxpy/static workgroup=(1024,)/Float64/64 1 allocs: 16 B 1 allocs: 16 B 1
saxpy/static workgroup=(1024,)/Float64/65536 4 allocs: 0.0625 kB 4 allocs: 0.0625 kB 1
time_to_load 0.2 k allocs: 11.8 kB 0.2 k allocs: 11.8 kB 1

Benchmark Plots

A plot of the benchmark results have been uploaded as an artifact to the workflow run for this PR.
Go to "Actions"->"Benchmark a pull request"->[the most recent run]->"Artifacts" (at the bottom).

Base automatically changed from tb/ci-reverse to main September 30, 2026 04:57
@maleadt
PoCL: compute the typed index queries with % T
ce7bd83 
`T(x)` is a checked conversion, which leaves a `throw_inexacterror` branch in
every kernel using a typed index query. `% T` doesn't, but needs
SPIRVIntrinsics 1.1.3 (JuliaGPU/OpenCL.jl#511): with older versions,
InstCombine folds the truncated 3-D builtin load into a vector type SPIR-V
can't represent.
@maleadt
maleadt force-pushed the tb/pocl-spirvintrinsics-1.1.3 branch from f74eae2 to ce7bd83 Compare September 30, 2026 04:57
@maleadt
maleadt merged commit 557fc62 into main Sep 30, 2026
54 of 55 checks passed
@maleadt
maleadt deleted the tb/pocl-spirvintrinsics-1.1.3 branch September 30, 2026 04:57
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@maleadt