From 8e929c9941bb2504d4b237935cd888edfc82ecc4 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Sun, 11 Oct 2026 13:50:28 +0200 Subject: [PATCH] Enable SPV_KHR_float_controls2 when the device accepts it Base SPIR-V can't carry the `contract` fast-math flag, so `muladd` reaches the driver as a separate multiply and add (and the SPIR-V back-end marks every kernel `ContractionOff`), which drivers never fuse into an FMA (JuliaGPU/KernelAbstractions.jl#865). SPV_KHR_float_controls2 carries the flag as `FPFastMathMode AllowContract`. Enabling the extension makes the back-end declare it in every kernel, which drivers that don't accept it reject, so add a `spirv_extensions` device property (cl_khr_spirv_queries) and enable the extension only when the device lists it. Explicitly passed `extensions` replace this default, and programs built through spirv2clc don't get it. Also add a missing comma in `propertynames(::Device)`, which made it throw. Assisted-by: Claude Code (Opus 5.5) --- lib/cl/device.jl | 15 +++++++++++++-- src/compiler/compilation.jl | 13 ++++++++++++- test/device.jl | 3 +++ test/execution.jl | 21 +++++++++++++++++++++ 4 files changed, 49 insertions(+), 3 deletions(-) diff --git a/lib/cl/device.jl b/lib/cl/device.jl index 17806f82..b693ce07 100644 --- a/lib/cl/device.jl +++ b/lib/cl/device.jl @@ -112,6 +112,17 @@ end return String[string(s) for s in split(bs)] end + # cl_khr_spirv_queries: the SPIR-V extensions the device accepts. Empty if the device + # doesn't support the query. + if s == :spirv_extensions + "cl_khr_spirv_queries" in d.extensions || return String[] + size = Ref{Csize_t}() + clGetDeviceInfo(d, CL_DEVICE_SPIRV_EXTENSIONS_KHR, 0, C_NULL, size) + result = Vector{Cstring}(undef, size[] รท sizeof(Cstring)) + clGetDeviceInfo(d, CL_DEVICE_SPIRV_EXTENSIONS_KHR, size[], result, C_NULL) + return String[unsafe_string(ext) for ext in result] + end + if s == :platform result = Ref{cl_platform_id}() clGetDeviceInfo(d, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), result, C_NULL) @@ -187,9 +198,9 @@ function Base.propertynames(::Device, private::Bool=false) :max_work_group_size, :max_parameter_size, :profiling_timer_resolution, # boolean properties :has_image_support, :has_local_mem, :host_unified_memory, :available, - :compiler_available, :non_uniform_work_group_support + :compiler_available, :non_uniform_work_group_support, # other - :extensions, :platform, :device_type, :max_work_item_size, :sub_group_sizes, + :extensions, :spirv_extensions, :platform, :device_type, :max_work_item_size, :sub_group_sizes, :max_image2d_shape, :max_image3d_shape, ) return private ? (fieldnames(Device)..., props...) : props diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index f40add91..608abf2b 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -184,6 +184,12 @@ function spirv_driver(dev::cl.Device) end end +# SPIR-V extensions that let the back-ends express more of the IR, enabled when the device +# accepts them: +# - SPV_KHR_float_controls2 carries per-instruction fast-math flags, such as `contract` (e.g. +# from `muladd`), which base SPIR-V can't express, so drivers never fuse those into FMAs. +const OPTIONAL_SPIRV_EXTENSIONS = ["SPV_KHR_float_controls2"] + @noinline function _compiler_config(dev, backend; kernel=true, name=nothing, always_inline=false, debug_level=Base.JLOptions().debug_level, sub_group_size::Union{Nothing,Int}=_sub_group_size(dev), @@ -210,8 +216,13 @@ end error("Device does not support cl_intel_required_subgroup_size") end + if extensions === nothing + # spirv2clc doesn't implement them, so only for programs built from SPIR-V + extensions = source ? String[] : filter(in(dev.spirv_extensions), OPTIONAL_SPIRV_EXTENSIONS) + end + # GPUCompiler adds the extensions that the atomics need - spirv_ext = extensions === nothing ? "" : join(("+$ext" for ext in extensions), ",") + spirv_ext = join(("+$ext" for ext in extensions), ",") # create GPUCompiler objects target = SPIRVCompilerTarget(; version=SPIRV_VERSION, supports_fp16, supports_fp64, diff --git a/test/device.jl b/test/device.jl index 31f5ca34..ade22527 100644 --- a/test/device.jl +++ b/test/device.jl @@ -31,6 +31,7 @@ end :driver_version, :version, :extensions, + :spirv_extensions, :platform, :name, :device_type, @@ -78,6 +79,8 @@ end if length(v) > 0 @test isa(v, Array{String, 1}) end + elseif k == :spirv_extensions + @test isa(v, Vector{String}) elseif k == :platform @test v == cl.platform() elseif k == :max_work_item_sizes diff --git a/test/execution.jl b/test/execution.jl index 96e3ab8e..714d0326 100644 --- a/test/execution.jl +++ b/test/execution.jl @@ -456,6 +456,27 @@ end end end +@testset "fast-math contraction" begin + # `muladd` is `fmul contract` + `fadd contract`, which only SPV_KHR_float_controls2 can + # carry into SPIR-V (JuliaGPU/KernelAbstractions.jl#865) + function contract_kernel(c, a, b) + i = get_global_id() + @inbounds c[i] = muladd(a[i], b[i], c[i]) + return + end + T = CLDeviceArray{Float32, 1, AS.CrossWorkgroup} + # (enabled when the device accepts it, except on the OpenCL C source backend) + enabled = occursin("SPV_KHR_float_controls2", + OpenCL.compiler_config(cl.device()).target.extensions) + spirv = sprint(io -> OpenCL.code_native(io, contract_kernel, Tuple{T, T, T}; kernel = true)) + @test occursin("FPFastMathMode AllowContract", spirv) == enabled + + # explicitly passing the extensions disables the default ones + spirv = sprint(io -> OpenCL.code_native(io, contract_kernel, Tuple{T, T, T}; + kernel = true, extensions = String[])) + @test !occursin("AllowContract", spirv) +end + @testset "has_feature folding" begin # has_feature must fold at compile time: the feature-bitset global gets baked in and DCE'd, # leaving only the branch matching the device.