def Sparkle.Backend.CudaSim.emitCudaType (ty : IR.Type.HWType) : String

Emit a CUDA/device-compatible scalar type. Delegates to CSim's emitTypeName, so scalars ≤ 64 bit are the native C integer types and wide (> 64 bit) values are uint32_t[⌈w/32⌉] arrays — identical to the CSim struct layout, which is what the device code operates on.

(#33 originally mapped wide types onto CUDA uint3/uint4, but those cap at 128 bit and don't match the CSim struct's word-array layout; the RV32 core's 578-bit bundle is exactly the case that broke. Reusing CSim's layout removes that limit.)

Equations

• Sparkle.Backend.CudaSim.emitCudaType ty = Sparkle.Backend.CSim.emitTypeName ty

def Sparkle.Backend.CudaSim.deviceStructName (m : IR.AST.Module) : String

The device state struct name. Under CSim the device state IS the plain-C struct <cls> that emitModule emits (inputs + outputs + observable wires + registers), so there is no separate _state_t type any more — the batch kernel and host API point straight at CSim's struct.

(#33 emitted its own <cls>_state_t and a copy-in/.evalTick()/copy-out wrapper because the old CppSim backend was a C++ class with a method. CSim replaced that class with a struct + free sparkle_<cls>_eval_tick function, which is already device-shaped — so the wrapper and the extra struct are gone.)

Equations

• Sparkle.Backend.CudaSim.deviceStructName m = toString "struct " ++ toString (Sparkle.Backend.CSim.sanitizeName m.name)

def Sparkle.Backend.CudaSim.deviceEvalTick (m : IR.AST.Module) : String

The device eval-tick symbol: CSim's own sparkle_<cls>_eval_tick, made callable from the device by emitting it with funcQual = hostDev.

Equations

• Sparkle.Backend.CudaSim.deviceEvalTick m = toString "sparkle_" ++ toString (Sparkle.Backend.CSim.sanitizeName m.name) ++ toString "_eval_tick"

def Sparkle.Backend.CudaSim.emitCudaDeviceCode (m : IR.AST.Module) : String

Emit the device-side design code: exactly CSim's toCDesign, but with every module function qualified __host__ __device__. With CSim's default funcQual = "" this would be the CPU backend byte-for-byte, so the device code cannot diverge from CSim's semantics — it is CSim compiled for the device.

This supersedes #33's emitCudaStateStruct + emitCudaDeviceEvalTick (class-wrapper) path, and because it reuses CSim's word-array layout it also handles wide (> 64-bit) state, which the old uint3/uint4 mapping could not (the RV32 578-bit bundle).

Equations

• Sparkle.Backend.CudaSim.emitCudaDeviceCode m = Sparkle.Backend.CSim.toCDesign { topModule := m.name, modules := [m] } none Sparkle.Backend.CudaSim.hostDev✝

def Sparkle.Backend.CudaSim.emitCudaDeviceCodeD (d : IR.AST.Design) : String

Emit the device code for a whole Design — every module, in dependency order, host+device qualified. This is what a hierarchical design needs (e.g. a systolic array whose top instantiates N×N PE sub-modules): the PE's struct + sparkle_PE_eval must be emitted too, and the top's eval_tick already contains the generated wire-copy between instances (CSim lowers each .inst to inst.a_in = <neighbour wire>; …; eval(&inst); <wire> = inst.a_out;). emitCudaDeviceCode m is the single-module case of this.

Equations

• Sparkle.Backend.CudaSim.emitCudaDeviceCodeD d = Sparkle.Backend.CSim.toCDesign d none Sparkle.Backend.CudaSim.hostDev✝

def Sparkle.Backend.CudaSim.emitCudaBatchKernel (m : IR.AST.Module) : String

Emit a CUDA global kernel that runs N independent simulation instances. Each CUDA thread handles one instance (thread index → state array index).

Launch pattern: dim3 blocks((N + 255) / 256, 1, 1); dim3 threads(256, 1, 1); className_batch_kernel<<<blocks, threads>>>(d_states, N, numCycles);

Equations

• One or more equations did not get rendered due to their size.

def Sparkle.Backend.CudaSim.emitCudaJITHostAPI (m : IR.AST.Module) : String

Emit the host-side extern "C" functions that mirror the existing CppSim JIT ABI. These are compiled into the .so loaded by dlopen.

New symbols added by the CUDA backend: jit_cuda_alloc — allocate device state array (N instances) jit_cuda_free — free device state array jit_cuda_set_input — write one input into host-pinned staging buffer jit_cuda_get_output — read one output from host-pinned staging buffer jit_cuda_run — copy H→D, launch kernel for numCycles, copy D→H jit_cuda_reset — reset all instances on device

Equations

• One or more equations did not get rendered due to their size.

def Sparkle.Backend.CudaSim.toCudaSim (m : IR.AST.Module) (_cppHeaderName : String := "") : String

Generate a self-contained .cu for a single Module (no sub-instances). _cppHeaderName is accepted for source compatibility with #33/#37 call sites but is unused — nothing is #included.

Equations

• One or more equations did not get rendered due to their size.

def Sparkle.Backend.CudaSim.toCudaSimDesign (d : IR.AST.Design) (_cppHeaderName : String := "") : String

Generate a self-contained .cu for a full Design. Emits EVERY module (not just the top): for a hierarchical design the top's eval_tick calls the sub-modules' eval, so their structs + functions must be present in the same translation unit. The batch kernel and host API target the top, whose fused device struct embeds every instance — so poking the top's input ports and stepping the kernel drives the whole hierarchy.

This is the path a systolic array / PE-mesh design takes: the PE-to-PE wire-copy is generated by CSim inside the top's eval_tick, not written by hand.

Equations

• One or more equations did not get rendered due to their size.