KernelInterface 1.0: a better API for writing kernels

Aperta
  • Qualcosa non va?
#810 3 commenti 0 reazioni 0 assegnatari Vedi su GitHub

I maintainer di solito rispondono entro 1 giorno

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
5/5
Tempo stimato
Più di una settimana
Idoneità per principianti
25/100
Tipo di issue
Funzionalità
Chiarezza
Abbastanza chiara
Stato di attività
Attiva
Stack tecnologico
julia
Ambito
backend-api-design

Direzione di ricerca

Start with the existing KI examples and examples/histogram.jl, which the issue identifies as relying on KA; compare their use of KI and KA against the proposed contracts. The issue lists several breaking API design choices and documentation items but does not specify a settled scope or tests, so confirm which proposal is in scope before starting. Done should mean an agreed change with behavior and documentation matching that decision.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

kernels

KernelInterface right now is mostly about the contract between KernelAbstractions and the back-ends. KI now validates and sizes launches itself, so back-ends implement KI.launch with checked 3-D sizes instead of the whole launch path. The semantics that differed between back-ends are written down (work-group limits, typed index queries, sub-groups, capability defaults), and a testsuite checks them. On top of that, #801 lets KA launch @kernel kernels on any KI back-end, which shrinks CUDA's KA extension from 127 to 37 lines.

KI is also meant for users who want to write kernels without KA's macros, with more control or less overhead. 0.3 doesn't serve them as well yet. Some of the fixes are breaking, but since we know of no direct KI users outside the back-ends yet, I propose releasing KI 0.4 in a couple of weeks. It should be out before KA 0.10 is registered, so that KA 0.10 only ever targets 0.4.

Proposed for 0.4

  1. ndrange has different semantics in KA and KI. KA masks work-items outside the ndrange by default. KI rounds it up to whole work-groups and leaves bounds checks to the kernel. That is the right choice at this level: the kernel can guard its memory accesses while all work-items still reach barriers. But without an explicit bounds check, a kernel ported from KA and launched with KI.@launch ndrange=... can access memory out of bounds when the range is padded. Rename the keyword, or at least make the difference hard to miss in the docs.

  2. No private-memory primitive in KI. KI has localmemory, but per-work-item scratch storage is only exposed through KA.Scratchpad(ctx, T, Val(dims)), which takes KA's hidden context. Back-ends override it and ignore the context (CUDA returns an MArray, PoCL a stack allocation). Add a portable KI primitive without the context, with defined lifetime and initialization, and lower @private to it. Back-ends then no longer need a KA hook for it.

  3. @Const has no KI equivalent. @Const promises that the memory an argument refers to is neither written by the kernel nor aliased by other arguments. Back-ends implement it with an Adapt rule for KA.ConstAdaptor. KI users could use the same, e.g. to get ld.global.nc loads on CUDA. This could become a KI primitive with the same promise, leaving the annotation and the argument traversal in KA. It is less clear-cut than private memory, so it could also wait.

  4. supports_subgroups implies a fixed width. It covers the sub-group queries and sub_group_barrier, and also promises that every compiled kernel uses the width reported by sub_group_size(backend). A back-end that can't guarantee the width has to report no sub-group support at all, which is what happens on OpenCL with NVIDIA. Kernels that read the width at run time don't need that guarantee. Split it into support for sub-group operations and a separate guarantee of a fixed width.

  5. Launch keywords. Keywords that KI doesn't know go to KI.launch on a Kernel call, but to kernel_function on KI.@launch. So back-end launch options such as CUDA's stream need launch=false and a second call. Adding a KI keyword later also changes what happens to a back-end keyword of the same name: #800 proposes dynamic local memory as a Kernel keyword, and calling it shmem would take over a keyword that is currently forwarded to CUDA. Either separate compiler and launch options explicitly (e.g. launch_options=(; stream)), or add a rule that new KI keywords only come in breaking releases.

Non-breaking (can go in any release)

  • An example that uses KI without KA. Both KI examples load KA, and examples/histogram.jl takes its atomics from KA.
  • Make clear that supports_atomics does not guarantee local-memory atomics, even though the histogram example uses them.
  • A guide for kernel authors that brings the existing contracts together: for each function, who calls it and who implements it, and which behavior is unspecified (shuffles from lanes that don't exist, how work-items map onto sub-groups). It should also point out that the index queries default to Int. Since #797, KA picks Int32 by itself when the indices fit, so a KI port of a KA kernel can end up with 64-bit index arithmetic where KA used 32-bit.
Lingua principale
Julia
Stelle
533
Fork
90
Merge medio
1g 11h
PR unite (30g)
66

Preparare l'ambiente

Questo progetto non fornisce container di sviluppo, Dockerfile né guida per i contributori, quindi l'ambiente è a tuo carico: parti dal suo README e consulta la nostra guida al primo contributo per i passaggi generali.

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di JuliaGPU/KernelAbstractions.jl

Tutte le issue di JuliaGPU/KernelAbstractions.jl

Issue simili

Altre issue su Julia

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.