Hardware-accelerated linear KA index is wrong #346

Activity

1. maleadt commented on Jul 29, 2025

   

   maleadt

   on Jul 29, 2025

   Last edited by maleadt

   Member

   More actions

   Oh fascinating. This reproduces with a much simpler example:

   using OpenCL
   import KernelAbstractions as KA
   
   @kernel function foo(max, actual)
       ctx = KA.@context
       I = KA.__index_Global_Cartesian(ctx)
       i = LinearIndices(KA.__ndrange(ctx))[I]
       j = KA.__index_Global_Linear(ctx)
   
       if j > 20
           @inbounds actual[1] = i
           OpenCL.@println "thread $(get_local_id(1)) block $(get_group_id(1)): I = ($(I[1]), $(I[2])) => i = $i, while j = $j"
       end
   end
   
   function main(n = 10, m = 2)
       actual = OpenCL.zeros(Int)
       foo(OpenCLBackend())(n*m, actual; ndrange = (n, m), workgroupsize = 4)
       Array(actual)[1]
   end

   julia> main()
   
   # some additional output
   groups = length(KA.blocks(iterspace)) = 6
   items = length(KA.workitems(iterspace)) = 4
   global_size = 24
   local_size = 4
   
   # kernel output
   thread 1 block 6: I = (9, 2) => i = 19, while j = 21
   thread 2 block 6: I = (10, 2) => i = 20, while j = 22
   

   So the issue is the global cartesian and linear indices disagreeing. For the (10, 2) ndrange, KA.jl ends up launching 24 threads (with a hard-coded workgroup size of 4). Thread 21 (local id 1, group id 6), which I think should be considered out of bounds, according to KA.jl is pointing to element element (9, 2) with linear index 19. This is because OpenCL does (6-1) * 4 + 1 = 21, while KA.jl does (8-1) * 2 + 1 based on the ndrange of (10, 2). I'm not sure how this is supposed to work, as it's mixing concepts that are incompatible. I'll let @vchuravy chime in when he's back from vacation.

   I guess this is also why JuliaGPU/CUDA.jl#2612 fails.

   👍React with 👍1Reacted by Simeon David Schaub
2. maleadt

   changed the title [-]Wrong `@index(Global)` calculation for multidimensional ndrange?[/-] [+]Hardware-accelerated linear KA index is wrong[/+] on Jul 29, 2025
3. maleadt

   added

   kernelsWriting and compiling kernels: device intrinsics, codegen, KernelAbstractions back-end.Writing and compiling kernels: device intrinsics, codegen, KernelAbstractions back-end.

   on Oct 19, 2026
4. maleadt

   added the

   Bug

   issue type on Oct 19, 2026
5. maleadt commented on Oct 19, 2026

   

   maleadt

   on Oct 19, 2026

   Member

   More actions

   Fixed by #348, which stopped using get_global_id for the linear index, and the index functions now come from KernelAbstractions 0.10 itself (#519). The MWE now gives indices covering exactly 1:1856400.
6. maleadt

   closed this as completedon Oct 19, 2026