.NET Accelerator MVP

“`org

1 0. Document Contract

This document is not merely an architecture proposal.

It is the implementation workbook for the first .NET Accelerator MVP.

The intended implementer is a junior developer who:

can build C# and C++;
can use git;
can run a debugger;
understands arrays, pointers, and basic control flow;
does NOT need prior CoreCLR experience;
does NOT need prior compiler implementation experience;
does NOT need prior CUDA programming experience.

The junior developer MUST NOT make architectural decisions while implementing this plan.

If a ticket appears to require an architectural decision not answered here, the ticket is considered underspecified and must be escalated rather than invented locally.

The MVP is deliberately narrow.

Its purpose is to prove this exact statement:

A normal C# method can be compiled to IL by the normal C# compiler, identified inside a forked CoreCLR process, imported from its MethodDesc/IL representation into a small accelerator IR, translated to PTX, launched on an NVIDIA GPU, and produce the same result as a CPU reference implementation.

The first successful kernel is elementwise float addition.

The MVP does NOT initially attempt transparent offload, TensorPrimitives, Span<T>, Vector<T>, GC objects on the GPU, RyuJIT GPU code generation, or automatic vectorization.

Those are explicitly later phases.

2 1. Definition of “Shovel Ready”

Every implementation ticket in this document defines:

Objective
Why the ticket exists
Prerequisites
Exact files to create or edit
Existing runtime code to imitate
Exact interfaces or data structures
Implementation algorithm
Build command
Test command
Expected output
Expected failure modes
Debug procedure
Definition of done
Commit boundary

A junior developer should normally implement one ticket and create one commit.

Do NOT combine tickets unless instructed.

3 2. Locked Architectural Decisions

These decisions are final for MVP 1.

Do not reopen them during implementation.

3.1 2.1 Host operating system

MVP development platform:

Ubuntu Linux
x86-64
CoreCLR
NVIDIA GPU

Windows support is deferred.

macOS is not applicable to this NVIDIA MVP.

ARM64 is deferred.

3.2 2.2 GPU backend

Use:

NVIDIA CUDA Driver API
PTX text generation

Do NOT use:

CUDA Runtime API
CUDA C++
NVCC
NVRTC
Vulkan
SPIR-V
DirectX
ComputeSharp
ILGPU
LLVM
MLIR

The compiler directly emits PTX text.

The NVIDIA driver JITs PTX to native device code.

3.3 2.3 CoreCLR fork

The MVP WILL modify dotnet/runtime.

The MVP WILL NOT initially modify RyuJIT.

All accelerator code remains behind an explicit runtime API.

Therefore normal methods continue through:

IL
 |
 v
RyuJIT
 |
 v
x64 code

Accelerator compilation explicitly goes through:

RuntimeMethodHandle
 |
 v
MethodDesc
 |
 v
IL
 |
 v
Accelerator Importer
 |
 v
Accelerator IR
 |
 v
PTX
 |
 v
CUDA Driver
 |
 v
GPU

3.4 2.4 No automatic kernel discovery

There is no automatic attribute scanning.

There is no JIT interception.

Host code explicitly requests accelerator compilation for a MethodInfo.

Conceptually:

MethodInfo method = typeof(Kernels).GetMethod("Add")!;

ulong kernel = AcceleratorRuntime.CompileKernel(
    method.MethodHandle,
    ...intrinsic metadata tokens...);

3.5 2.5 No public BCL API

Do not attempt to design the final public .NET API.

The experimental managed bridge lives inside System.Private.CoreLib as an internal class.

Tests call that internal class through reflection.

This avoids:

changing public reference assemblies;
API review questions;
API compatibility failures;
prematurely designing System.Runtime.Accelerators.

The public API comes after compiler feasibility is demonstrated.

3.6 2.6 No Span<T> in MVP

The first kernel ABI is deliberately primitive.

A kernel accepts:

ulong inputADevicePointer
ulong inputBDevicePointer
ulong outputDevicePointer
int elementCount

The C# kernel does NOT dereference these pointers directly.

Instead it calls compiler-recognized intrinsic methods.

3.7 2.7 Exactly three compiler intrinsics initially

The first importer recognizes only:

AcceleratorIntrinsics.GlobalIndexX()

AcceleratorIntrinsics.LoadF32(
    ulong deviceAddress,
    int elementIndex)

AcceleratorIntrinsics.StoreF32(
    ulong deviceAddress,
    int elementIndex,
    float value)

These functions exist as ordinary C# stub methods in the test assembly.

Their metadata tokens are passed explicitly to the accelerator compiler.

The importer therefore does NOT need to implement arbitrary metadata call resolution in MVP.

3.8 2.8 No SSA requirement for MVP

The first IR is a typed CFG/value IR.

It is NOT full SSA.

This is intentional.

A compiler value receives an immutable ValueId, but local variables may be tracked by the importer as current ValueIds.

MVP restrictions ensure merge points do not require phi nodes.

Phi nodes are a follow-up.

3.9 2.9 No optimizer

The initial pipeline is:

IL
 |
 v
decode
 |
 v
validate
 |
 v
import
 |
 v
verify IR
 |
 v
emit PTX

There is:

no DCE
no CSE
no constant propagation pass
no LICM
no vectorization
no inlining

The CUDA driver may optimize generated PTX.

3.10 2.10 Explicit GPU residency

The host explicitly allocates device memory.

There is no automatic RAM -> VRAM movement.

This remains an architectural principle.

3.11 2.11 One process-global CUDA device

MVP uses:

device ordinal: 0
primary CUDA context
default CUDA stream

No multi-GPU.

No streams API.

No asynchronous copies.

No concurrent kernel launch management.

3.12 2.12 Synchronous execution

Every kernel launch is followed by:

cuCtxSynchronize

before returning to managed code.

Performance optimization comes later.

4 3. Repository Baseline

4.1 3.1 Fork

Fork:

dotnet/runtime

Clone the fork.

Example:

git clone git@github.com:<your-user>/runtime.git
cd runtime

Add upstream:

git remote add upstream https://github.com/dotnet/runtime.git
git fetch upstream

4.2 3.2 Pin the baseline

Do NOT continuously develop against moving upstream/main.

At project start:

git checkout main
git pull upstream main
git rev-parse HEAD

Record that SHA in:

docs/design/features/accelerator-mvp-baseline.txt

File contents:

UPSTREAM_COMMIT=<40-character-sha>
MVP_VERSION=1
PLATFORM=linux-x64
BACKEND=nvidia-ptx

All implementation tickets are based on this SHA.

4.3 3.3 Branch

Create:

git checkout -b feature/accelerator-mvp

Do not develop directly on main.

5 4. Canonical Local Build

5.1 4.1 Runtime

From repository root:

./build.sh -subset clr -configuration Checked

5.2 4.2 Libraries

./build.sh -subset libs -configuration Release

5.3 4.3 Core_Root

./src/tests/build.sh \
    -arch x64 \
    -checked \
    -generatelayoutonly

Canonical Core_Root:

export CORE_ROOT="$PWD/artifacts/tests/coreclr/linux.x64.Checked/Tests/Core_Root"

Verify:

test -x "$CORE_ROOT/corerun"
test -f "$CORE_ROOT/libcoreclr.so"

Both commands must exit zero.

5.4 4.4 Never install the fork system-wide

Do NOT:

sudo make install

Do NOT replace the machine-wide dotnet runtime.

All tests use:

$CORE_ROOT/corerun

or the generated runtime test launcher.

6 5. Required Hardware Validation

Before touching CoreCLR source, run:

nvidia-smi

Definition of success:

command exits zero;
at least one NVIDIA GPU appears;
driver is loaded.

The accelerator runtime itself will dynamically load:

libcuda.so.1

Do not link the runtime against a CUDA runtime library.

7 6. Final MVP File Layout

Create the following native directory:

src/coreclr/vm/accelerator/

Final expected files:

src/coreclr/vm/accelerator/
  acceleratorstatus.h
  acceleratornative.h
  acceleratornative.cpp

  cudatypes.h
  cudadriver.h
  cudadriver.cpp

  acceleratorir.h
  acceleratorir.cpp
  acceleratorirprinter.h
  acceleratorirprinter.cpp
  acceleratorirverifier.h
  acceleratorirverifier.cpp
  acceleratorinterpreter.h
  acceleratorinterpreter.cpp

  ildecoder.h
  ildecoder.cpp
  acceleratorimporter.h
  acceleratorimporter.cpp

  ptxcodegen.h
  ptxcodegen.cpp

Managed runtime bridge:

src/coreclr/System.Private.CoreLib/src/
  System/Runtime/CompilerServices/
    AcceleratorRuntime.CoreCLR.cs

Tests:

src/tests/Accelerator/
  Smoke/
    Smoke.csproj
    Program.cs
    RuntimeBridge.cs

  Kernel/
    Kernel.csproj
    Program.cs
    RuntimeBridge.cs
    AcceleratorIntrinsics.cs
    Kernels.cs

Documentation/debug scripts:

eng/accelerator/
  update-coreroot.sh
  dump-env.sh

docs/design/features/
  accelerator-mvp-baseline.txt
  accelerator-mvp.md

8 7. Native Status Contract

Create:

src/coreclr/vm/accelerator/acceleratorstatus.h

Exact enum:

#pragma once

enum class AcceleratorStatus : int
{
    Ok = 0,

    DriverUnavailable = 1,
    DriverInitializationFailed = 2,
    DeviceUnavailable = 3,
    ContextFailure = 4,

    InvalidArgument = 10,
    OutOfMemory = 11,

    ModuleLoadFailed = 20,
    FunctionLookupFailed = 21,
    KernelLaunchFailed = 22,

    InvalidMethod = 30,
    UnsupportedMethod = 31,
    UnsupportedIL = 32,
    InvalidIL = 33,

    InvalidIR = 40,
    PTXGenerationFailed = 41,

    InternalError = 100
};

Never invent additional status codes ad hoc.

Add a code only when a ticket explicitly requires it.

9 8. CUDA ABI Contract

Create:

src/coreclr/vm/accelerator/cudatypes.h

The MVP declares only the CUDA types it uses.

#pragma once

#include <cstddef>
#include <cstdint>

using CUresult = int;
using CUdevice = int;

using CUcontext = void*;
using CUmodule = void*;
using CUfunction = void*;
using CUstream = void*;

using CUdeviceptr = std::uint64_t;

static constexpr CUresult CUDA_SUCCESS = 0;

Do not reproduce the complete CUDA headers.

9.1 8.1 Required CUDA symbols

CudaDriver must load exactly these logical operations:

cuInit
cuDeviceGetCount
cuDeviceGet
cuDeviceGetName

cuDevicePrimaryCtxRetain
cuDevicePrimaryCtxRelease
cuCtxSetCurrent
cuCtxSynchronize

cuMemAlloc
cuMemFree
cuMemcpyHtoD
cuMemcpyDtoH

cuModuleLoadData
cuModuleUnload
cuModuleGetFunction

cuLaunchKernel

cuGetErrorName
cuGetErrorString

For versioned driver symbols, the loader may attempt the versioned symbol first and the unversioned name second.

Example:

cuMemAlloc_v2
then
cuMemAlloc

The same strategy applies to versioned memcpy/free operations where necessary.

10 9. CUDA Driver Ownership Rules

There is exactly one process-global:

CudaDriver

It owns:

PAL dynamic-library handle for libcuda.so.1
CUdevice for ordinal 0
retained primary CUcontext
loaded function pointers

It does NOT own:

device allocations
CUmodule instances returned to callers
CUfunction independently of CUmodule

Rules:

Initialize lazily.
Initialization is idempotent.
Device ordinal is always 0.
Retain the device primary context once.
Make the primary context current before operations.
Release it only in explicit Accelerator_Shutdown during MVP.
Process termination is acceptable cleanup fallback during failed tests.

There is no background thread.

There is no CUDA callback.

11 10. Kernel Handle Contract

A native kernel handle is:

struct AcceleratorKernelHandle
{
    CUmodule module;
    CUfunction function;
};

Externally it is represented as:

uintptr_t / ulong

The pointer refers to a native heap allocation of AcceleratorKernelHandle.

Ownership:

Compile/Load
  |
  v
caller owns handle
  |
  v
Accelerator_DestroyKernel
  |
  +--> cuModuleUnload
  +--> delete AcceleratorKernelHandle

CUfunction is not independently freed.

12 11. Device Allocation Contract

Device allocation is represented externally as:

ulong

Internally it is:

CUdeviceptr

Ownership rule:

Accelerator_Allocate
   -> caller owns pointer

Accelerator_Free
   -> ownership ends

Double free is undefined in MVP tests and must not occur.

Zero-byte allocation is rejected with InvalidArgument.

13 12. Managed Internal Bridge Contract

Create:

src/coreclr/System.Private.CoreLib/src/System/Runtime/CompilerServices/
AcceleratorRuntime.CoreCLR.cs

Namespace:

namespace System.Runtime.CompilerServices;

Type:

internal static class AcceleratorRuntime

This is NOT public API.

Required internal methods:

internal static bool IsSupported();

internal static int GetDeviceCount();

internal static ulong Allocate(nuint byteCount);

internal static void Free(ulong devicePointer);

internal static void CopyToDevice(
    ulong destination,
    byte[] source);

internal static void CopyFromDevice(
    byte[] destination,
    ulong source);

internal static ulong LoadPtxKernel(
    byte[] nullTerminatedUtf8Ptx,
    byte[] nullTerminatedUtf8KernelName);

internal static void DestroyKernel(
    ulong kernelHandle);

internal static void LaunchSet42(
    ulong kernelHandle,
    ulong outputDevicePointer);

internal static ulong CompileKernel(
    RuntimeMethodHandle method,
    int globalIndexToken,
    int loadF32Token,
    int storeF32Token);

internal static void LaunchAddF32(
    ulong kernelHandle,
    ulong a,
    ulong b,
    ulong output,
    int length);

internal static void Shutdown();

Private extern methods use QCall.

Example pattern:

[DllImport(
    RuntimeHelpers.QCall,
    EntryPoint = "Accelerator_IsSupported")]
private static extern int IsSupportedNative();

Use primitive QCall arguments only.

Managed wrappers perform:

byte[] pinning;
conversion from RuntimeMethodHandle.Value to IntPtr;
AcceleratorStatus checking;
throwing InvalidOperationException when status != Ok.

Do not pass managed object references directly into native accelerator code.

14 13. Test Reflection Bridge

Normal test projects cannot reference the internal AcceleratorRuntime type.

Therefore create:

RuntimeBridge.cs

in each test directory.

Required implementation pattern:

using System.Reflection;

internal static class RuntimeBridge
{
    private static readonly Type RuntimeType =
        typeof(object).Assembly.GetType(
            "System.Runtime.CompilerServices.AcceleratorRuntime",
            throwOnError: true)!;

    private static MethodInfo Method(string name) =>
        RuntimeType.GetMethod(
            name,
            BindingFlags.Static |
            BindingFlags.NonPublic)!
        ?? throw new MissingMethodException(name);

    internal static bool IsSupported() =>
        (bool)Method("IsSupported").Invoke(null, null)!;

    internal static int GetDeviceCount() =>
        (int)Method("GetDeviceCount").Invoke(null, null)!;

    // Add wrappers ticket-by-ticket.
}

Do NOT make AcceleratorRuntime public merely to simplify tests.

15 14. Accelerator IR Specification v1

15.1 14.1 Types

Exactly these types exist:

Void
Bool
I32
I64
F32
DevicePtr

DevicePtr is represented as 64 bits but is distinct from I64 in the verifier.

15.2 14.2 Value IDs

using ValueId = uint32_t;

ValueId 0 is invalid.

Every value-producing instruction receives a monotonically increasing ValueId.

15.3 14.3 Block IDs

using BlockId = uint32_t;

BlockId 0 is valid and is the entry block.

15.4 14.4 Opcodes

Exactly:

Parameter
ConstantI32

GlobalIndexX

LoadF32
StoreF32

AddF32

CompareGeI32

Branch
BranchIf

Return

Do not add integer multiplication merely because PTX address generation needs it.

Address scaling for LoadF32/StoreF32 is backend semantics.

15.5 14.5 Parameter ABI

Kernel parameter numbers:

0 : DevicePtr : a
1 : DevicePtr : b
2 : DevicePtr : output
3 : I32       : length

No other signature is supported in MVP 1.

15.6 14.6 LoadF32 semantics

LoadF32(base, index)

means:

address = base + (index * 4 bytes)
return *(float*)address

15.7 14.7 StoreF32 semantics

StoreF32(base, index, value)

means:

address = base + (index * 4 bytes)
*(float*)address = value

15.8 14.8 GlobalIndexX semantics

globalIndexX =
    blockIndexX * blockDimensionX
    + threadIndexX

15.9 14.9 CFG restrictions

Every block must end in exactly one of:

Branch
BranchIf
Return

No instruction may appear after a terminator.

Branch targets must exist.

Entry block is block 0.

15.10 14.10 MVP merge restrictions

To avoid phi nodes:

evaluation stack must be empty on every CFG edge;
locals reaching a block from multiple predecessors must have identical ValueIds;
otherwise importer returns UnsupportedIL.

This is not permanent language semantics.

It is an MVP compiler limitation.

16 15. IR Verifier Contract

Verifier checks:

[ ] function has at least one block
[ ] block 0 exists
[ ] every block has one terminator
[ ] no instruction exists after terminator
[ ] every referenced ValueId exists
[ ] every referenced BlockId exists
[ ] operand type matches opcode
[ ] CompareGeI32 operands are I32
[ ] AddF32 operands are F32
[ ] LoadF32 base is DevicePtr
[ ] LoadF32 index is I32
[ ] StoreF32 base is DevicePtr
[ ] StoreF32 index is I32
[ ] StoreF32 value is F32
[ ] BranchIf condition is Bool
[ ] Return has no value

Verifier returns:

AcceleratorStatus::Ok
or
AcceleratorStatus::InvalidIR

Checked builds additionally assert invariant failures internally.

17 16. IR Text Format

IR dump must be deterministic.

Example:

kernel AddF32 {
  block B0:
    %1:DevicePtr = param 0
    %2:DevicePtr = param 1
    %3:DevicePtr = param 2
    %4:I32 = param 3
    %5:I32 = global_index_x
    %6:Bool = cmp_ge_i32 %5, %4
    branch_if %6, B2, B1

  block B1:
    %7:F32 = load_f32 %1, %5
    %8:F32 = load_f32 %2, %5
    %9:F32 = add_f32 %7, %8
    store_f32 %3, %5, %9
    branch B2

  block B2:
    return
}

Whitespace and IDs must remain deterministic so golden tests are stable.

18 17. PTX Contract v1

Emit a module beginning with:

.version 7.0
.target sm_50
.address_size 64

Kernel entry:

.visible .entry AddF32(
    .param .u64 param0,
    .param .u64 param1,
    .param .u64 param2,
    .param .u32 param3
)

Parameter interpretation:

param0 = a device pointer
param1 = b device pointer
param2 = output device pointer
param3 = length

GlobalIndexX lowers to PTX values derived from:

%ctaid.x
%ntid.x
%tid.x

LoadF32:

byteOffset = index * 4
address = base + byteOffset
ld.global.f32

StoreF32 uses:

st.global.f32

Branch conditions use predicate registers.

Do not optimize PTX register count during MVP.

Allocate monotonically numbered registers.

19 18. Launch Contract

Block size is always:

256 threads

For N elements:

gridX = (N + 255) / 256

If N == 0:

do NOT call cuLaunchKernel.

Return success immediately.

Launch dimensions:

grid  = (gridX, 1, 1)
block = (256, 1, 1)
shared memory = 0
stream = nullptr/default

After cuLaunchKernel:

cuCtxSynchronize()

must succeed before returning.

20 19. Canonical Test Kernel Source

Create:

src/tests/Accelerator/Kernel/AcceleratorIntrinsics.cs

Exact source intent:

using System;

internal static class AcceleratorIntrinsics
{
    public static int GlobalIndexX() =>
        throw new PlatformNotSupportedException();

    public static float LoadF32(
        ulong baseAddress,
        int index) =>
        throw new PlatformNotSupportedException();

    public static void StoreF32(
        ulong baseAddress,
        int index,
        float value) =>
        throw new PlatformNotSupportedException();
}

Create:

Kernels.cs

with:

internal static class Kernels
{
    internal static void AddF32(
        ulong a,
        ulong b,
        ulong output,
        int length)
    {
        int i = AcceleratorIntrinsics.GlobalIndexX();

        if (i >= length)
            return;

        float value =
            AcceleratorIntrinsics.LoadF32(a, i) +
            AcceleratorIntrinsics.LoadF32(b, i);

        AcceleratorIntrinsics.StoreF32(
            output,
            i,
            value);
    }
}

Do not change this method to make importer implementation easier without updating the compiler contract.

21 20. Ticket Sequence Overview

Implement in this exact order:

A000 baseline
A001 Core_Root
A002 smoke test project
A003 no-op QCall

A010 CUDA library load
A011 CUDA initialization
A012 primary context
A013 allocation/free
A014 memcpy round-trip
A015 handwritten PTX set42
A016 handwritten PTX vector add

A020 IR model
A021 IR verifier
A022 IR printer
A023 IR interpreter
A024 PTX generation for hand-built IR

A030 method handle bridge
A031 IL extraction
A032 IL decoder
A033 CFG construction
A034 IL stack/local importer
A035 intrinsic recognition
A036 AddF32 IL -> IR

A040 compiled PTX load
A041 compiled AddF32 launch
A042 differential correctness suite
A043 negative compiler tests
A044 diagnostic dumps

A050 subgroup IR extension
A051 shuffle implementation
A052 reduction experiment

A060 performance baseline

MVP-A ends at A044.

MVP-B research validation ends at A052.

22 21. A000 — Pin Upstream Baseline

22.1 Objective

Create reproducible project baseline metadata.

22.2 Files

Create:

docs/design/features/accelerator-mvp-baseline.txt

22.3 Steps

git rev-parse HEAD
uname -a
nvidia-smi

Record:

UPSTREAM_COMMIT=
OS=
ARCH=x64
GPU=
NVIDIA_DRIVER=

22.4 Test

git diff --check

22.5 Done

[ ] exact 40-character upstream commit stored
[ ] GPU model stored
[ ] NVIDIA driver version stored
[ ] file committed alone

Commit:

accelerator: record MVP baseline

23 22. A001 — Build Private Core_Root

23.1 Objective

Prove the unmodified fork can build and run.

23.2 Commands

./build.sh -subset clr -configuration Checked

./build.sh -subset libs -configuration Release

./src/tests/build.sh \
    -arch x64 \
    -checked \
    -generatelayoutonly

Set:

export CORE_ROOT="$PWD/artifacts/tests/coreclr/linux.x64.Checked/Tests/Core_Root"

Verify:

"$CORE_ROOT/corerun" --help >/dev/null 2>&1 || true
test -f "$CORE_ROOT/libcoreclr.so"
test -f "$CORE_ROOT/System.Private.CoreLib.dll"

23.3 Failure procedure

If CLR build fails:

do NOT start accelerator work;
inspect artifacts/log;
verify dotnet/runtime Linux prerequisites;
repair environment;
rerun baseline.

23.4 Done

[ ] CoreCLR Checked builds
[ ] libraries Release build
[ ] Core_Root exists
[ ] corerun exists

No source commit required unless documentation/script added.

24 23. A002 — Add Accelerator Smoke Test

24.1 Objective

Add one isolated runtime test process.

Create:

src/tests/Accelerator/Smoke/Smoke.csproj
src/tests/Accelerator/Smoke/Program.cs

Use this project pattern:

<Project Sdk="Microsoft.NET.Sdk">
  <PropertyGroup>
    <RequiresProcessIsolation>true</RequiresProcessIsolation>
    <ReferenceXUnitWrapperGenerator>false</ReferenceXUnitWrapperGenerator>
    <CLRTestExecutionArguments />
  </PropertyGroup>

  <ItemGroup>
    <Compile Include="Program.cs" />
  </ItemGroup>
</Project>

Program:

using System;

internal static class Program
{
    public static int Main()
    {
        Console.WriteLine("ACCELERATOR_SMOKE_OK");
        return 100;
    }
}

CoreCLR tests conventionally use 100 for success in many standalone tests; follow the local generated test runner convention if it requires that value.

Build:

./src/tests/build.sh \
    -arch x64 \
    -checked \
    -test:Accelerator/Smoke/Smoke.csproj

Run generated test launcher.

If unsure of exact output path:

find artifacts/tests/coreclr/linux.x64.Checked \
    -path '*Accelerator/Smoke/Smoke.sh' \
    -print

Expected output contains:

ACCELERATOR_SMOKE_OK

24.2 Done

[ ] test builds
[ ] test launches under forked CoreCLR
[ ] process exits according to runtime-test success convention

Commit:

tests: add accelerator smoke test

25 24. A003 — Wire First QCall

25.1 Objective

Prove managed CoreLib -> CoreCLR native accelerator bridge.

Create:

src/coreclr/vm/accelerator/acceleratornative.h
src/coreclr/vm/accelerator/acceleratornative.cpp

Native function:

extern "C"
INT32 QCALLTYPE Accelerator_IsSupported()
{
    QCALL_CONTRACT_NO_GC_TRANSITION;

    return 0;
}

Add header include to:

src/coreclr/vm/qcallentrypoints.cpp

Add:

#include "accelerator/acceleratornative.h"

Add registration:

DllImportEntry(Accelerator_IsSupported)

to s_QCall.

Add files to:

src/coreclr/vm/CMakeLists.txt

under VM_SOURCES_WKS / VM_HEADERS_WKS.

Create managed CoreLib class:

AcceleratorRuntime.CoreCLR.cs

Private QCall:

[DllImport(
    RuntimeHelpers.QCall,
    EntryPoint = "Accelerator_IsSupported")]
private static extern int IsSupportedNative();

internal static bool IsSupported()
    => IsSupportedNative() != 0;

Add RuntimeBridge reflection wrapper to Smoke test.

Program prints:

Console.WriteLine(
    $"ACCELERATOR_SUPPORTED={RuntimeBridge.IsSupported()}");

Rebuild:

./build.sh -subset clr -configuration Checked

Regenerate Core_Root because CoreLib changed:

./src/tests/build.sh \
    -arch x64 \
    -checked \
    -generatelayoutonly

Expected:

ACCELERATOR_SUPPORTED=False

25.2 Failure: EntryPointNotFound

Check:

DllImportEntry exists.
acceleratornative.h is included.
exact spelling is Accelerator_IsSupported.
Core_Root contains newly rebuilt libcoreclr.so.

25.3 Failure: managed method missing through reflection

Check:

AcceleratorRuntime.CoreCLR.cs

was compiled into CoreLib and namespace is exactly:

System.Runtime.CompilerServices

25.4 Done

[ ] managed reflection reaches CoreLib class
[ ] QCall reaches native function
[ ] returns False
[ ] no GPU code exists yet

Commit:

coreclr: add accelerator QCall bridge

26 25. A010 — Dynamically Load NVIDIA Driver

26.1 Objective

Load libcuda.so.1 without linking against it.

Create:

cudatypes.h
cudadriver.h
cudadriver.cpp

Use CoreCLR PAL dynamic loader helpers.

Do not add a normal ELF link dependency on libcuda.

CudaDriver fields:

class CudaDriver
{
private:
    void* _library;
    bool _loadAttempted;
    bool _loaded;

public:
    static CudaDriver& Instance();

    bool TryLoad();
};

Exact state semantics:

before TryLoad:
    loadAttempted=false
    loaded=false

successful:
    loadAttempted=true
    loaded=true

failed:
    loadAttempted=true
    loaded=false

subsequent calls:
    return cached state

Library:

libcuda.so.1

Change Accelerator_IsSupported temporarily to:

return CudaDriver::Instance().TryLoad() ? 1 : 0;

Expected on GPU development machine:

ACCELERATOR_SUPPORTED=True

Test CPU/no NVIDIA driver separately when possible:

ACCELERATOR_SUPPORTED=False

Do not throw if driver absent.

26.2 Debug

Set LLDB breakpoint:

CudaDriver::TryLoad

If load fails despite nvidia-smi working:

inspect actual loader call and error.

26.3 Done

[ ] no static CUDA link dependency
[ ] driver loads
[ ] absence returns false
[ ] repeated calls are stable

Commit:

accelerator: dynamically load CUDA driver

27 26. A011 — Resolve CUDA Initialization Symbols

27.1 Objective

Resolve and invoke minimal driver discovery API.

Load:

cuInit
cuDeviceGetCount
cuDeviceGet
cuDeviceGetName
cuGetErrorName
cuGetErrorString

Add function-pointer typedefs.

Example:

using PFN_cuInit =
    CUresult (*)(unsigned int);

using PFN_cuDeviceGetCount =
    CUresult (*)(int*);

using PFN_cuDeviceGet =
    CUresult (*)(CUdevice*, int);

using PFN_cuDeviceGetName =
    CUresult (*)(char*, int, CUdevice);

Initialization algorithm:

TryLoad library
 |
resolve required symbols
 |
cuInit(0)
 |
cuDeviceGetCount
 |
if count <= 0 -> DeviceUnavailable
 |
cuDeviceGet(device ordinal 0)
 |
cuDeviceGetName

Add QCall:

Accelerator_GetDeviceCount

Expected smoke output:

ACCELERATOR_SUPPORTED=True
ACCELERATOR_DEVICE_COUNT=<number >= 1>

Optional diagnostic print under environment variable only:

DOTNET_AcceleratorVerbose=1

may print GPU name.

Never unconditionally printf from runtime code.

27.2 Done

[ ] cuInit(0) called once
[ ] device count visible to managed test
[ ] ordinal 0 obtained
[ ] GPU name can be inspected in verbose mode

Commit:

accelerator: initialize CUDA driver

28 27. A012 — Retain Primary CUDA Context

28.1 Objective

Establish context required for memory/module/launch operations.

Resolve:

cuDevicePrimaryCtxRetain
cuDevicePrimaryCtxRelease
cuCtxSetCurrent
cuCtxSynchronize

CudaDriver gains:

CUdevice _device;
CUcontext _context;
bool _contextRetained;

Algorithm:

EnsureInitialized
 |
if contextRetained -> cuCtxSetCurrent(context)
 |
else:
    cuDevicePrimaryCtxRetain
    cuCtxSetCurrent
    contextRetained=true

Add:

EnsureContext()

Every operation requiring CUDA context begins by calling it.

Add Accelerator_Shutdown.

Shutdown:

if contextRetained:
    cuDevicePrimaryCtxRelease(device)
    context=null
    contextRetained=false

Do not unload libcuda during MVP shutdown.

28.2 Test

Call through managed reflection:

IsSupported
GetDeviceCount
Shutdown

twice in separate processes.

No crash.

28.3 Done

[ ] primary context retained
[ ] current context established
[ ] explicit shutdown works
[ ] second process still works

Commit:

accelerator: establish CUDA primary context

29 28. A013 — Implement Device Allocate and Free

Resolve:

cuMemAlloc_v2 or compatible symbol
cuMemFree_v2 or compatible symbol

QCalls:

Accelerator_Allocate
Accelerator_Free

Native signature concept:

extern "C"
INT32 QCALLTYPE Accelerator_Allocate(
    UINT64 byteCount,
    UINT64* result);

extern "C"
INT32 QCALLTYPE Accelerator_Free(
    UINT64 devicePointer);

Rules:

byteCount == 0 -> InvalidArgument
result == null -> InvalidArgument
CUDA allocation failure -> OutOfMemory or ContextFailure as appropriate
success -> result contains CUdeviceptr

Managed wrapper converts nonzero status to InvalidOperationException.

Smoke test:

ulong p = RuntimeBridge.Allocate(4);
Console.WriteLine(p != 0 ? "ALLOC_OK" : "ALLOC_BAD");
RuntimeBridge.Free(p);

Expected:

ALLOC_OK

Run test 100 times from shell.

No crash.

29.1 Done

[ ] 4-byte allocation works
[ ] pointer nonzero
[ ] free works
[ ] zero allocation rejected
[ ] repeated process execution works

Commit:

accelerator: implement device allocation

30 29. A014 — Implement Host/Device Copies

Resolve:

cuMemcpyHtoD_v2
cuMemcpyDtoH_v2

Native QCalls accept:

device pointer
host pointer as IntPtr
byte count

Managed CoreLib wrapper pins byte[].

Exact managed algorithm:

internal static unsafe void CopyToDevice(
    ulong destination,
    byte[] source)
{
    ArgumentNullException.ThrowIfNull(source);

    fixed (byte* p = source)
    {
        Check(
            CopyToDeviceNative(
                destination,
                (IntPtr)p,
                (nuint)source.Length));
    }
}

Reverse similarly.

Test:

byte[] expected =
{
    0x12, 0x34, 0x56, 0x78
};

byte[] actual = new byte[4];

ulong p = Allocate(4);

CopyToDevice(p, expected);
CopyFromDevice(actual, p);

Free(p);

if (!expected.AsSpan().SequenceEqual(actual))
    return failure;

Expected:

ROUNDTRIP_OK

Test sizes:

30.1 Done

[ ] H->D works
[ ] D->H works
[ ] all five sizes round-trip
[ ] pinned pointer lifetime cannot escape native call

Commit:

accelerator: implement synchronous memory copies

31 30. A015 — Launch Handwritten PTX set42

31.1 Objective

Prove the complete lower runtime half before building a compiler.

Add symbol resolution:

cuModuleLoadData
cuModuleUnload
cuModuleGetFunction
cuLaunchKernel

Add:

LoadPtxKernel
DestroyKernel
LaunchSet42

Use this PTX fixture:

.version 7.0
.target sm_50
.address_size 64

.visible .entry set42(
    .param .u64 output_ptr
)
{
    .reg .b64 %rd<2>;
    .reg .b32 %r<2>;

    ld.param.u64 %rd1, [output_ptr];
    mov.u32 %r1, 42;
    st.global.u32 [%rd1], %r1;

    ret;
}

Managed test:

allocate 4 bytes
load PTX
get function set42
launch
synchronize
copy 4 bytes back
BitConverter.ToInt32 == 42
destroy kernel
free allocation

Expected:

SET42_RESULT=42
SET42_OK

This is a mandatory gate.

DO NOT begin IR implementation until this passes.

31.2 Failure isolation

ModuleLoadFailed:

dump exact PTX bytes.

Check:

.version
.target
.address_size
null termination

FunctionLookupFailed:

confirm exact kernel name:

set42

KernelLaunchFailed:

inspect:

parameter pointer
grid 1
block 1
context current

Wrong result:

run generated PTX independently if desired before touching compiler code.

31.3 Done

[ ] CoreCLR invokes driver
[ ] PTX module loads
[ ] kernel function resolves
[ ] GPU writes integer 42
[ ] host reads 42

Commit:

accelerator: execute handwritten PTX kernel

32 31. A016 — Handwritten PTX Vector Add

32.1 Objective

Validate the exact runtime ABI later used by generated code.

Kernel ABI:

u64 a
u64 b
u64 output
u32 length

Use PTX implementing:

i = ctaid.x * ntid.x + tid.x

if i >= length:
    return

output[i] = a[i] + b[i]

Launch:

block=256
grid=(N+255)/256

Test N:

Input generation:

a[i] = i * 0.25f;
b[i] = i * -0.5f;

Expected CPU:

expected[i] = a[i] + b[i];

Require bitwise equality for simple float addition where generated operation ordering is identical.

Gate output:

HANDWRITTEN_ADD_OK

DO NOT start compiler backend until green.

Commit:

accelerator: validate vector add kernel ABI

33 32. A020 — Implement Accelerator IR Model

Create:

acceleratorir.h
acceleratorir.cpp

Required structs:

enum class IRType
{
    Void,
    Bool,
    I32,
    I64,
    F32,
    DevicePtr
};

enum class IROpcode
{
    Parameter,
    ConstantI32,

    GlobalIndexX,

    LoadF32,
    StoreF32,

    AddF32,

    CompareGeI32,

    Branch,
    BranchIf,

    Return
};

struct IRValue
{
    ValueId id;
    IRType type;
};

struct IRInstruction
{
    IROpcode opcode;

    ValueId result;

    ValueId operand0;
    ValueId operand1;
    ValueId operand2;

    int32_t immediate;

    BlockId target0;
    BlockId target1;
};

struct IRBlock
{
    BlockId id;
    std::vector<IRInstruction> instructions;
};

struct IRFunction
{
    std::vector<IRBlock> blocks;
};

You may adapt container types to CoreCLR coding conventions, but semantic fields must remain equivalent.

Create helper builder methods rather than manually populating instruction fields in callers.

Examples:

ValueId AddParameter(IRType type, int index);

ValueId AddGlobalIndexX();

ValueId AddLoadF32(
    ValueId base,
    ValueId index);

ValueId AddAddF32(
    ValueId left,
    ValueId right);

void AddStoreF32(
    ValueId base,
    ValueId index,
    ValueId value);

Test only hand-built IR.

No compiler.

No PTX.

33.1 Done

[ ] AddF32 sample can be represented
[ ] all ValueIds deterministic
[ ] blocks deterministic

Commit:

accelerator: introduce accelerator IR

34 33. A021 — Implement IR Verifier

Implement the exact verifier contract from section 15.

Create unit-style native tests if convenient, or expose a temporary test QCall.

Preferred test cases:

valid AddF32 -> Ok

missing terminator -> InvalidIR

unknown ValueId -> InvalidIR

AddF32(I32,I32) -> InvalidIR

LoadF32(I32,I32) -> InvalidIR

Branch to missing block -> InvalidIR

Every IR passed to backend MUST first verify successfully.

Do not allow backend to “do its best” with malformed IR.

Commit:

accelerator: verify accelerator IR

35 34. A022 — Implement Deterministic IR Printer

Create:

acceleratorirprinter.h
acceleratorirprinter.cpp

Print exact semantic structure.

Add environment-controlled dump:

DOTNET_AcceleratorDumpIR=1

Do NOT integrate with CLRConfig yet.

Reading process environment directly inside experimental accelerator code is acceptable for MVP Linux.

Expected example is the format in section 16.

Golden comparison should normalize final newline only.

Commit:

accelerator: add deterministic IR diagnostics

36 35. A023 — Implement IR Interpreter

36.1 Objective

Create CPU semantic oracle for IR independently of PTX.

Interpreter invocation accepts:

IRFunction
parameter values
globalIndexX
host-side backing buffers

For MVP interpreter tests, DevicePtr may map to a small logical buffer ID rather than literal host pointer.

Required operations:

Parameter
ConstantI32
GlobalIndexX
LoadF32
StoreF32
AddF32
CompareGeI32
Branch
BranchIf
Return

Run AddF32 IR once for each logical global index:

for i in [0,N):
    interpreter.Execute(function, globalIndexX=i)

Compare CPU array.

Test sizes:

Expected:

IR_INTERPRETER_ADD_OK

This becomes the canonical semantic reference.

Commit:

accelerator: add accelerator IR interpreter

37 36. A024 — Emit PTX From Hand-Built IR

Create:

ptxcodegen.h
ptxcodegen.cpp

Input:

IRFunction

Output:

std::string PTX

There is no optimization.

Bool      -> .pred
I32       -> .b32/.s32
I64       -> .b64
DevicePtr -> .b64
F32       -> .f32

Allocate result registers in ValueId order.

Generate labels:

BB0
BB1
BB2

from BlockId.

Build the exact AddF32 IR program by hand.

Generate PTX.

Load generated PTX through already-working runtime backend.

Launch.

Expected:

HAND_BUILT_IR_ADD_GPU_OK

This is mandatory gate #2.

At this point:

IR -> PTX -> GPU

works without IL.

Commit:

accelerator: lower accelerator IR to PTX

38 37. A030 — Pass RuntimeMethodHandle Into CoreCLR

Add:

CompileKernel(
    RuntimeMethodHandle method,
    int globalIndexToken,
    int loadF32Token,
    int storeF32Token)

Managed wrapper passes:

method.Value

to native as IntPtr.

Native receives:

INT_PTR methodDescValue

and converts:

MethodDesc* method =
    reinterpret_cast<MethodDesc*>(
        methodDescValue);

Validation:

method != nullptr
method represents IL method
method has IL header

For this ticket do NOT compile.

Only expose diagnostic information such as:

IL byte count

through verbose logging.

Kernel test obtains:

MethodInfo method =
    typeof(Kernels).GetMethod(
        nameof(Kernels.AddF32),
        BindingFlags.Static |
        BindingFlags.NonPublic)!;

Expected verbose output:

Accelerator method: Kernels.AddF32
IL size: <positive integer>

38.1 Failure

If MethodDesc conversion crashes:

stop.

Do not guess.

Compare the representation to existing RuntimeMethodHandle CoreCLR code in:

src/coreclr/System.Private.CoreLib/src/System/RuntimeHandles.cs
src/coreclr/vm/runtimehandles.cpp
src/coreclr/vm/runtimehandles.h

The ticket is not complete until MethodDesc validation is safe.

Commit:

accelerator: accept managed method handles

39 38. A031 — Extract IL Method Body

Use CoreCLR’s existing IL method representation.

Relevant runtime facilities:

MethodDesc::GetILHeader
COR_ILMETHOD
COR_ILMETHOD_DECODER

Importer input structure:

struct MethodIL
{
    const uint8_t* code;
    uint32_t codeSize;
};

MVP rejects:

no IL body
dynamic method
runtime-generated method
method with EH sections

Exception handling is unsupported.

Do NOT parse EH.

Return UnsupportedMethod.

Diagnostic:

DOTNET_AcceleratorDumpIL=1

prints:

IL_0000: xx
IL_0001: xx
...

For this ticket bytes are sufficient; mnemonics come next.

Expected:

code != nullptr
codeSize > 0

Commit:

accelerator: extract kernel IL bodies

40 39. A032 — Decode MVP IL Opcodes

Create:

ildecoder.h
ildecoder.cpp

Decoder produces:

struct DecodedILInstruction
{
    uint32_t offset;
    uint16_t opcode;

    int32_t intOperand;
    uint32_t tokenOperand;
    int32_t branchTarget;
};

Support only required kernel opcodes.

Expected likely set:

nop

ldarg.0
ldarg.1
ldarg.2
ldarg.3

ldloc.0
ldloc.1
ldloc.2
ldloc.3

stloc.0
stloc.1
stloc.2
stloc.3

ldc.i4.*
ldc.i4.s
ldc.i4

call

add

bge
bge.s

br
br.s

ret

Before implementing, dump AddF32 IL with existing tooling.

If compiler produced one additional benign opcode required by this exact kernel, add it deliberately to this list and document it in the commit.

Do NOT implement all ECMA-335 opcodes.

Unknown opcode:

UnsupportedIL

Diagnostic must include:

method
IL offset
numeric opcode

Commit:

accelerator: decode minimal kernel IL subset

41 40. A033 — Construct Basic Blocks

Algorithm:

Decode all instructions.
Start block set with offset 0.
For each branch:
- add branch target as block start;
- add instruction after conditional branch as block start if present.
Add instruction after unconditional branch only if reachable later.
Add instruction after ret only if independently targeted.
Sort unique block starts.
Assign BlockId in ascending IL offset order.
Assign every decoded instruction to corresponding block.

Keep mapping:

std::map<uint32_t, BlockId> ilOffsetToBlock;

Validation:

branch target must point to instruction boundary;
target must be inside method body.

Invalid branch -> InvalidIL.

Golden debug dump:

B0 IL_0000
B1 IL_0012
B2 IL_0028

Do not import semantics yet.

Commit:

accelerator: construct kernel IL CFG

42 41. A034 — Implement Evaluation Stack and Locals

Importer state:

struct ImportState
{
    std::vector<ValueId> stack;

    ValueId locals[4];

    bool localAssigned[4];
};

Only four locals supported initially.

Evaluation stack helpers:

void Push(ValueId);
ValueId Pop();
ValueId Peek();

Underflow -> InvalidIL.

Local rules:

stloc.N:
    pop value
    locals[N] = value
    assigned[N] = true

ldloc.N:
    if !assigned -> UnsupportedIL
    push locals[N]

Argument types are hardcoded by kernel ABI:

arg0 DevicePtr
arg1 DevicePtr
arg2 DevicePtr
arg3 I32

Create parameter IR instructions at entry.

ldarg.N pushes corresponding parameter ValueId.

MVP branch rule:

stack MUST be empty at every control-flow edge.

If nonempty:

UnsupportedIL:
stack merge not supported by MVP importer

At merge blocks:

incoming local state must match exactly.

If local ValueIds differ between incoming edges:

UnsupportedIL:
local phi required

Do not implement phi.

Commit:

accelerator: import kernel stack and locals

43 42. A035 — Recognize Accelerator Intrinsics

Managed test obtains:

int globalIndexToken =
    typeof(AcceleratorIntrinsics)
        .GetMethod(nameof(
            AcceleratorIntrinsics.GlobalIndexX))!
        .MetadataToken;

int loadF32Token =
    typeof(AcceleratorIntrinsics)
        .GetMethod(nameof(
            AcceleratorIntrinsics.LoadF32))!
        .MetadataToken;

int storeF32Token =
    typeof(AcceleratorIntrinsics)
        .GetMethod(nameof(
            AcceleratorIntrinsics.StoreF32))!
        .MetadataToken;

These three tokens are passed to CompileKernel.

Importer behavior on call token:

GlobalIndexX:

pop 0
emit GlobalIndexX -> I32
push result

LoadF32:

IL argument order means stack has:

baseAddress
index

Importer:

index = Pop()
base = Pop()
emit LoadF32(base,index)
push F32 result

StoreF32:

value = Pop()
index = Pop()
base = Pop()
emit StoreF32(base,index,value)
push nothing

Any other call token:

UnsupportedIL

Do NOT inspect method names.

Do NOT perform metadata call resolution.

Commit:

accelerator: recognize MVP kernel intrinsics

44 43. A036 — Import Complete AddF32 Kernel

Implement semantic lowering for remaining IL.

add:

right = Pop
left = Pop

if left/right are F32:
    result = AddF32
else:
    UnsupportedIL

bge/bge.s:

right = Pop
left = Pop

condition = CompareGeI32(left,right)

BranchIf(condition, target, fallthrough)

br:

Branch(target)

ret:

MVP kernel return type is void.

Evaluation stack must be empty.

Emit Return.

After import:

run IR verifier

Dump IR.

Expected IR structurally equivalent to section 16.

Then execute the IR interpreter.

Expected:

IL_TO_IR_INTERPRETER_ADD_OK

This is mandatory gate #3.

Commit:

accelerator: import AddF32 C# kernel

45 44. A040 — Compile Imported IL to PTX

CompileKernel now performs:

MethodDesc
 |
GetILHeader
 |
Decode
 |
Build CFG
 |
Import
 |
Verify IR
 |
PTX codegen
 |
cuModuleLoadData
 |
cuModuleGetFunction
 |
return AcceleratorKernelHandle

Kernel PTX name is fixed for MVP:

managed_kernel

Do not derive a mangled name from MethodDesc yet.

PTX emitter should accept entrypoint name parameter.

Environment variable:

DOTNET_AcceleratorDumpPTX=1

prints complete generated PTX to stderr.

No disk cache.

No memory cache.

Every CompileKernel call recompiles.

Commit:

accelerator: compile managed kernels to PTX

46 45. A041 — Launch Compiled AddF32

Implement:

LaunchAddF32

Host algorithm:

N
 |
create host float arrays
 |
convert to byte[] / pin
 |
allocate GPU buffers
 |
copy inputs
 |
CompileKernel(Kernels.AddF32)
 |
LaunchAddF32
 |
copy output
 |
compare CPU
 |
DestroyKernel
 |
Free buffers

cuLaunchKernel parameter storage must be host variables:

CUdeviceptr aArg = ...;
CUdeviceptr bArg = ...;
CUdeviceptr outputArg = ...;
int32_t lengthArg = ...;

void* parameters[] =
{
    &aArg,
    &bArg,
    &outputArg,
    &lengthArg
};

Do NOT put device pointers themselves directly into void* parameters.

Kernel launch uses address of argument storage.

Expected:

MANAGED_IL_GPU_ADD_OK

This is the first primary demo.

Commit:

accelerator: execute C# IL kernel on NVIDIA GPU

47 46. A042 — Differential Correctness Matrix

Create test loop for N:

Data families:

ascending
descending
alternating signs
random deterministic seed 12345
zero
subnormal/special-value set

Special values include:

0
-0
float.NaN
float.PositiveInfinity
float.NegativeInfinity
float.MaxValue
float.MinValue
float.Epsilon

For addition comparison:

if expected is NaN:

actual must be NaN

otherwise compare bit representation where semantics permit.

Log failing:

N
index
a bits
b bits
expected bits
actual bits

Never log only decimal floats.

Expected:

DIFFERENTIAL_ADD_PASS=<number>
DIFFERENTIAL_ADD_FAIL=0

Commit:

tests: add accelerator differential correctness matrix

48 47. A043 — Negative Compiler Tests

Each unsupported construct must fail deterministically.

Add kernels for:

new object
throw
array access
unknown method call
loop requiring local merge/phi
unsupported return value
more than four parameters
unsupported parameter type
try/catch

Expected:

AcceleratorStatus::UnsupportedIL
or
AcceleratorStatus::UnsupportedMethod

No crash.

No assertion in Checked build for user-invalid kernel.

Diagnostic contains:

method name if available
IL offset where applicable
reason

Example:

ACCELERATOR COMPILATION FAILED
method: BadKernels.Allocate
IL: IL_0001
reason: unsupported opcode newobj

Commit:

tests: validate accelerator compiler rejection paths

49 48. A044 — Developer Diagnostics and Debug Recipes

Add:

DOTNET_AcceleratorVerbose
DOTNET_AcceleratorDumpIL
DOTNET_AcceleratorDumpIR
DOTNET_AcceleratorDumpPTX

All disabled by default.

49.1 Debug recipe: native runtime crash

lldb -- \
  "$CORE_ROOT/corerun" \
  <test.dll>

Recommended breakpoints:

Accelerator_IsSupported
CudaDriver::TryLoad
CudaDriver::EnsureInitialized
CudaDriver::EnsureContext
Accelerator_CompileKernel
AcceleratorImporter::Import
PTXCodegen::Generate
Accelerator_LaunchAddF32

49.2 Failure classification

Crash before Accelerator_CompileKernel
    -> managed/QCall/runtime plumbing

Importer error
    -> IL/importer

IR interpreter wrong
    -> importer/IR semantics

IR interpreter correct
PTX GPU wrong
    -> PTX backend

Handwritten PTX wrong too
    -> CUDA runtime/launch

Handwritten PTX works
generated PTX fails
    -> codegen

Generated PTX standalone works
CoreCLR launch fails
    -> marshalling/launch integration

49.3 MVP-A completion gate

At this point all must be true:

[ ] Core_Root reproducible
[ ] CUDA driver loaded dynamically
[ ] primary context works
[ ] GPU memory round-trip works
[ ] handwritten set42 works
[ ] handwritten vector add works
[ ] IR interpreter vector add works
[ ] hand-built IR -> PTX -> GPU works
[ ] managed MethodDesc IL extraction works
[ ] C# AddF32 -> IR works
[ ] C# AddF32 -> PTX -> GPU works
[ ] differential matrix passes
[ ] rejection tests pass
[ ] all intermediates dumpable

This is the definition of MVP-A.

50 49. A050 — Add Subgroup IR Operations

Do not begin before MVP-A is green.

Extend IR types/opcodes with:

LaneId
ActiveMask
ShuffleDownF32

Semantics:

LaneId:

0..31 for NVIDIA warp

ShuffleDownF32(value, delta):

read value from lane + delta when valid

No generic shuffle API.

No scan.

No ballot API yet unless required for active lanes.

Add interpreter subgroup model with exactly 32 logical lanes.

Test:

input lane value = lane id

shuffle_down delta=1

lane 0 receives 1
lane 1 receives 2
...

Boundary lane semantics must be explicit.

Commit:

accelerator: introduce subgroup shuffle IR

51 50. A051 — PTX Warp Shuffle Backend

Lower ShuffleDownF32 using supported PTX warp shuffle operation.

Test 32 lanes.

Then test partial warp:

N=17

The active-lane semantics must not read garbage from inactive lanes.

If active masking requires an additional IR operation, add:

ActiveMask

in the same ticket only if specified by backend semantics.

Required tests:

N=1
N=2
N=17
N=31
N=32
N=33

Commit:

accelerator: lower subgroup shuffles to PTX

52 51. A052 — Reduction Proof of Concept

This is the VectorWare-motivated validation milestone.

Implement a compiler-known reduction helper rather than general automatic reduction recognition.

Logical operation:

each lane owns one float

for offset:
    16
    8
    4
    2
    1

value += shuffle_down(value, offset)

Lane 0 holds warp sum.

Start with exactly 32 elements.

Then add partial warp mask behavior.

Do NOT implement multi-block reduction first.

Test:

values = 1..32
expected = 528

Then randomized vectors.

Compare CPU sum with documented FP tolerance.

MVP-B success statement:

The same accelerator IR now expresses both ordinary elementwise work and a warp-level collective operation, and the NVIDIA backend lowers that collective operation to GPU lane communication.

This is the first concrete validation of the SIMD/subgroup thesis.

Commit:

accelerator: prototype warp reduction

53 52. A060 — Performance Baseline

Only use Release runtime for performance.

Build:

./build.sh -subset clr -configuration Release

Do not report Checked measurements.

Measure separately:

CoreCLR startup
CUDA initialization
device allocation
H->D
CompileKernel
PTX module JIT/load
kernel launch
synchronize
D->H
warm launch only
total cold
total warm

Workloads:

AddF32

N:
1K
10K
100K
1M
10M
100M if memory permits

CPU references:

plain scalar loop
TensorPrimitives.Add where practical

GPU references:

handwritten PTX AddF32
generated PTX AddF32

The generated kernel should be structurally compared to handwritten PTX.

Performance MVP does NOT require GPU to win at every N.

Questions to answer:

What is fixed launch overhead?

What is initial PTX JIT overhead?

At what N does resident-data GPU execution beat CPU?

How much slower is generated PTX than handwritten PTX?

Is generated memory access coalesced?

What percentage of total time is transfer?

54 53. Fast Native Development Loop

After managed CoreLib bridge stabilizes, most changes are native.

Run incremental:

./build.sh -subset clr -configuration Checked

Then update the active Core_Root.

Create:

eng/accelerator/update-coreroot.sh

Initial script:

#!/usr/bin/env bash
set -euo pipefail

ROOT="$(git rev-parse --show-toplevel)"

SRC="$ROOT/artifacts/bin/coreclr/linux.x64.Checked"
DST="$ROOT/artifacts/tests/coreclr/linux.x64.Checked/Tests/Core_Root"

cp "$SRC/libcoreclr.so" "$DST/libcoreclr.so"

If JIT is not modified, do NOT copy libclrjit.so.

If System.Private.CoreLib changes, regenerate layout instead of using this script.

55 54. Mandatory Testing Rings

55.1 Ring 0 — unmodified baseline

Build/run before accelerator source change.

55.2 Ring 1 — runtime bridge

QCall works.

55.3 Ring 2 — CUDA driver

Driver/device/context.

55.4 Ring 3 — memory

allocate/copy/free.

55.5 Ring 4 — handwritten PTX

set42 and AddF32.

55.6 Ring 5 — IR interpreter

no GPU compiler dependency.

55.7 Ring 6 — IR -> PTX

hand-built IR.

55.8 Ring 7 — IL importer

IL -> IR interpreter.

55.9 Ring 8 — full GPU

IL -> IR -> PTX -> GPU.

55.10 Ring 9 — differential fuzz/matrix

CPU/GPU agreement.

55.11 Ring 10 — existing CoreCLR tests

No regressions to unrelated runtime behavior.

Do not skip directly from Ring 2 to Ring 8.

56 55. Regression Testing

Accelerator functionality must be dormant when unused.

Run relevant standard CoreCLR tests after MVP integration.

At minimum:

./src/tests/build.sh \
    -arch x64 \
    -checked \
    -dir:JIT

A broader runtime suite should be run before declaring fork stable.

Because MVP does not modify RyuJIT, normal generated machine code should remain unchanged.

If later JIT changes are introduced, add:

DOTNET_JitDump
DOTNET_JitDisasm
SuperPMI

validation at that time.

Do NOT introduce SuperPMI complexity into the first non-JIT MVP unnecessarily.

57 56. Error Reporting Standard

Never return “false” for compiler failure without reason.

Internal error object:

struct AcceleratorError
{
    AcceleratorStatus status;
    uint32_t ilOffset;
    const char* message;
};

For MVP this may be transient/thread-local.

Required message examples:

CUDA driver unavailable

CUDA device 0 unavailable

unsupported IL opcode 0x73 at IL_0012

evaluation stack underflow at IL_0009

branch target 43 is not an instruction boundary

kernel requires phi node at IL_0020

unsupported call token 0x06000007

invalid IR: AddF32 expected F32,F32

PTX module rejected by CUDA driver

Junior developers must be able to diagnose test failures without stepping into a debugger for every malformed kernel.

58 57. No-Crash Rule

Input problems must never crash CoreCLR.

The following must be rejected:

null MethodDesc
method with no IL
invalid opcode
bad stack
unsupported call
unsupported branch merge
invalid IR
invalid kernel handle
invalid device pointer where detectable
bad PTX
CUDA launch failure

Checked assertions are for internal impossible states.

They are NOT input validation.

59 58. Memory Safety Checklist

Every native pointer must have documented ownership.

59.1 libcuda handle

Owner:

CudaDriver singleton

59.2 primary context

Owner:

CudaDriver singleton retains one reference

59.3 CUdeviceptr

Owner:

managed caller after Allocate

Released:

Free

59.4 AcceleratorKernelHandle*

Owner:

managed caller after LoadPtxKernel/CompileKernel

Released:

DestroyKernel

59.5 PTX string

Owner:

temporary compiler/native std::string

Driver module load must complete before string storage disappears.

59.6 byte[] pin

Owner:

managed wrapper fixed block

Native call must complete synchronously before unpin.

No native code stores host pointers after QCall return.

60 59. First Demo Script

The project’s first real demo must print each stage.

Example:

.NET Accelerator MVP

Runtime:
  CoreCLR fork: yes

CUDA:
  supported: yes
  device count: 1
  device: NVIDIA ...

Kernel:
  managed method: Kernels.AddF32
  IL bytes: 37

Accelerator IR:
  verified

PTX:
  generated
  module loaded

Execution:
  elements: 1000000
  block size: 256
  grid size: 3907

Verification:
  mismatches: 0

RESULT: PASS

With:

DOTNET_AcceleratorDumpIL=1 \
DOTNET_AcceleratorDumpIR=1 \
DOTNET_AcceleratorDumpPTX=1 \
<test launcher>

the demo additionally prints every intermediate representation.

61 60. Explicit MVP-A Non-Goals

The following are NOT permitted to delay MVP-A:

System.Numerics.Vector<T>
Vector128<T>
Vector256<T>
Vector512<T>

TensorPrimitives

Span<T>
ReadOnlySpan<T>

arrays in kernels

managed objects

GC integration

exceptions in kernels

virtual calls

interfaces

delegates

generic math

generic kernels

helper method calls

kernel inlining

automatic kernel discovery

attributes

automatic CPU/GPU placement

memory migration

unified memory

async launch

CUDA streams

multiple GPUs

AMD

Vulkan

SPIR-V

Windows

NativeAOT

RyuJIT integration

automatic vectorization

LINQ

Parallel.For

kernel caching

disk cache

PGO

tensor cores

Any pull request attempting these before MVP-A should be deferred.

62 61. Post-MVP Expansion Order

Only after MVP-A:

1. subgroup/shuffle IR
2. reduction
3. proper phi nodes / SSA
4. arbitrary scalar arithmetic
5. pointer/byref lowering
6. unmanaged structs
7. helper-method calls
8. generic instantiations
9. Span-like device views
10. kernel cache
11. device-resident Tensor abstraction
12. TensorPrimitives backend
13. RyuJIT/shared IR investigation
14. automatic loop extraction
15. automatic placement cost model

Do not jump directly from vector-add MVP to transparent TensorPrimitives offload.

63 62. Decision Gate: RyuJIT Integration

Do not modify RyuJIT until the following data exists:

[ ] at least three C# kernels compile
[ ] interpreter validation is stable
[ ] PTX backend is stable
[ ] subgroup reduction works
[ ] IL importer pain points are documented
[ ] generated PTX quality measured

Then evaluate:

63.1 Option A

IL
 +--> RyuJIT -> CPU
 |
 +--> AcceleratorCompiler -> GPU

63.2 Option B

IL
 |
RyuJIT importer
 |
shared IR
 +--> CPU backend
 +--> GPU backend

63.3 Option C

IL
 |
analysis
 |
data-parallel IR
 +--> CPU SIMD
 +--> GPU subgroup

No option is selected in MVP.

The MVP produces evidence required to select it.

64 63. Decision Gate: TensorPrimitives

TensorPrimitives integration is allowed only after:

[ ] explicit device memory abstraction exists
[ ] kernel cache exists
[ ] Add/Multiply/Reduce are stable
[ ] transfer cost measured
[ ] resident GPU buffer benchmark exists

The first TensorPrimitives experiment should NOT transparently move Span<T> data to the GPU.

It should operate on explicitly device-resident memory.

65 64. Definition of MVP-A Done

MVP-A is complete only if this source method:

internal static void AddF32(
    ulong a,
    ulong b,
    ulong output,
    int length)
{
    int i = AcceleratorIntrinsics.GlobalIndexX();

    if (i >= length)
        return;

    float value =
        AcceleratorIntrinsics.LoadF32(a, i) +
        AcceleratorIntrinsics.LoadF32(b, i);

    AcceleratorIntrinsics.StoreF32(
        output,
        i,
        value);
}

can:

be compiled by the ordinary C# compiler;
remain ordinary IL in the test assembly;
be located through RuntimeMethodHandle;
be represented by a CoreCLR MethodDesc;
have its IL read by the accelerator compiler;
be decoded by our importer;
become valid Accelerator IR;
execute correctly in our IR interpreter;
generate PTX;
load through the NVIDIA driver;
execute on device-resident buffers;
match CPU results over the differential test matrix.

No source-generated kernel is permitted.

No handwritten PTX equivalent is used for the final managed-kernel execution.

No CUDA C++ is permitted.

66 65. Definition of MVP-B Done

MVP-B is complete when, in addition to MVP-A:

[ ] subgroup lane operation represented in IR
[ ] PTX shuffle lowering exists
[ ] 32-lane reduction works
[ ] partial warp behavior tested
[ ] CPU interpreter provides subgroup oracle

This validates the central hypothesis:

A portable data-parallel compiler representation can map collective vector-like operations onto GPU warp/subgroup operations rather than requiring the source language to expose CUDA directly.

67 66. Junior Developer Escalation Rules

The implementer MUST stop and escalate when:

an existing CoreCLR API contradicts an interface assumed here;
a MethodDesc is not safely obtainable from RuntimeMethodHandle.Value;
current C# IL for the canonical AddF32 kernel requires an opcode not listed here and the semantics are nontrivial;
a CFG merge requires phi nodes in the canonical kernel;
the CUDA driver ABI available on the baseline machine cannot load the specified entry points;
generated PTX semantics differ from the IR interpreter for unclear reasons;
fixing a problem seems to require modifying RyuJIT;
fixing a problem seems to require GC knowledge;
a change requires exposing a public CoreLib API;
a test failure can only be fixed by disabling an assertion.

Escalation is success.

Inventing undocumented architecture inside a ticket is failure.

68 67. Reviewer Checklist Per Commit

Every implementation commit must answer:

Does this commit implement only one ticket?

Does it preserve previous tests?

Is ownership explicit?

Can failure be reproduced?

Is there a deterministic test?

Does unsupported input fail safely?

Did the change add unconditional logging?

Did the change accidentally affect normal JIT behavior?

Can this component be tested without later components?

69 68. Expected Commit History

The first successful history should approximately look like:

accelerator: record MVP baseline

tests: add accelerator smoke test

coreclr: add accelerator QCall bridge

accelerator: dynamically load CUDA driver
accelerator: initialize CUDA driver
accelerator: establish CUDA primary context
accelerator: implement device allocation
accelerator: implement synchronous memory copies

accelerator: execute handwritten PTX kernel
accelerator: validate vector add kernel ABI

accelerator: introduce accelerator IR
accelerator: verify accelerator IR
accelerator: add deterministic IR diagnostics
accelerator: add accelerator IR interpreter
accelerator: lower accelerator IR to PTX

accelerator: accept managed method handles
accelerator: extract kernel IL bodies
accelerator: decode minimal kernel IL subset
accelerator: construct kernel IL CFG
accelerator: import kernel stack and locals
accelerator: recognize MVP kernel intrinsics
accelerator: import AddF32 C# kernel

accelerator: compile managed kernels to PTX
accelerator: execute C# IL kernel on NVIDIA GPU

tests: add accelerator differential correctness matrix
tests: validate accelerator compiler rejection paths

accelerator: add developer diagnostics

If the history instead contains one commit named:

implement GPU support

the project has lost its debugging strategy.

70 69. The Four Hard Gates

Do not proceed past a gate while it is red.

70.1 Gate 1

CoreCLR -> handwritten PTX -> GPU -> 42

Proves runtime/CUDA layer.

70.2 Gate 2

hand-built IR -> generated PTX -> GPU -> correct AddF32

Proves IR/backend layer.

70.3 Gate 3

C# IL -> IR interpreter -> correct AddF32

Proves IL/importer layer.

70.4 Gate 4

C# IL -> IR -> generated PTX -> GPU -> correct AddF32

Proves complete MVP.

Every bug after Gate 4 can be bisected across already-known-good boundaries.

71 70. Why This Version Is Junior-Executable

The implementer is NOT being asked:

"design GPU support for .NET"

They are being asked:

A003:
wire this QCall.

A010:
load this shared library.

A011:
resolve these symbols.

A013:
allocate these bytes.

A015:
execute this exact PTX.

A020:
implement these exact IR types.

A024:
lower these exact IR operations.

A031:
extract these IL bytes.

A032:
decode these exact opcodes.

A035:
map these exact three tokens.

A041:
pass these exact four kernel arguments.

A042:
compare these exact test cases.

The difficult architectural decisions have already been made.

The workflow intentionally constructs four independently testable systems:

1. CoreCLR <-> CUDA runtime plumbing

2. Accelerator IR interpreter

3. Accelerator IR -> PTX backend

4. IL -> Accelerator IR importer

Only after each subsystem independently works are they connected.

That is the core implementation strategy.

72 71. Final MVP Principle

The first milestone is not:

Make .NET magically use the GPU.

The first milestone is:

Make one ordinary C# method become inspectable IL, inspectable accelerator IR, inspectable PTX, and a correct GPU result while running inside a private forked CoreCLR.

Once that pipeline is reliable, everything interesting becomes incremental:

more IL
more types
more intrinsics
subgroups
reductions
vectors
generic math
device tensors
TensorPrimitives
automatic extraction
automatic placement

But until:

IL
 |
IR
 |
PTX
 |
GPU

is boring, deterministic, inspectable, and testable, none of the higher-level abstractions should be attempted.

That boring pipeline is the MVP. “`

lqdev/spec.org

.NET Accelerator MVP

1 0. Document Contract

2 1. Definition of “Shovel Ready”

3 2. Locked Architectural Decisions

3.1 2.1 Host operating system

3.2 2.2 GPU backend

3.3 2.3 CoreCLR fork

3.4 2.4 No automatic kernel discovery

3.5 2.5 No public BCL API

3.6 2.6 No Span<T> in MVP

3.7 2.7 Exactly three compiler intrinsics initially

3.8 2.8 No SSA requirement for MVP

3.9 2.9 No optimizer

3.10 2.10 Explicit GPU residency

3.11 2.11 One process-global CUDA device

3.12 2.12 Synchronous execution

4 3. Repository Baseline

4.1 3.1 Fork

4.2 3.2 Pin the baseline

4.3 3.3 Branch

5 4. Canonical Local Build

5.1 4.1 Runtime

5.2 4.2 Libraries

5.3 4.3 Core_Root

5.4 4.4 Never install the fork system-wide

6 5. Required Hardware Validation

7 6. Final MVP File Layout

8 7. Native Status Contract

9 8. CUDA ABI Contract

9.1 8.1 Required CUDA symbols

10 9. CUDA Driver Ownership Rules

11 10. Kernel Handle Contract

12 11. Device Allocation Contract

13 12. Managed Internal Bridge Contract

14 13. Test Reflection Bridge

15 14. Accelerator IR Specification v1

15.1 14.1 Types

15.2 14.2 Value IDs

15.3 14.3 Block IDs

15.4 14.4 Opcodes

15.5 14.5 Parameter ABI

15.6 14.6 LoadF32 semantics

15.7 14.7 StoreF32 semantics

15.8 14.8 GlobalIndexX semantics

15.9 14.9 CFG restrictions

15.10 14.10 MVP merge restrictions

16 15. IR Verifier Contract

17 16. IR Text Format

18 17. PTX Contract v1

19 18. Launch Contract

20 19. Canonical Test Kernel Source

21 20. Ticket Sequence Overview

22 21. A000 — Pin Upstream Baseline

22.1 Objective

22.2 Files

22.3 Steps

22.4 Test

22.5 Done

23 22. A001 — Build Private Core_Root

23.1 Objective

23.2 Commands

23.3 Failure procedure

23.4 Done

24 23. A002 — Add Accelerator Smoke Test

24.1 Objective

24.2 Done

25 24. A003 — Wire First QCall

25.1 Objective

25.2 Failure: EntryPointNotFound

25.3 Failure: managed method missing through reflection

25.4 Done

26 25. A010 — Dynamically Load NVIDIA Driver

26.1 Objective

26.2 Debug

26.3 Done

27 26. A011 — Resolve CUDA Initialization Symbols

27.1 Objective

27.2 Done

28 27. A012 — Retain Primary CUDA Context