From 2cac7291a19dd6d09e96221140612ee2e33d2323 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Wed, 5 Mar 2025 18:45:08 +0100 Subject: [PATCH 1/2] Avoid promition to Int32 in work_items functions --- lib/intrinsics/src/SPIRVIntrinsics.jl | 7 +++++++ lib/intrinsics/src/work_item.jl | 16 ++++++++-------- 2 files changed, 15 insertions(+), 8 deletions(-) diff --git a/lib/intrinsics/src/SPIRVIntrinsics.jl b/lib/intrinsics/src/SPIRVIntrinsics.jl index 5de144eb..2c64921f 100644 --- a/lib/intrinsics/src/SPIRVIntrinsics.jl +++ b/lib/intrinsics/src/SPIRVIntrinsics.jl @@ -7,6 +7,13 @@ import ExprTools import SpecialFunctions +# helper type for writing UInt32/Int32 literals +# TODO: upstream this +struct Literal{T} end +Base.:(*)(x::Number, ::Type{Literal{T}}) where {T} = T(x) +const i32 = Literal{Int32} +const u32 = Literal{UInt32} + include("pointer.jl") include("utils.jl") diff --git a/lib/intrinsics/src/work_item.jl b/lib/intrinsics/src/work_item.jl index ee3cf73e..7d279a1a 100644 --- a/lib/intrinsics/src/work_item.jl +++ b/lib/intrinsics/src/work_item.jl @@ -12,17 +12,17 @@ export get_work_dim, @device_function get_work_dim() = @builtin_ccall("get_work_dim", UInt32, ()) % Int -@device_function get_global_size(dimindx::Integer=1) = @builtin_ccall("get_global_size", UInt, (UInt32,), dimindx-1) % Int -@device_function get_global_id(dimindx::Integer=1) = @builtin_ccall("get_global_id", UInt, (UInt32,), dimindx-1) % Int + 1 +@device_function get_global_size(dimindx::Integer = 1u32) = @builtin_ccall("get_global_size", UInt, (UInt32,), dimindx - 1u32) % Int +@device_function get_global_id(dimindx::Integer = 1u32) = @builtin_ccall("get_global_id", UInt, (UInt32,), dimindx - 1u32) % Int + 1 -@device_function get_local_size(dimindx::Integer=1) = @builtin_ccall("get_local_size", UInt, (UInt32,), dimindx-1) % Int -@device_function get_enqueued_local_size(dimindx::Integer=1) = @builtin_ccall("get_enqueued_local_size", UInt, (UInt32,), dimindx-1) % Int -@device_function get_local_id(dimindx::Integer=1) = @builtin_ccall("get_local_id", UInt, (UInt32,), dimindx-1) % Int + 1 +@device_function get_local_size(dimindx::Integer = 1u32) = @builtin_ccall("get_local_size", UInt, (UInt32,), dimindx - 1u32) % Int +@device_function get_enqueued_local_size(dimindx::Integer = 1u32) = @builtin_ccall("get_enqueued_local_size", UInt, (UInt32,), dimindx - 1) % Int +@device_function get_local_id(dimindx::Integer = 1u32) = @builtin_ccall("get_local_id", UInt, (UInt32,), dimindx - 1u32) % Int + 1 -@device_function get_num_groups(dimindx::Integer=1) = @builtin_ccall("get_num_groups", UInt, (UInt32,), dimindx-1) % Int -@device_function get_group_id(dimindx::Integer=1) = @builtin_ccall("get_group_id", UInt, (UInt32,), dimindx-1) % Int + 1 +@device_function get_num_groups(dimindx::Integer = 1u32) = @builtin_ccall("get_num_groups", UInt, (UInt32,), dimindx - 1u32) % Int +@device_function get_group_id(dimindx::Integer = 1u32) = @builtin_ccall("get_group_id", UInt, (UInt32,), dimindx - 1u32) % Int + 1 -@device_function get_global_offset(dimindx::Integer=1) = @builtin_ccall("get_global_offset", UInt, (UInt32,), dimindx-1) % Int + 1 +@device_function get_global_offset(dimindx::Integer = 1u32) = @builtin_ccall("get_global_offset", UInt, (UInt32,), dimindx - 1u32) % Int + 1 @device_function get_global_linear_id() = @builtin_ccall("get_global_linear_id", UInt, ()) % Int + 1 @device_function get_local_linear_id() = @builtin_ccall("get_local_linear_id", UInt, ()) % Int + 1 From 50950029cb7ef9084018ca3c0b74e1b1604f97d5 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Thu, 17 Apr 2025 15:43:16 +0200 Subject: [PATCH 2/2] fixup! Avoid promition to Int32 in work_items functions --- lib/intrinsics/Project.toml | 2 ++ lib/intrinsics/src/SPIRVIntrinsics.jl | 7 +------ 2 files changed, 3 insertions(+), 6 deletions(-) diff --git a/lib/intrinsics/Project.toml b/lib/intrinsics/Project.toml index a4c6c753..ca2c3294 100644 --- a/lib/intrinsics/Project.toml +++ b/lib/intrinsics/Project.toml @@ -5,11 +5,13 @@ version = "0.2.0" [deps] ExprTools = "e2ba6199-217a-4e67-a87a-7c52f15ade04" +GPUToolbox = "096a3bc2-3ced-46d0-87f4-dd12716f4bfc" LLVM = "929cbde3-209d-540e-8aea-75f648917ca0" SpecialFunctions = "276daf66-3868-5448-9aa4-cd146d93841b" [compat] ExprTools = "0.1" +GPUToolbox = "0.2.0" LLVM = "9.1" SpecialFunctions = "1.3, 2" julia = "1.10" diff --git a/lib/intrinsics/src/SPIRVIntrinsics.jl b/lib/intrinsics/src/SPIRVIntrinsics.jl index 2c64921f..e3b6e464 100644 --- a/lib/intrinsics/src/SPIRVIntrinsics.jl +++ b/lib/intrinsics/src/SPIRVIntrinsics.jl @@ -7,12 +7,7 @@ import ExprTools import SpecialFunctions -# helper type for writing UInt32/Int32 literals -# TODO: upstream this -struct Literal{T} end -Base.:(*)(x::Number, ::Type{Literal{T}}) where {T} = T(x) -const i32 = Literal{Int32} -const u32 = Literal{UInt32} +using GPUToolbox include("pointer.jl") include("utils.jl")