Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions .gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,10 @@
mandelbrot.png
result.png
newton.png
lena-sobel.bmp
strategy_heatmap.png
# machine-local package feed override
/nuget.config
*.pdb
bin
*.gguf
Expand Down
338 changes: 338 additions & 0 deletions e -i HEAD~3

Large diffs are not rendered by default.

37 changes: 37 additions & 0 deletions src/1.Simple/Builtin/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
BUILTIN FUNCTION EXAMPLE
=========================

What this sample shows
-----------------------
Normally, when Hybridizer compiles your C# code for the GPU or CPU, it
translates every instruction itself. But sometimes you want to tell
Hybridizer: "for this specific .NET method, don't translate it yourself —
just replace it directly with this native function instead."

This is what a "builtin" is: a manual mapping between an existing .NET
method and a native function that already exists in CUDA (or AVX/OMP).

The example
-----------
This sample sums an array of 1024 integers (0, 1, 2, ... 1023) in
parallel, using System.Threading.Interlocked.Add to safely add each
value into a shared result variable without threads overwriting each
other.

Normally, Interlocked.Add doesn't mean anything on a GPU. To make it
work, the file "sample.builtins" tells Hybridizer: "whenever you see
Interlocked.Add being called, replace it with CUDA's native atomicAdd
function instead." atomicAdd is a function built into CUDA that does
exactly the same thing — it lets many threads add to the same variable
at once without conflicts.

So this sample is really about showing that you can reuse familiar
.NET methods (like Interlocked.Add) and have Hybridizer swap them out
for the right native equivalent on the target hardware — without you
having to rewrite your code.

Expected output
----------------
sum = 523776

(this is the sum of all integers from 0 to 1023)
1 change: 0 additions & 1 deletion src/1.Simple/HelloWorld/Program.cs
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,6 @@ static void Main(string[] args)

// run .Net method
Run(N, adotnet, b);

// verify the results
for (int k = 0; k < N; ++k)
{
Expand Down
46 changes: 46 additions & 0 deletions src/1.Simple/HelloWorld/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
HELLOWORLD EXAMPLE
====================

What this sample shows
-----------------------
This is the most basic Hybridizer sample: it takes a piece of ordinary
C# code, runs it on the GPU, then runs the exact same code on the CPU
(through plain .NET), and checks that both give the same result.

The idea is to prove that Hybridizer really does what it promises: you
write your logic once in C#, and it works identically whether it
executes on the GPU or on the CPU — no separate GPU-specific code
needed.

The example
-----------
The method "Run" simply adds two big arrays of random numbers together,
element by element (a[i] += b[i]), using Parallel.For so each element
can be computed independently and in parallel. Two arrays are used:
16 million doubles each (about 268 MB), which is a size that fits on
basically any CUDA-compatible GPU.

The same "Run" method is called twice:
1. Once through "wrapped.Run(...)", which sends it to the GPU via
Hybridizer.
2. Once as a normal, unmodified C# method call, which runs on the CPU
as regular .NET code.

The program then compares the two resulting arrays value by value. If
every single value matches, it means the GPU computed exactly the same
thing as the CPU would have — confirming the translation from C# to
GPU code was correct.

Expected output
----------------
Expected output
----------------
GPU Results :
0.732961, 1.245..., 0.891..., (... 16,777,216 comma-separated values ...)
CPU Results :
0.732961, 1.245..., 0.891..., (... 16,777,216 comma-separated values ...)
DONE

(the two lists should be identical, value by value; if any value
differs, the program prints "ERROR !" and stops instead of reaching
"DONE")
46 changes: 46 additions & 0 deletions src/1.Simple/InOut/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
INOUT EXAMPLE
===============

What this sample shows
-----------------------
When Hybridizer sends an array to the GPU, by default it has to copy it
in both directions: from CPU to GPU before the kernel runs, and from
GPU back to CPU after — just in case the array was both read and
modified. But often, an array is only ever read (input) or only ever
written (output), and copying it in the unused direction is wasted
time.

The [In] and [Out] attributes let you tell Hybridizer exactly how an
array is used, so it can skip the unnecessary copy. This sample
measures how much time that actually saves.

The example
-----------
The same computation, "dst[i] = src[i] + i", is run twice, using two
almost identical methods:

1. "NoAttributes": takes "dst" and "src" as plain arrays, with no hint
about how they're used. Hybridizer copies both arrays in both
directions to be safe.

2. "Attributes": the exact same computation, but "dst" is marked [Out]
(only ever written to, never read) and "src" is marked [In] (only
ever read, never written). This tells Hybridizer it can skip copying
"dst" to the GPU beforehand, and skip copying "src" back afterward.

Both versions process 16,777,216 (2^24) random integers, and a
Stopwatch measures the execution time of each version separately.

Expected output
----------------
Selecting device <GPU name> with compute capability <XX>
running generated CUDA (no attributes)
no in/out attribute time : <X> ms
running generated CUDA (attributes)
in/out attributes time : <Y> ms
OK

(the exact GPU name and timings depend on your hardware, but "Y"
should generally be lower than "X", showing the benefit of the [In]/
[Out] attributes; the program prints an error and stops instead of
"OK" if a CUDA error is detected)
13 changes: 13 additions & 0 deletions src/1.Simple/Intrinsics/Program.cs
Original file line number Diff line number Diff line change
Expand Up @@ -53,10 +53,23 @@ static void Main(string[] args)
const int N = 1024 * 1024 * 32;
half2[] input = new half2[N];


half2 before = input[0];

HybRunner runner = SatelliteLoader.Load();
dynamic wrapped = runner.Wrap(new Program());

// Chronomètre (point 2)
var sw = System.Diagnostics.Stopwatch.StartNew();
wrapped.Compute(input, N);
cuda.ERROR_CHECK(cuda.DeviceSynchronize());
sw.Stop();

// Vérification du résultat (point 1)
Console.WriteLine($"Value Before : {before}");
Console.WriteLine($"Value After exp12 : {input[0]}");
Console.WriteLine($"GPU Time for {N:N0} elements : {sw.ElapsedMilliseconds} ms");

}
}
}
96 changes: 96 additions & 0 deletions src/1.Simple/Intrinsics/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,96 @@
INTRINSICS EXAMPLE
=====================

What this sample shows
-----------------------
GPUs can compute using different levels of numeric precision. Standard
precision uses 32-bit floats (float), but modern GPUs also support
16-bit "half precision" floats, which take half the memory and can be
computed faster — at the cost of a much smaller range of representable
values (roughly up to 65504) and less numeric accuracy.

This sample demonstrates "mixed precision" arithmetic: using the half2
type, which packs two 16-bit half-precision numbers together, and
performing a custom exponential function on them entirely in half
precision.

Hardware requirement
---------------------
Mixed precision (half2) requires a GPU with compute capability 5.3 or
higher — this covers Volta, Pascal, and Jetson TX1 GPUs onwards. On
older GPUs (Maxwell, Kepler, Fermi), the generated code won't even
compile. Note that Pascal itself is now old enough that NVIDIA has
dropped it from the latest drivers.

The example
-----------
The sample defines its own approximation of the exponential function
("exp"), computed as a 14-term polynomial (a Taylor series
approximation), operating entirely on half2 values. This custom "exp"
is then applied 12 times in a row ("exp12") to every element of a large
array (33,554,432 elements), fully in parallel on the GPU.

Because "exp12" applies the exponential 12 times in a row, even a tiny
starting value grows extremely fast — well beyond what 16-bit half
precision can represent (which tops out around 65504). This makes the
sample a good illustration of a real limitation of mixed precision:
values can silently overflow into Infinity or NaN (Not a Number) if
you're not careful about the range of numbers you're working with.

Expected output
----------------
INTRINSICS EXAMPLE
=====================

What this sample shows
-----------------------
GPUs can compute using different levels of numeric precision. Standard
precision uses 32-bit floats (float), but modern GPUs also support
16-bit "half precision" floats, which take half the memory and can be
computed faster — at the cost of a much smaller range of representable
values (roughly up to 65504) and less numeric accuracy.

This sample demonstrates "mixed precision" arithmetic: using the half2
type, which packs two 16-bit half-precision numbers together, and
performing a custom exponential function on them entirely in half
precision.

Hardware requirement
---------------------
Mixed precision (half2) requires a GPU with compute capability 5.3 or
higher — this covers Volta, Pascal, and Jetson TX1 GPUs onwards. On
older GPUs (Maxwell, Kepler, Fermi), the generated code won't even
compile. Note that Pascal itself is now old enough that NVIDIA has
dropped it from the latest drivers.

The example
-----------
The sample defines its own approximation of the exponential function
("exp"), computed as a 14-term polynomial (a Taylor series
approximation), operating entirely on half2 values. This custom "exp"
is then applied 12 times in a row ("exp12") to every element of a large
array (33,554,432 elements), fully in parallel on the GPU.

Because "exp12" applies the exponential 12 times in a row, even a tiny
starting value grows extremely fast — well beyond what 16-bit half
precision can represent (which tops out around 65504). This makes the
sample a good illustration of a real limitation of mixed precision:
values can silently overflow into Infinity or NaN (Not a Number) if
you're not careful about the range of numbers you're working with.

Expected output
----------------
Value before : Hybridizer.Runtime.CUDAImports.half2
Value after exp12 : Hybridizer.Runtime.CUDAImports.half2
GPU time for 33 554 432 elements : ~400 ms

(the exact timing depends on your GPU; starting from a small value like
0.001, the repeated exponential is expected to overflow to Infinity for
every element well before the 12th iteration — this is the expected
and instructive behavior of this sample, not a bug)


(the exact timing depends on your GPU; starting from a small value like
0.001, the repeated exponential is expected to overflow to Infinity for
every element well before the 12th iteration — this is the expected
and instructive behavior of this sample, not a bug)
5 changes: 0 additions & 5 deletions src/1.Simple/Intrinsics/README.txt

This file was deleted.

34 changes: 18 additions & 16 deletions src/1.Simple/Malloc/Program.cs
Original file line number Diff line number Diff line change
Expand Up @@ -31,18 +31,9 @@ public static double apply(double[] stencil, double[] src, int i)
[EntryPoint]
public static void test([Out] double[] dest, [In] double[] src, int N)
{
double[] stencil =
[
-4.0,
-3.0,
-2.0,
-1.0,
0.0,
1.0,
2.0,
3.0,
4.0,
];
double[] stencil = new double[9];
for(int i = 0; i < 9; ++i) { stencil[i] = i - 4.0; }

for (int k = 4 + threadIdx.x + blockIdx.x * blockDim.x; k < N - 4; k += blockDim.x * gridDim.x)
{
dest[k] = apply(stencil, src, k);
Expand All @@ -51,21 +42,32 @@ public static void test([Out] double[] dest, [In] double[] src, int N)

static void Main(string[] args)
{
const int N = 1024*1024*32;
const int N = 1024 * 1024 * 32;
double[] src = new double[N];
double[] dst = new double[N];
Random rand = new();
for(int i = 0; i < N; ++i)
for (int i = 0; i < N; ++i)
{
src[i] = rand.NextDouble();
dst[i] = src[i];
}

cuda.GetDeviceProperties(out cudaDeviceProp prop, 0);

HybRunner runner = SatelliteLoader.Load().SetDistrib(prop.multiProcessorCount, 512);
dynamic wrapper = runner.Wrap(new Program());

var sw = System.Diagnostics.Stopwatch.StartNew();
wrapper.test(dst, src, N);
cuda.ERROR_CHECK(cuda.DeviceSynchronize());
sw.Stop();

Console.WriteLine($"GPU time for {N:N0} elements : {sw.ElapsedMilliseconds} ms");

Console.WriteLine("Result Sample (from i = 4 to 13) :");
for (int i = 4; i < 14; ++i)
Console.Write($"{dst[i]:F4}, ");
Console.WriteLine();
}
}
}
}
44 changes: 44 additions & 0 deletions src/1.Simple/Malloc/README.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,44 @@
MALLOC EXAMPLE
=================

What this sample shows
-----------------------
Normally, you cannot use .NET's "new" keyword to allocate objects from
code that runs on the GPU — GPU threads don't have the same kind of
memory management as the CPU. There is one exception though: arrays.
Each GPU thread is allowed to allocate its own small, local array with
"new", which gets automatically freed once that thread is done with it.

This sample is a toy example (the author's own comment says "no
physical meaning at all") whose only purpose is to demonstrate that
this thread-local array allocation works — not to compute anything
meaningful.

The example
-----------
Every GPU thread computes one output value by:
1. Allocating its own small array of 9 numbers ("stencil"), filled with
the values -4, -3, -2, ... up to 4. This "new double[9]" happens
individually inside each thread — it is not shared or precomputed
on the CPU.
2. Using that array as a set of 9 weights, applied to 9 neighboring
values around its position in the "src" array (4 values before,
itself, and 4 values after) — a classic "stencil" pattern used in
things like signal processing or numerical simulations.
3. Storing the weighted sum into "dest".

This is repeated for 33,554,432 (32 * 1024 * 1024) values, each
computed independently and in parallel, with every thread creating and
discarding its own little array along the way.

Expected output
----------------
GPU time for 33 554 432 elements : ~450 ms
Results Samples (from i = 4 to 13) :
-1.2345, 0.8821, ...

(exact timing and values depend on your GPU and the random seed; since
"src" is filled with random numbers between 0 and 1, the sample output
values will differ on every run — this sample doesn't verify
correctness against a CPU reference, so the numbers are only shown to
confirm the GPU computation actually produced something)
Loading