Q: How to debug kernels - KA.@print?

[svilupp]

First of all, thank you for the awesome package!

How do you print / introspect your kernels?

Eg,

• how do you introspect the grid/hierarchy in your object? (I have read the Apple docs but I'm struggling to imagine how that hierarchy looks in my matrix)
• see results of calls like dispatch_simdgroups_per_threadgroup() or sgitg=simdgroup_index_in_threadgroup()?
• how do you check that your inner loops are correct?
• any other tips and tricks that help you debug your kernel?

I couldn't find any examples in the repo and @print doesn't seem to be implemented?

Context:
I was hoping to replicate a kernel using SIMD groups (example in .metal, but I cannot figure out the indexing with simdgroup_index_in_threadgroup etc, because I'm not even sure how the grid looks.

[maleadt]

Hi,

I've converted this to a discussion, since it's not really a bug report. For general questions, Github Discussions (or Slack or Discourse) are generally better places.

The main problem here is that we don't have a device-side printing function right now. Other platforms, like CUDA and oneAPI, offer an I/O intrinsic, but Metal doesn't. So we need to implement our own I/O, which is fairly hard. AMDGPU.jl has it, so we could port that solution, but that will take time.

In the mean time, you should use an array to write debug values into. Very clunky, I know, but this isn't unique to Metal.jl (writing kernels in Metal C suffers from the same restriction). To simplify this, it's recommended to incrementally write your kernels. Begin with simple operations, and validate the output at every step.

[svilupp]

If anyone finds this, I found it super helpful to write my kernels by first accumulating/adding "1s" in each loop/thread/group when I'm figuring out the indices. Then I can immediately check if the "work" has been distributed correctly and I see the right number of "units of works" in each cell of the output array