Avoid promition to Int32 in work_items functions

Author: vchuravy

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")
 

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