Design question: where should SPIR-V/OpenCL-specific GPU lowering live in Codon?

[BI71317]

Hi! I have been experimenting with a simple SPIR-V (OpenCL target) port on top of Codon’s existing GPU pipeline, mainly to understand what would be required to support a non-NVPTX backend.

From what I can tell so far, one difficulty is that although both NVPTX and SPIR-V can consume LLVM IR, they do not expect identical IR conventions. In particular, address spaces / pointer types / kernel-facing conventions do not line up cleanly, and the current lowering path appears to be heavily shaped around NVPTX assumptions.

One obvious idea would be to adjust the high-level IR -> LLVM IR lowering so that it can emit SPIR-V-friendly LLVM IR instead of NVPTX-oriented LLVM IR.

However, I’m not sure this is the right layer for the change, because there seem to be multiple source -> LLVM IR paths, and some of the information needed for SPIR-V-friendly lowering may need to exist earlier in the pipeline.

For example, By using Raw LLVM IR in Source Level code, While LLVM IR requires specific address space (SPIR-V LLVM IR requires AS1 ptr while nvptx uses AS0 (default) ptr), Source Level code returns only default address space pointer.

# in gpu.codon
@pure
@llvm
def _ptr_to_type(p: cobj, T: type) -> T:
    ret ptr %p


@pure
@llvm
def _type_to_ptr(p: cobj, T: type) -> T:
    ret ptr %p

Like these cases, as target-specific pointer/address-space distinctions matter, it feels like those distinctions may need to be represented before the final LLVM IR lowering step, rather than patched in only at the end.

So my main question is about project direction:

- Would you expect SPIR-V/OpenCL support to be approached by adding a target-specific rewrite close to LLVM IR emission?

- Or would you prefer introducing some more explicit target-aware representation earlier in Codon IR / the source-level typing model?

- More generally, if support for another GPU backend were explored, where would you expect the main abstraction boundary to be?

I’m not claiming I have the right design yet — I mostly want to understand what direction would be considered idiomatic for Codon before going further with the experiment.

If helpful, I can also write up the concrete mismatches I’ve run into so far.

[arshajii]

Hi @BI71317 -- my first thought would be to try to keep it at the LLVM level if possible. I don't know too much about SPIR-V, but e.g. would it be possible to have an LLVM pass that takes Codon's output LLVM IR and converts it to SPIR-V compatible IR (which would include e.g. adding addrspace to pointers etc.)? This is effectively what we do to some extent already with the NVPTX backend; we replace certain functions, remove certain instructions, etc. SPIR-V has more requirements presumably, but as far as I can tell this approach might still be viable... If you can list some of those mismatches, I think it would help us to evaluate this approach.

[BI71317]

Hi @arshajii

I looked into this a bit more, and from what I can tell, one of the main issues seems to be that NVPTX and SPIR-V don’t accept “default” pointers under quite the same contract.

For NVPTX, LLVM/NVVM explicitly treats addrspace(0) as the generic/default pointer space. So lowering plain pointers that way seems to fit the backend model fairly naturally. As far as I can tell, this also matches Codon’s current implementation, where pointers are lowered into addrspace 0 in llvisitor.cpp.

For SPIR-V/OpenCL though, Generic seems to be a different kind of thing. Rather than being a concrete object storage class, OpenCL defines it more like a placeholder pointer space over concrete named spaces such as global, local, and private.

For example, passing a generic pointer argument directly to a kernel entry function failed for me, while passing a global pointer succeeded:

/home/swchoi/src/test_code/kernel_return_semantics/kernel_spirv_addrspace_generic.cl:1:46: error: pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space
    1 | __kernel void generic_probe(__generic float *scratch, const int n) {
      |                                              ^
1 error generated.

So my current understanding is that, at least in the SPIR-V/OpenCL path I tested, more specific address-space information seems to be required earlier in the pipeline, whereas NVPTX is more tolerant of keeping pointers generic for longer.

Also, my impression is that ptr vs ptr addrspace(N) is not just a matter of adding annotations. They are effectively different pointer types in IR, so once a pointer type changes, that change tends to propagate through loads/stores/GEPs, call sites, return types, and so on.

For example,

callee:

define void @use(ptr %p) {
entry:
  %x = load i32, ptr %p, align 4
  ret void
}

caller:

define void @caller(ptr %q) {
entry:
  call void @use(ptr %q)
  ret void
}

If we change the pointer type to ptr addrspace(1), it’s not just the callee signature that changes — the body changes too, and the caller/call site needs to be updated as well:

callee

define void @use(ptr addrspace(1) %p) {
entry:
  %x = load i32, ptr addrspace(1) %p, align 4
  ret void
}

caller

define void @caller(ptr addrspace(1) %q) {
entry:
  call void @use(ptr addrspace(1) %q)
  ret void
}

So my current impression is that LLVM-level rewriting is probably still viable at ABI boundaries, but some of the harder mismatches seem to come from storage-class information that may already be lost by the time we get to LLVM IR.

If it would be useful, I’d be happy to put together a more concrete list of the mismatches I ran into.