Module ring_kernel 

Ring kernel configuration and code generation for persistent actor kernels.

This module provides the infrastructure for generating CUDA code for persistent actor kernels that process messages in a loop.

Overview

Ring kernels are persistent GPU kernels that:

• Run continuously until terminated
• Process messages from input queues
• Produce responses to output queues
• Use HLC (Hybrid Logical Clocks) for causal ordering
• Support kernel-to-kernel (K2K) messaging

Generated Kernel Structure

extern "C" __global__ void ring_kernel_NAME(
    ControlBlock* __restrict__ control,
    unsigned char* __restrict__ input_buffer,
    unsigned char* __restrict__ output_buffer,
    void* __restrict__ shared_state
) {
    // Preamble: thread setup
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    // Persistent message loop
    while (!atomicLoad(&control->should_terminate)) {
        if (!atomicLoad(&control->is_active)) {
            __nanosleep(1000);
            continue;
        }

        // Message processing...
        // (user handler code inserted here)
    }

    // Epilogue: mark terminated
    if (tid == 0) {
        atomicStore(&control->has_terminated, 1);
    }
}

Structs

CudaDispatchTable

Dispatch table for multi-handler kernel code generation.

CudaHandlerInfo

Handler registration for CUDA code generation.

KernelReductionConfig

Configuration for global reductions in a ring kernel.

RingKernelConfig

Configuration for a ring kernel.

Enums

RingKernelIntrinsic

Intrinsic functions available in ring kernel handlers.

Functions

generate_control_block_struct

Generate CUDA ControlBlock struct definition.

generate_extended_h2k_message_struct

Generate the ExtendedH2KMessage struct for handler dispatch.

generate_handler_dispatch_code

Generate CUDA switch statement for handler dispatch.

generate_hlc_struct

Generate CUDA HLC helper struct.

generate_k2k_structs

Generate CUDA K2K routing structs and helper functions.

generate_message_envelope_structs

Generate CUDA MessageEnvelope structs matching ringkernel-core layout.

generate_multi_handler_kernel

Generate a complete multi-handler kernel with dispatch table.