Forem: Michael Ivertowski

DotCompute RC2 — Cross-Backend GPU Compute for .NET

Michael Ivertowski — Thu, 06 Nov 2025 10:56:15 +0000

Another framework? Yes.

Another abstraction layer that hates your cache and your startup time? No.

DotCompute’s first release candidate is ready. It gives you GPU and CPU acceleration from C# with:

One kernel definition → multiple backends (CPU SIMD, CUDA, Metal, OpenCL)
Native AOT-first design for .NET 9+
Modern C# API using [Kernel] attributes and source generators
Automatic backend selection based on data size and available hardware

If you’re a .NET dev who wants to use the GPU without writing a second language (or giving up AOT), this is for you.

What DotCompute is (in one paragraph)

DotCompute is a universal compute framework for .NET 9+: you write kernels in C#, mark them with [Kernel], and let the runtime decide whether to run them on CPU (AVX2/AVX512), CUDA, Metal or fallback CPU. The compiler pipeline is built around source generators and Roslyn analyzers, with cross-backend validation and telemetry baked in. :contentReference[oaicite:4]{index=4}

No DSL, no separate shader projects, no “almost C#” language.

Installing the RC

Minimal setup:

dotnet new console -n MyFirstKernel
cd MyFirstKernel

dotnet add package DotCompute.Core
dotnet add package DotCompute.Abstractions
dotnet add package DotCompute.Memory
dotnet add package DotCompute.Backends.CPU
dotnet add package DotCompute.Backends.CUDA   # optional, NVIDIA
dotnet add package DotCompute.Backends.Metal  # optional, Apple Silicon
dotnet add package DotCompute.Generators
dotnet add package DotCompute.Runtime

Target .NET 9 + C# 13 in your .csproj:

<PropertyGroup>
  <TargetFramework>net9.0</TargetFramework>
  <LangVersion>13.0</LangVersion>
</PropertyGroup>

Straight from the getting-started guide; no magic project templates required. (mivertowski.github.io)

Your first kernel (real one, not a toy)

You can get from zero to a working GPU-capable kernel in roughly one coffee:

using DotCompute;

public static class Kernels
{
    /// <summary>
    /// result[i] = a[i] + b[i]
    /// </summary>
    [Kernel]
    public static void VectorAdd(
        ReadOnlySpan<float> a,
        ReadOnlySpan<float> b,
        Span<float> result)
    {
        int idx = Kernel.ThreadId.X;

        if (idx < result.Length)
        {
            result[idx] = a[idx] + b[idx];
        }
    }
}

Key rules:

static, void, [Kernel]
ReadOnlySpan<T> for inputs, Span<T> for outputs
Use Kernel.ThreadId for indexing
Always bounds-check (the GPU will not feel bad for you) (mivertowski.github.io)

Wire it up with the runtime:

using DotCompute.Abstractions;
using DotCompute.Runtime;
using Microsoft.Extensions.DependencyInjection;
using Microsoft.Extensions.Hosting;

var host = Host.CreateDefaultBuilder(args)
    .ConfigureServices(services =>
    {
        services.AddDotComputeRuntime();
    })
    .Build();

var orchestrator = host.Services.GetRequiredService<IComputeOrchestrator>();

const int size = 1_000_000;
var a = Enumerable.Range(0, size).Select(i => (float)i).ToArray();
var b = Enumerable.Range(0, size).Select(i => (float)i * 2).ToArray();
var result = new float[size];

await orchestrator.ExecuteKernelAsync(
    kernelName: "VectorAdd",
    args: new object[] { a, b, result });

On build, the source generator emits:

A CPU SIMD implementation
A CUDA kernel (if the backend is present)
A Metal shader (if you’re on Apple Silicon)
Registration glue so the runtime can find and run the kernel (mivertowski.github.io)

The orchestrator then picks CPU or GPU based on data size and hardware. Small arrays stay on CPU, big arrays go to the fastest device that exists. (mivertowski.github.io)

Why you might care

A few use cases where DotCompute actually solves problems instead of being yet another abstraction:

High-throughput numerics

Vector ops, matrix multiplies, simulations—offload to CUDA/Metal when the payload makes it worth it, stay on CPU for “toy sized” data without changing code. (GitHub)
Hybrid services

.NET microservices that sometimes need a GPU burst: keep the service in C#, plug DotCompute in as a library instead of building a separate Python sidecar.
Cross-platform compute

Same codebase running on a Linux server with NVIDIA, a MacBook (Metal), or just CPU. The framework is designed as AOT-friendly for lean deployment. (GitHub)
Future: persistent kernels and GPU “actors”

The same infrastructure underpins RingKernels and persistent kernels—launch once, push messages forever. That gets its own post, but RC already lays the groundwork.

Production-leaning bits (even in an RC)

The RC isn’t just a “demo build”. Today you already get: (GitHub)

Unified memory management with pooling to cut allocations
Cross-backend validation to compare CPU vs GPU results
Telemetry & profiling hooks for execution timing
Native AOT support with trimming-friendly design

It’s still a release candidate, not a court-admissible standard library, but it’s aimed at real workloads, not just blog benchmarks.

Where to start

Docs / Quickstart:

Getting Started guide (10-minute kernel):

👉 https://mivertowski.github.io/DotCompute/docs/articles/getting-started.html (mivertowski.github.io)
Source / issues / stars (the usual ritual):

👉 https://github.com/mivertowski/DotCompute (GitHub)
Coming next on the blog:
- The beauty of persistent kernels (RingKernels deep dive)
- Real-world scenarios (graphs, audits, simulations)
- Under-the-hood: how the generator pipeline works

If you try DotCompute and something breaks, that’s… good. For now.

File an issue, tell it what you attempted, and we can make the RC a bit less “R” and a bit more “just ship it”.

CUDA Kernel Execution Debugging Journey

Michael Ivertowski — Thu, 04 Sep 2025 12:41:23 +0000

Short version: we went from 8/70 passing CUDA tests to a stable, auditable path by fixing NVRTC name resolution, argument marshaling, and unified-memory sync in DotCompute. No mysticism—just careful pointers and fewer foot-guns.

TL;DR

NVRTC will happily mangle your kernel names. Resolve them explicitly.
CUDA expects pointers to values for kernel params (yes, even device pointers).
Unified memory needs synchronization before CPU access.
Track every unmanaged allocation; free it.
Tests climbed from 8/70 → 41/70 → aiming 95%+.

The Starting Point

The hardware test suite was a parade of classics:

“named symbol not found” at launch
GCHandle pinning failures (non-blittable types)
CUDA 700 (illegal address) on kernel calls
Unified memory access violations
A demoralizing 8/70 green checks

What Actually Fixed Things

1) NVRTC name mangling (resolved)

Problem: extern "C" didn’t save us—NVRTC produced mangled names while our loader looked for unmangled ones.

Fix: Register name expressions before compile, then retrieve lowered names after.

// Register before compilation
foreach (var funcName in functionNames)
{
    var nameExpr = $"&{funcName}";
    NvrtcInterop.nvrtcAddNameExpression(program, nameExpr);
}

// After compilation, retrieve the lowered (mangled) name
var lowered = NvrtcInterop.GetLoweredName(program, nameExpr);
mangledNames[funcName] = lowered;

Impact: Jumps us to 41+ / 70 passing tests.

2) Marshaling unified memory arguments

Problem: CudaUnifiedMemoryBuffer<T> isn’t blittable; direct pinning fails.

Fix: Reflect out the device pointer and pass that (see §3 for the pointer-to-value nuance).

if (argValue.GetType().Name.StartsWith("CudaUnifiedMemoryBuffer"))
{
    var prop = argType.GetProperty("DevicePointer",
        BindingFlags.Public | BindingFlags.NonPublic | BindingFlags.Instance);

    if (prop?.GetValue(argValue) is IntPtr devicePtr)
    {
        unsafe
        {
            var storage = Marshal.AllocHGlobal(sizeof(IntPtr));
            *(IntPtr*)storage = devicePtr;    // store value
            unmanagedAllocations.Add(storage); // remember to free
            return storage;                    // return pointer to value
        }
    }
}

3) The critical param passing bug

Problem: We passed device pointer values directly to cuLaunchKernel. CUDA wants an array of pointers to values.

Fix: Allocate space, write the value, pass a pointer to that space.

unsafe
{
    var storage = Marshal.AllocHGlobal(sizeof(IntPtr));
    *(IntPtr*)storage = devicePtr;   // write the value
    unmanagedAllocations.Add(storage);
    return storage;                  // CUDA reads the value from here
}

Symptom this kills: persistent CUDA 700 on matrix-mult tests.

4) Unmanaged memory hygiene

Problem: Leaks from tiny per-arg allocations.

Fix: Track and free religiously.

var unmanagedAllocations = new List<IntPtr>();

try
{
    var argPtr = PrepareKernelArgument(arg, handles, unmanagedAllocations);
    // ... launch ...
}
finally
{
    foreach (var p in unmanagedAllocations)
        if (p != IntPtr.Zero) Marshal.FreeHGlobal(p);
}

5) Unified memory: sync before you touch

Problem: CPU reading stale/remote pages.

Fix: Gate host spans behind an explicit sync.

public override Span<T> GetSpan()
{
    EnsureOnHost(); // synchronize first
    return new Span<T>((void*)_hostPtr, Length);
}

Deep Dives (the “why”)

How CUDA reads kernel arguments

Wrong: argPointers[i] = devicePtr;
Right: argPointers[i] = &devicePtr;

cuLaunchKernel dereferences your argPointers to fetch the actual values. Device pointers are values too—treat them like any other scalar.

C# overload resolution trap

Without the generic arg, a params overload may win by accident:

LaunchAsync(config, buf1, buf2, buf3, size);      // may hit wrong overload
LaunchAsync<float>(config, buf1, buf2, buf3, size); // forces intended path

Lessons (written on a sticky note)

700 ≠ haunted memory — often just wrong argument plumbing.
P/Invoke is sharp — mind blittability, lifetimes, and double-indirection.
GPU is async by default — sync before the CPU peeks.
Reflection is a tool, not a lifestyle — but it saved us here.
Iterate mercilessly — fix → run → commit → repeat.

Status & Performance

Start: 8 / 70 (11.4%)
After NVRTC fix: 41 / 70 (58.6%)
Targeting 95%+ with the remaining stragglers (edge cases & perf polish).

Production Checklist

[ ] Every CUDA/NVRTC call checks return codes
[ ] All unmanaged allocations tracked & freed
[ ] Timers/metrics around compile, HtoD/DtoH, and kernel time
[ ] Launch config validated against device caps
[ ] Tests for edge cases, stress, and multi-GPU

What’s Next

Dynamic parallelism (flag plumbing + tests)
Faster transfers for bidi workloads
CUDA Graphs to amortize launch overhead
Texture/constant memory where patterns fit
Nsight-based profiling in CI

Appendix: File-level changes

CudaKernelCompiler.cs — NVRTC name resolution
CudaKernelLauncher.cs — argument marshaling, lifetime fixes
CudaUnifiedMemoryBuffer.cs — host-access synchronization
CudaKernelExecutionTests.cs — overload resolution fixed

Key APIs we leaned on

nvrtcAddNameExpression, nvrtcGetLoweredName
Marshal.AllocHGlobal, Marshal.FreeHGlobal, GCHandle.Alloc
cudaDeviceSynchronize (and friends)

Handy error codes

NVRTC_ERROR_COMPILATION — syntax/flags/headers
700 — illegal address (often arg passing)
716 — misaligned address
719 — launch failure

Credit where due: ClaudeCode helped pattern-match the error soup, surface the right docs, and keep the loop tight. The wins are still boring engineering ones—our favorite kind.

Event-Sourced State Machines in Orleans

Michael Ivertowski — Wed, 03 Sep 2025 08:36:14 +0000

The good thing about agentic coding is, you can easily awake dead repos out of experiments back to life. Event sourced grain state machines is such an example.

Short version: take a clean Stateless state machine, host it as a grain, and let events be the source of truth. You get audit, replay, and fewer “what just happened?” moments.

Why this combo?

Stateless gives a tidy, strongly-typed FSM API. GitHub
Orleans gives virtual actors, timers/reminders, and consistent single-threaded execution per grain. Microsoft Learn
Orleans.StateMachineES wraps them and adds event sourcing, snapshots, and nice enterprise extras like streams & versioning. GitHub

Install

dotnet add package Orleans.StateMachineES

(See README for exact version and notes.) GitHub

1) Define states, triggers, and events

public enum OrderState { Created, Paid, Shipped, Cancelled }
public enum OrderTrigger { Pay, Ship, Cancel }

// Event payloads (source of truth)
public abstract record OrderEvent(DateTime OccurredAt);
public record OrderPaid(decimal Amount, DateTime OccurredAt) : OrderEvent(OccurredAt);
public record OrderShipped(string Carrier, DateTime OccurredAt) : OrderEvent(OccurredAt);
public record OrderCancelled(string Reason, DateTime OccurredAt) : OrderEvent(OccurredAt);

// Grain state (can be rebuilt from events)
public record OrderGrainState(OrderState State, decimal TotalPaid, string? LastCarrier);

2) Implement an event-sourced state machine grain

using Orleans.StateMachineES;
using Stateless;
using static OrderState;
using static OrderTrigger;

public interface IOrderGrain : IGrainWithStringKey
{
    Task<OrderState> GetState();
    Task Pay(decimal amount, string idempotencyKey);
    Task Ship(string carrier, string idempotencyKey);
    Task Cancel(string reason, string idempotencyKey);
}

[StorageProvider(ProviderName = "Default")]
public sealed class OrderGrain
    : EventSourcedStateMachineGrain<OrderState, OrderTrigger, OrderGrainState>,
      IOrderGrain
{
    protected override void ConfigureEventSourcing(EventSourcingOptions opt)
    {
        // Event sourcing knobs (see README)
        opt.AutoConfirmEvents = true;           // apply+persist in one go
        opt.SnapshotInterval = 100;             // every N events
        opt.EnableIdempotency = true;           // dedupe by key
    }

    protected override StateMachine<OrderState, OrderTrigger> BuildStateMachine()
    {
        var sm = new StateMachine<OrderState, OrderTrigger>(Created);

        sm.Configure(Created)
            .Permit(Pay, Paid)
            .Permit(Cancel, Cancelled);

        sm.Configure(Paid)
            .Permit(Ship, Shipped)
            .Permit(Cancel, Cancelled);

        sm.Configure(Shipped); // terminal for this sample
        sm.Configure(Cancelled); // terminal

        return sm;
    }

    // Map incoming triggers to domain events
    public async Task Pay(decimal amount, string key)
        => await FireAsync(Pay, key, new OrderPaid(amount, DateTime.UtcNow));

    public async Task Ship(string carrier, string key)
        => await FireAsync(Ship, key, new OrderShipped(carrier, DateTime.UtcNow));

    public async Task Cancel(string reason, string key)
        => await FireAsync(Cancel, key, new OrderCancelled(reason, DateTime.UtcNow));

    public Task<OrderState> GetState() => GetStateAsync();

    // Apply events to rebuild state
    protected override OrderGrainState Apply(OrderGrainState state, object @event) => @event switch
    {
        OrderPaid e      => state with { State = Paid,     TotalPaid = state.TotalPaid + e.Amount },
        OrderShipped e   => state with { State = Shipped,  LastCarrier = e.Carrier },
        OrderCancelled _ => state with { State = Cancelled },
        _ => state
    };
}

Notes

EventSourcedStateMachineGrain<…> follows Orleans’ event-sourcing model (akin to JournaledGrain) but wrapped for FSMs. You rebuild state by applying events; snapshots keep load times sane. Microsoft Learn+1
FireAsync(trigger, idempotencyKey, event) (API per README) drives transitions and persists the event once confirmed. Idempotency avoids double-fire under retries. GitHub

3) Minimal client call flow

var grain = client.GetGrain<IOrderGrain>("ORD-42");
await grain.Pay(99.00m, "pay#txn-123");
await grain.Ship("DHL", "ship#awb-987");
var s = await grain.GetState(); // Shipped

4) Timeouts with timers & reminders (SLA guards)

Use a reminder to auto-cancel unshipped orders after e.g. 24h—reminders survive deactivation and are durable. Microsoft Learn

public override async Task OnActivateAsync(CancellationToken ct)
{
    await base.OnActivateAsync(ct);
    await RegisterOrUpdateReminder("ship-deadline", TimeSpan.FromHours(24), TimeSpan.FromDays(365));
}

public async Task OnReminder(string reminderName)
{
    if (reminderName == "ship-deadline" && (await GetState()) == OrderState.Paid)
        await Cancel("Shipping deadline exceeded", "auto-cancel#ship-deadline");
}

(Replace with the reminder hook your host template uses; the idea stands.) Microsoft Learn

5) Streaming transitions (optional, but great for audits)

Publish transitions to Orleans Streams for a live audit trail and dashboards. The fork advertises stream integration; wire an IAsyncStream<OrderEvent> and OnTransition hook to push events. GitHub

6) Storage choice (TL;DR)

StateStorage provider: persists snapshots of state; leaner loads.
LogStorage provider: replays full event history each activation; great for rigorous audit/rebuild scenarios. Pick based on load profile and audit needs. Microsoft Learn

7) Quick benchmark sanity check

If you care about throughput, enable AutoConfirmEvents. The README’s example shows a notable bump (it reduces round-trips/allocations on the hot path). Your mileage will vary—measure in your cluster. GitHub

8) A small checklist before prod

Idempotency keys for all external-facing triggers.
Snapshot interval appropriate for event volume.
Durable deadlines use reminders, not timers. Microsoft Learn
Log state transitions (streams or your sink of choice).
Tests around guards and parameterized triggers (Stateless makes this easy). GitHub

Appendix: bare-bones non-ES grain (for contrast)

public sealed class SimpleOrderGrain
    : StateMachineGrain<OrderState, OrderTrigger>, IOrderGrain
{
    protected override StateMachine<OrderState, OrderTrigger> BuildStateMachine()
        => new StateMachine<OrderState, OrderTrigger>(OrderState.Created)
            .Configure(OrderState.Created).Permit(OrderTrigger.Pay, OrderState.Paid)
            .Permit(OrderTrigger.Cancel, OrderState.Cancelled)
            .Configure(OrderState.Paid).Permit(OrderTrigger.Ship, OrderState.Shipped);
}

Works. Just not auditable.