diff --git a/.gitignore b/.gitignore index 8b53d68..c3f3c8d 100644 --- a/.gitignore +++ b/.gitignore @@ -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 diff --git a/e -i HEAD~3 b/e -i HEAD~3 new file mode 100644 index 0000000..4dfa671 --- /dev/null +++ b/e -i HEAD~3 @@ -0,0 +1,338 @@ +5236c12 (HEAD -> Modifications-Milan, origin/Modifications-Milan) Added README and upgraded Mandelbrot +d6fa2f3 Amélioration ou Ajout des README +99127d1 (status) Première Vague Changements Codes +b39d705 Rajotut de StopWatch et Mise en Page +5b07fa1 (master) Merge pull request #115 from hybridizer-io/fix/malloc-avoid-initializearray +8db5964 chore: gitignore sample outputs and local nuget.config +f5069ae fix(Malloc): initialize stencil per-element to avoid RuntimeHelpers.InitializeArray +fe7da95 Feat: wmma gemm sample +704749b chore: reference Hybridizer.Runtime.CUDAImports 3.7.1 across all samples +3d7038d feat(WmmaGemm): switch to BASIC-mode hybridizer CLI (Essentials 7.6.0) +87865d6 refactor(WmmaGemm): drop duplicate SatelliteLoader, document BASIC-mode wiring +3cb5d92 feat(6.Advanced/WmmaGemm): half-precision GEMM sample with tensor cores +3e1d1f0 Merge pull request #113 from hybridizer-io/feat/tiny-llama +de88cfa tiny-llama +1952689 update readme +4195583 (tag: v3.4.0) Update to cudaimports latest and autodetect nvcc/nvrtc +5326c15 feat: add StrategyBacktest finance sample and MonteCarloHeatEquation visualization utilities +dd6f9d7 fix: faster mandelbrot +d6ff2db feat: in-out-attributes +01d9576 remove CI as we need a CUDA driver to build solution +36be162 build in CI +535f082 feat: in-out attributes usage +faca63f feat: In-Out Attribute usage +0a0e75c reference cudaimports 3.0 +6fa7a25 feat: migrate to new version of hybridizer +2b278e2 fix generic reduction +fddf8a6 feat: add FPS indicator +dcd501b (tag: v3.0.0) cleanup folrder tree +673a6fc migrate last samples +b74d134 migrate 4.Finance +ebe1cfd migrate 3.Maths samples +981c739 migrate imaging samples +0680e99 migrate simple directory to new hybridizer +3b262a9 update license +1f153c5 delete base.config +4877680 update readme +b06f225 delete jupyter my ass +2fd55a8 move to new version : initial cleanup +7b7e5fb add a sample for generic memory access (no special meaning, just feature demonstration) +cac7bf8 support for cuda12.2 except NPP_ImageSegmentation which needs further investigation +52e3908 (tag: v2.0.1, tag: v2.0) Merge pull request #104 from altimesh/december-update +e47184e update medatada +3904ae9 finish unification of VS solutions and CUDA versions into - one solution - two config builds +8d69084 everything runs smoothly +abbd0ac all projects compile and generates CUDA which compiles +62be07d samples close to OK w/ cuda 11.6 +889d4f6 samples cuda 11.6 close to OK +8267a11 (tag: 1.2.10616) fix jitter version no more float3/float4 constructors -- replaced by make_floatn +8e56162 maxlabel is inclusice +0cbe2d7 (tag: 1.2.10607) Merge pull request #87 from altimesh/september_update +7ebe84b add NPP sample +2b0b123 (tag: 1.2.10484) cuda 10.1 and visual 2019 +be9007f (tag: 1.2.10341) Merge pull request #77 from altimesh/february-update +42b99c6 fix generic functions sample +bb0fd2b remove dependency to native library for ConjugateGradient implement scalar prod in c# +e3efa70 fix reference to cudaimports +43e43ea add generic function sample +846b29f remove dead comments + add a todo +c6fbe66 add virtual function sample +be487dd fix cooperative groups sample +3c56134 add a cooperative group (multi block) sample +3d36f0c update malloc sample (new double[] instead of malloc) +aff612e Merge pull request #72 from altimesh/january_update +4e96d89 fix #71 (jitter version) +cbdf8c9 Update README.md +c336bef (tag: 1.2.10271) Merge pull request #70 from altimesh/january_update +e5d007f add one exp to intrinsic sample +e58d188 delete old files from cuda 10 samples (makefiles, run.sh and native reduction) which are now useless +c9c7367 improve intrinsics sample (half2) +7707beb improve printf sample +f673ead add simple cufft example +8ec691d fix shared memory amount in generic/lambda reduction +c2b9c97 cuda 8 does not support c++14 -- rollback to c++11 remove generated files +13f63d1 move c++ projects from c++11 to c++14 +47f2536 (tag: 1.2.10227) Merge pull request #67 from altimesh/december_update +7e73ced Fix samples CUDA 10.0 +9addea9 atomic expressions in generic and lambda reduction +bf6b2f8 (tag: 1.2.10129) Merge pull request #65 from altimesh/november_update +a46f38e Merge branch 'november_update' of https://github.com/altimesh/hybridizer-basic-samples into november_update +a8a1d13 add simpledecorator sample +d829b36 add sample jupyter hybridizer zero-copy +a3d6b18 Update README.md +910afb5 fix sobel compute +7621986 add generic reduction sample add lambda reduction sample +4295377 (tag: 1.2.10108) - add winform sample (mandelbrot renderer) - delete some generated files +50bdf8b (tag: 1.2.10021) rename basicsamples to basicsamples_cuda8 add cuda10 samples +8f531a0 Merge branch 'master' of https://github.com/altimesh/hybridizer-basic-samples +4b2313e fix mandelbrot renderer (winform sample) + add info button +606e35c Merge pull request #62 from altimesh/fduguet-avx512-labs +83ed8b3 Adding test mechanics +7a97fe1 Adding test mechanics +14f5446 faster blackscholes sample +f542c52 hint path for CUDAImports on linux +ba9e142 removing unused comment +1031ca1 removing unused comment +d13c33f fix name of naive-add.cs +a6646c3 (tag: 1.1.9833) Merge pull request #59 from altimesh/september_update +1e43543 fix cudaruntimeversion in CUDA9.2 samples +cc70b32 fix project dependencies +f33f08e reflect package changes (lib path) + fix postbuild event in cuda projects +8b8c9c2 remove unused option (/arch:AVX512) +80d60d9 update compiler flags for optimization in AVX* +f13ddbd disable flavors at compile time +8f8d406 add winform sample (AVX* + CUDA) : Mandelbrot +afaabc3 Create Jupyter/Labs-AVX512/README.md +3743933 Revert "rename Labs to Labs-CUDA" +f93926b rename Labs to Labs-CUDA +c34907f Fix path in jupyter notebook +d1a99e5 modify path +eb17907 add jupyter notebook for Texture and Surface memory +132dc99 (tag: 1.1.9554) Merge pull request #54 from altimesh/july_update +e4c1564 fix FPS display (typo) +6029608 fix reference path for vb project +7c438a6 Merge branch 'july_update' of https://github.com/altimesh/hybridizer-basic-samples into july_update +fee6756 add kernel duration and FPS counter in mandelbulb +6f38599 fix typo +03fe150 add Conjugate Gradient method with ResidentArray to Labs +377e6b9 Merge branch 'july_update' of https://github.com/altimesh/hybridizer-basic-samples into july_update +01d930b resolve path on median Filter notebook +fa7eef8 jupyter must be run in a visual studio command prompt +5724f95 better readme for jupyter notebooks +b99cebe fix hybridizer path in mandelbulb example +3a76450 MANDELBULB : cleanup code -- animate +91b97f0 mandelbulg opengl (quick and dirty) +3394e07 fix interlocked builtin file +3a6452c reference CUDAImports directly, as it's now in the GAC +ec0de96 Mandelbulb : display image once done +60434fa add mandelbulb sample (simple version no OPENGL) +3d9ebfa add reduction sample with shared memory and atomics +1f05ee2 Merge branch 'july_update' of https://github.com/altimesh/hybridizer-basic-samples into july_update +52e3a27 simple recursion sample +ccf71d5 add two dimensional array notebook +c81b79d improve constant memory notebook +741df0f notebook for Constant memory completed +53d769b Add the beginning of the Constant memory notebook +96892c2 fix out of bound +cedf9b1 fix typo +aba3dc3 add Parallel.For and fix typo +88af8cb add median filter lab +625e5b4 fix typo +6ef7571 add missing file +9e69911 Add Stream jupyter notebook +9f5446e Update README.md +11c9949 (tag: 1.1.9368) Merge pull request #52 from altimesh/june_update +9834c26 remove AnyCPU config from 9.2 solution +c91a607 delete large checkpoint file +13aaf32 remove typo +8874322 Merge branch 'master' into june_update +0b5ef24 downgrade toolset to v140 as visual Studio broke cuda support once again +66effe0 remove CUDA_GL_Interop.cs as it's integrated in CUDAImports as of latest version +fac71ac clean 02-lab +cb21f56 clean Hello-world 2 +e2afbfd cleanup Hello World lab +3a16128 add vector-add lab +28dff92 HYB-1123 Adding hello world +7029f30 fix readme and remove deprecated material +6116a11 remove dead code +3c2adda Merge branch 'june_update' of https://github.com/altimesh/hybridizer-basic-samples into june_update +316ae4d add NBody simulation for opengl interop +aad7dc7 Change hybrunner.Cuda("dllName") to hybrunner.Cuda() +7edbd88 fix vb sample +8f2b254 add vb sample +89bd942 better stream management (from wrapped object instead of runner, allowing to have only one runner per flavor) +2ec0ddf (tag: 1.1.9352) Update Streams_CUDA.vcxproj +79893a2 Update README.md +44382d4 (tag: 1.1.9225) Update README.md +ec9622c Merge pull request #50 from altimesh/april_update +87328b3 fix toolsets since cuda 9.2 and v141 are compatible fix build dependencies +515fe65 cuda 9.2 support +7839acc add builtin sample (atomicadd) +cd31e1a Update ISSUE_TEMPLATE.md +62c812a Changing pointer arithmetic to ref to avoid usage of unsafe. +29f698f changing pointer to ref parameters for cusparse. +222b8e4 more files in gitignore +07cecdc Update README.md +d6a1464 remove initial cd in some Makefile +3c30b00 Merge branch 'april_update' of https://github.com/altimesh/hybridizer-basic-samples into april_update +046aeda Test commit with removing of cd in some Makefile +ef61e39 Update Makefile +6cd4743 add constant memory sample move lena to share location +4d213c3 texture and surface sample +8836352 stream example +71b65d2 (tag: 1.1.5931.8943) Merge pull request #40 from altimesh/march_update +aa270e9 update makefile to explicitely reference the right version of nvcc +8977a2d remove sobel2D from linux samples add try/catch on process start for non interactive (no X11) linux machines +6fccf46 fix cuda version for cuda8 +573535c fix jitter path for cuda 9.0 +943029b fix project dependency +9be4d4d add missing builtins in config files +219a0ee replace spaces by tabs in makefiles +f6b359c Merge branch 'march_update' of https://github.com/altimesh/hybridizer-basic-samples into march_update +17ef166 update projects for linux support add makefiles and run scripts +0587260 Update README.md +a53d145 remove visual studio version specific projects -- reduce repository size and ease bugfixes and future developments +941213d (tag: 1.1.5787.8707) missing file +6d1db13 fix #37 (missing vcxproj) Added README to mixed precision sample (as it doesn't compile on pre Jetson TX1 Gpus. Fixed imports (props instead of targets) to fix warnings +a79b2af (tag: 1.1.5727.8609) Merge pull request #35 from altimesh/jan18update +606c714 fix CUDA 9.1 vs2017 project dependencies removed half2 example, as it doesn't compile on kepler. +3fe5b96 added a check to issue-template (installing vc++) +c3d5e41 fix vs2015 build removed fp16 sample in cuda 8 +5840cca add missing files for CUDA 8 +7b6ea2f moved all vs2017 projects to v140 added Intrinsics sample in visual 2017 examples +cb77afa add simple example of intrinsic includes (using half2 type) +477920d update all vcxproj to match new vsix version +63e9a53 (tag: 1.1.5685.8524) Merge pull request #33 from altimesh/cuda91 +4a197d8 fix project dependencies +b3daf59 cuda 9.1 support + minor bug fixes +1a4d803 Update README.md +cd6c816 Update README.md +d73179a Update ISSUE_TEMPLATE.md +27313f4 (tag: v1.1.5656.8398) fix some output dll paths +a033db1 add project dependency +f32f4c1 Merge pull request #30 from altimesh/cuda9andvisualstudiosupport +00ee0e1 add notes to readme -- warnings for update/build +b37745a fix projects - remove cublas and cusparse from cuda9 samples - fixed projects references - cleanup solution configurations +89cc5ac make all path relatives +22b31ef other intial changes +f07c9fd initial changes +6a15b0e Merge pull request #28 from altimesh/add-code-of-conduct-1 +c0e6332 Create CODE_OF_CONDUCT.md +56bb9e8 Create CONTRIBUTING.md +e6e6392 Create ISSUE_TEMPLATE.md +5d9c194 delete generated files +6dcd65d Revert wrong part of commit "add sobel_lock to demonstrate unsafe and direct calls to cuda runtime" +4db6c02 add missing /unsafe to Sobel_Lock.csproj +ff2ef2e use pinned memory instead of manual device allocation/memcpies +8056b4e add sobel_lock to demonstrate unsafe and direct calls to cuda runtime +9157a82 Merge branch 'master' of https://github.com/altimesh/hybridizer-basic-samples +065f751 use getlength on 2D arrays (sobel 2D) remove useless solution configurations +9594406 remove x86 and AnyCPU solution configurations +3a87426 Update LICENSE +15161bf Update LICENSE +baecd88 Create LICENSE +72f256a add link to visual studio gallery in readme.md +5c46bcd remove useless intrinsic function +f0be218 remove unsafe from cusparse sample +d1c7dc4 Adding link to wiki +4ba2c9a Naming changes in Hybridizer +5054505 Merge branch 'master' of https://github.com/altimesh/hybridizer-basic-samples.git +e636b5f ignoring x64 directory +604d923 ignoring perf sessions +973ac08 Changed product name to Essentials +e83fa98 fix import path (essentials instead of basic) +56cec05 rename basic to essentials +da36545 malloc/free sample +729a8da fix out dir of sobel_2D +fae8513 sobel 2D +137d3d3 basic samples -- minor upgrade and bugfixes +e6c5553 essentials wiki -- float4 images +b0ac370 add mandelbrot image +f4f1806 more info in readme +8019835 fix output directories +5b7579e vsix fix sobel +52d55f0 fix naivematrix +fdfd50d fix blackscholes float 4 +129ad0a improve performances on 2D kernels +21971e9 move sobel to right directory +30dc609 add subtitles +cf7bfc7 delete unused project +12acea1 revive cublas example +ddb66c5 add video explanation +abce1a7 montercarlo : finalize +6d22912 add resulting image + acceptable default configuration +187c9f2 montecarlo : add another problem type +b8779cc add another template parameter +ef0412f add montecarlo heat equation sample +e6e0a6e conjugate gradient on GPU +44cc01f Merge branch 'master' of https://github.com/altimesh/hybridizer-basic-samples.git +85e1c2e fix cusparse +51ec95e cublas +41d17ad Merge branch 'master' of https://github.com/altimesh/hybridizer-basic-samples.git +9e2fa6d fix basic samples (reference to cudaimports) +8b75c2d Merge pull request #27 from altimesh/axel-branch +c5baa92 Merge branch 'master' into axel-branch +e051a9e save +4f734ee upgrade basic samples to new version of basic (breaking changes!) +1575fb2 Update Program.cs +e0c72c9 Update Program.cs +52baac3 Update Program.cs +3dadc62 Update Program.cs +30afbdb Update Program.cs +6221e25 Merge pull request #26 from altimesh/axel-branch +1b3ba61 hyb basic sample : change on the FillMatrix method , uses of min and max parameters with default values. +c316ab1 Merge pull request #25 from altimesh/axel-branch +f4786fc hyb basic sample : Creation of a library that using cub, with complex, cache and two operation on vectors (sum and scalar product). +c5703ad Merge pull request #21 from altimesh/axel-branch +22e2d6f hyb basic sample : change namespace typo +cd783b8 Hyb Basic sample : creation of a Conjugate Gradient method with sparse matrix, only in C#. +1a6e849 Hyb Basic Sample: Creation of an Utils directory that gather some files reusable. +ca23168 Merge pull request #20 from altimesh/axel-branch +3ba0af8 hyb basic sample: Sparse matrix with cusparse. +0486e25 Hyb Basic sample: SparseMatrix, print of the bandwith on C# and CUDA +24c5d71 Merge pull request #19 from altimesh/axel-branch +8be4116 Hyb Basic samples: SparseMatrix do some changes like creation of SparseMatrixReader and Vector Reader in two other files and little fixes. +c1eb8f7 Hyb Basic samples: Creation of SparseMatrix with reading of .mtx files. +489f106 use fastmath for blackscholes float4 (94% of peak bandwidth) -- cuda sample is at 96% +750cf57 Merge pull request #18 from altimesh/axel-branch +4b3d6a6 Update Program.cs +40ecfea hyb basic sample: BlackScholes with Float4, create static class randomExtension and minor fix. +b859dbc hyb basic sample: BlackScholes with float 4 change typo. +146e728 hyb basic sample: creation of BlackScholes with float 4, same compute time than normal for cuda but it's better for c#. +cef3bf1 Merge pull request #17 from altimesh/axel-branch +b283fcb hyb basic sample: BlackScholes display of the calculation error beween cuda and c#. +d78a7e8 hyb basic sample: creation of BlackScholes sample. +6f48412 Merge pull request #16 from altimesh/axel-branch +0d5d82f hyb basic samples: creation and completion of shared memory matrix multiplication sample. +189c830 Merge pull request #15 from altimesh/axel-branch +a4095a4 hyb basic sample: change typo on NaiveMatrix class and the main program. +4c8476c hyb basic sample: separation fo the NaiveMatrix class and the Main program. +0881ccd Merge pull request #14 from altimesh/axel-branch +41ef923 hyb basic sample: minor fix in NaiveMatrix. +fb7d37a hyb basic sample: creation of a NaiveMatrix class, test if the matrix size matchs. +a0e2759 hyb basic sample : creation and completion of a naive matrix multiplication. +aa58036 Merge pull request #13 from altimesh/axel-branch +a72925b Improvement of the newton fractal program. +e8583d6 Merge pull request #12 from altimesh/axel-branch +d6d7301 hyb basic sample: creation and completion of newton fractal. +1410dc5 Update timer.h +89f1d67 Merge pull request #11 from altimesh/axel-branch +4cc5cb4 sobel cuda +f74dab5 Merge pull request #10 from altimesh/axel-branch +30d088d move sobel to imaging +9627e26 Merge pull request #9 from altimesh/axel-branch +c10489b add gitignore +4cea743 fix comment for review +8411d88 hyb basic sample : sobel ready to hybridyze +4171dc2 fix typo in helloworld +110cd3b test commit +494e4bd hyb basic samples : create a c# sobel filter for black and white image with an image to test it. +53b2dae hyb basic samples : attribute randoms doubles to arrays and test if the two results are equals. +0906723 hyb basic samples : change int to double in hello world sample +4dabfbd first commit: multi implementation high resolution timer missing : linux implementation +495ca13 Adding simple printf sample +ffea6eb Adding Mandelbrot float sample +98be4c0 Hello World sample +2a2b67d Adding short Getting Started tutorial +8cbeb7c Initial version +9dd2f3d Initial commit diff --git a/src/1.Simple/Builtin/README.md b/src/1.Simple/Builtin/README.md new file mode 100644 index 0000000..784fa30 --- /dev/null +++ b/src/1.Simple/Builtin/README.md @@ -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) \ No newline at end of file diff --git a/src/1.Simple/HelloWorld/Program.cs b/src/1.Simple/HelloWorld/Program.cs index 39a2a9e..e381e26 100644 --- a/src/1.Simple/HelloWorld/Program.cs +++ b/src/1.Simple/HelloWorld/Program.cs @@ -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) { diff --git a/src/1.Simple/HelloWorld/README.md b/src/1.Simple/HelloWorld/README.md new file mode 100644 index 0000000..7ba35ba --- /dev/null +++ b/src/1.Simple/HelloWorld/README.md @@ -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") \ No newline at end of file diff --git a/src/1.Simple/InOut/README.md b/src/1.Simple/InOut/README.md new file mode 100644 index 0000000..699135e --- /dev/null +++ b/src/1.Simple/InOut/README.md @@ -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 with compute capability +running generated CUDA (no attributes) +no in/out attribute time : ms +running generated CUDA (attributes) +in/out attributes time : 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) \ No newline at end of file diff --git a/src/1.Simple/Intrinsics/Program.cs b/src/1.Simple/Intrinsics/Program.cs index 08a0a74..bbb27f8 100644 --- a/src/1.Simple/Intrinsics/Program.cs +++ b/src/1.Simple/Intrinsics/Program.cs @@ -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"); + } } } \ No newline at end of file diff --git a/src/1.Simple/Intrinsics/README.md b/src/1.Simple/Intrinsics/README.md new file mode 100644 index 0000000..ce538dd --- /dev/null +++ b/src/1.Simple/Intrinsics/README.md @@ -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) \ No newline at end of file diff --git a/src/1.Simple/Intrinsics/README.txt b/src/1.Simple/Intrinsics/README.txt deleted file mode 100644 index 58d8849..0000000 --- a/src/1.Simple/Intrinsics/README.txt +++ /dev/null @@ -1,5 +0,0 @@ -This is sample of mixed precision arithmetic from C#. - -GPU support mixed precision starting from computing capability 5.3, which means Volta, Pascal, and Jetson TX1 - -Before that (Maxwell, Kepler, Fermi...), the generated code won't compile but wait, Pascal is now so old that it has beed removed from latest drivers... \ No newline at end of file diff --git a/src/1.Simple/Malloc/Program.cs b/src/1.Simple/Malloc/Program.cs index 62ba11e..5e14dc5 100644 --- a/src/1.Simple/Malloc/Program.cs +++ b/src/1.Simple/Malloc/Program.cs @@ -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); @@ -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(); } } -} \ No newline at end of file +} diff --git a/src/1.Simple/Malloc/README.md b/src/1.Simple/Malloc/README.md new file mode 100644 index 0000000..b01f126 --- /dev/null +++ b/src/1.Simple/Malloc/README.md @@ -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) \ No newline at end of file diff --git a/src/1.Simple/Mandelbrot/Program.cs b/src/1.Simple/Mandelbrot/Program.cs index 22f593f..ac8c8b6 100644 --- a/src/1.Simple/Mandelbrot/Program.cs +++ b/src/1.Simple/Mandelbrot/Program.cs @@ -72,8 +72,8 @@ static void Main() { const int redo = 20; - IntResidentArray light_net = new(N*N); - IntResidentArray light_cuda = new(N*N); + IntResidentArray light_net = new(N * N); + IntResidentArray light_cuda = new(N * N); #region c# Stopwatch w = new(); @@ -83,7 +83,8 @@ static void Main() ComputeImage(light_net, false); } w.Stop(); - Console.WriteLine($"elapsed time per image (C#) : {w.ElapsedMilliseconds/redo} ms"); + long elapsedCSharp = w.ElapsedMilliseconds; + Console.WriteLine($"elapsed time per image (C#) : {elapsedCSharp / redo} ms"); #endregion c# HybRunner runner = SatelliteLoader.Load().SetDistrib(32, 32, 16, 16, 1, 0); @@ -98,7 +99,24 @@ static void Main() light_cuda.RefreshHost(); // included for fair comparison } w.Stop(); - Console.WriteLine($"elapsed time per image (CUDA) : {w.ElapsedMilliseconds/redo} ms"); + long elapsedCuda = w.ElapsedMilliseconds; + Console.WriteLine($"elapsed time per image (CUDA) : {elapsedCuda / redo} ms"); + #endregion + + #region speedup + double speedup = (double)elapsedCSharp / elapsedCuda; + Console.WriteLine($"Speedup : {speedup:F2}x"); + #endregion + + #region verification + int mismatches = 0; + for (int i = 0; i < N * N; ++i) + { + if (light_net[i] != light_cuda[i]) mismatches++; + } + Console.WriteLine(mismatches == 0 + ? "CPU/GPU results match : OK" + : $"CPU/GPU results differ on {mismatches} pixels !"); #endregion #region save to image @@ -106,9 +124,9 @@ static void Main() for (int k = 0; k < maxiter; ++k) { - byte red = (byte) (127.0F * k / maxiter); + byte red = (byte)(127.0F * k / maxiter); byte green = (byte)(200.0F * k / maxiter); - byte blue = (byte) (90.0F * k / maxiter); + byte blue = (byte)(90.0F * k / maxiter); colors[k] = Color.FromRgb(red, green, blue); } colors[maxiter] = Color.Black; @@ -124,6 +142,7 @@ static void Main() } image.Save("mandelbrot.png", new PngEncoder()); + Console.WriteLine("Image saved to mandelbrot.png"); #endregion } } diff --git a/src/1.Simple/Mandelbrot/README.md b/src/1.Simple/Mandelbrot/README.md new file mode 100644 index 0000000..c2bb165 --- /dev/null +++ b/src/1.Simple/Mandelbrot/README.md @@ -0,0 +1,44 @@ +MANDELBROT EXAMPLE +===================== + +What this sample shows +----------------------- +This sample renders the Mandelbrot set — a well-known fractal image — +and uses it to compare performance between running the same C# code +on the CPU versus on the GPU through Hybridizer, on a computation that +is actually visual and easy to understand. + +The Mandelbrot set is computed by testing, for every pixel of the +image, how many iterations of a simple mathematical formula it takes +before the result "escapes" past a certain threshold (or reaching a +maximum of 32 iterations, meaning that pixel is considered part of the +set). Pixels that escape quickly are colored differently from pixels +that never escape, which is what produces the fractal's iconic pattern. + +The example +----------- +The image is a 4096 x 4096 grid of pixels. For each pixel, the +"IterCount" function runs the iteration formula and returns how many +steps it took (from 0 to 32). + +This whole computation is run twice, 20 times each (to get a stable +average), for a fair comparison: +1. On the CPU, using Parallel.For to spread the work across CPU cores. +2. On the GPU, using the exact same "Run" method through Hybridizer. + +Once computed, the result is turned into colors (each iteration count +maps to a distinct RGB color, black meaning the pixel never escaped) +and saved as "mandelbrot.png". + +Expected output +---------------- +elapsed time per image (C#) : ~180 ms +elapsed time per image (CUDA) : ~4 ms +Speedup : 45.00x +CPU/GPU results match : OK +Image saved to mandelbrot.png + +(exact timings depend heavily on your CPU and GPU; the speedup is +usually very significant here, since this is an "embarrassingly +parallel" workload — every pixel is fully independent from every +other pixel, which is exactly the kind of problem GPUs excel at) \ No newline at end of file diff --git a/src/1.Simple/Printf/README.md b/src/1.Simple/Printf/README.md new file mode 100644 index 0000000..02b136f --- /dev/null +++ b/src/1.Simple/Printf/README.md @@ -0,0 +1,46 @@ +PRINTF EXAMPLE +================= + +What this sample shows +----------------------- +Normally, when you write to the console in C#, that code runs on the +CPU. This sample shows that Hybridizer also lets GPU threads print +directly to the console themselves, using the same Console.Out.Write +syntax you'd use in ordinary C# — no special GPU-specific print +function needed. + +Each GPU thread prints its own identity (which thread it is, which +block it belongs to) along with the value it's working on, which is a +simple way to see the GPU's parallelism directly, instead of just +trusting that "it ran on many threads at once". + +The example +----------- +The array "a" contains 13 integers (1 through 13). The kernel is +launched with 4 blocks of 4 threads each — 16 threads in total, more +threads than there are elements in the array. + +Each thread checks its own global index (based on threadIdx and +blockIdx) and, if that index falls within the array (index < 13), it +prints a line with its thread number, its block number, and the value +of a[i] at that index. The 3 extra threads (16 - 13) simply have +nothing to do and print nothing. + +Because all 16 threads run independently and in parallel on the GPU, +the order in which their print lines actually appear in the console is +not guaranteed — it depends on how the GPU schedules the threads, and +can vary from one run to another. + +Expected output +---------------- +hello from thread = 0 in block = 0 a[i] = 1 +hello from thread = 1 in block = 0 a[i] = 2 +hello from thread = 2 in block = 0 a[i] = 3 +hello from thread = 3 in block = 0 a[i] = 4 +hello from thread = 0 in block = 1 a[i] = 5 +... (13 lines total, one per array element) +---- Done ---- + +(the exact order of the "hello from thread..." lines can differ +between runs — this is expected, since GPU threads execute +concurrently and are not printed in a fixed sequence) \ No newline at end of file diff --git a/src/1.Simple/Recursion/README.md b/src/1.Simple/Recursion/README.md new file mode 100644 index 0000000..529c44e --- /dev/null +++ b/src/1.Simple/Recursion/README.md @@ -0,0 +1,96 @@ +# Recursion — GPU-Accelerated Recursive Factorial (Device-Side Recursion) + +This sample demonstrates that **Hybridizer doesn't require flattening recursive algorithms into loops before running them on the GPU**: a simple recursive factorial (`Fact`) is marked as a `[Kernel]` device function and compiled as-is — each GPU thread evaluates its own recursive call chain, exactly like a CPU thread would. + +The program computes `Fact(b[i])` for **33,554,432** (`1024 × 1024 × 32`) input values on the GPU, then re-checks every single result against the same recursive method called directly on the CPU. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`Run`) and the `[Kernel]` recursive device function (`Fact`), unchanged from its CPU form +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — the input array is generated in-code (`b[i] = i % 11`, so every factorial computed is `Fact(0)` through `Fact(10)`, repeated across all 33M elements). + +## Output + +The program produces: + +- **`OK`** if every one of the 33,554,432 GPU results matches the CPU-computed reference +- An error line and exit code `6` at the **first** mismatch found, of the form: + \`\`\` + Error at : != + \`\`\` + +### Example Output + +\`\`\` +OK +\`\`\` + +## How It Works + +### Recursive kernel + +`Fact` is a plain recursive factorial, written exactly as it would be in ordinary C#: + +\`\`\`csharp +[Kernel] +public static int Fact(int N) +{ + if (N <= 1) return 1; + return N * Fact(N - 1); +} +\`\`\` + +The only constraint Hybridizer imposes on recursive `[Kernel]` functions is that **stack allocations are forbidden** inside them (no `new StackArray` or `stackalloc`) — recursion itself is fully supported. + +### Explicit work distribution + +`Run` calls `Fact` from inside a manual **grid-stride loop** rather than a `Parallel.For` — as the code comments note, calling a recursive function from `Parallel.For` isn't supported yet, so work distribution has to be written explicitly: + +\`\`\` +for (i = threadIdx.x + blockDim.x * blockIdx.x; i < N; i += blockDim.x * gridDim.x) + a[i] = Fact(b[i]); +\`\`\` + +This is also the sample where the GPU launch configuration is left at its **default**: `HybRunner.Load()` is called without `.SetDistrib(...)`, so Hybridizer falls back to `SetDistrib(multiProcessorCount * 16, 128)` automatically. + +### Validation + +After the GPU run (with explicit error checking via `cuda.GetLastError()` and `cuda.DeviceSynchronize()`), the program loops over all 33M elements on the CPU, calling the **exact same `Fact` method**, and compares it against the GPU-computed value — stopping immediately at the first mismatch, if any. + +| Factor | Explanation | +|---|---| +| **Per-thread call stack** | Each GPU thread gets its own CUDA call stack; recursion depth here is small and bounded (`N ≤ 10`), well within the default stack size | +| **Default launch config** | Demonstrates that a working kernel doesn't require manually tuning `SetDistrib` — Hybridizer's default grid/block sizing is enough here | +| **Full-result validation** | Every single element is checked against the CPU reference, not just a sample, since a stack-related recursion bug could show up on only a few threads | + +### Hybridizer: C# philosophy on the GPU + +This is one of the clearest demonstrations of Hybridizer's core value proposition: **you don't have to restructure your algorithm to fit the GPU's execution model**. A recursive method that "just works" on CPU — `Fact`, unchanged — is handed to Hybridizer with a single `[Kernel]` attribute, and it compiles down to genuine device-side recursive CUDA code. The only adjustment needed anywhere in this sample is writing the work-distribution loop explicitly, since `Parallel.For` doesn't yet support calling into recursive functions. \ No newline at end of file diff --git a/src/1.Simple/Reduction/Readme.md b/src/1.Simple/Reduction/Readme.md index b36741e..39e3ce0 100644 --- a/src/1.Simple/Reduction/Readme.md +++ b/src/1.Simple/Reduction/Readme.md @@ -1,4 +1,73 @@ -Reduction sample +# Reduction — GPU-Accelerated Parallel Sum Reduction -using shared memory and atomic operations on final stage -Interlocked.Add being mapped to atomicAdd through a custom builtin \ No newline at end of file +This sample demonstrates a **classic CUDA parallel reduction** using Hybridizer: summing a large array of integers on the GPU with the standard grid-stride accumulation + shared-memory tree reduction + atomic combine pattern, and checking the result against a CPU-computed reference (`Aggregate`). + +The program sums an array of **33,554,432** (`1024 × 1024 × 32`) integers, each randomly `1` (with 20% probability) or `0`, and prints both the GPU-computed sum and the expected sum for comparison. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`ReduceAdd`), including the shared-memory allocation and `__syncthreads` barriers +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — the input array is generated in-code with a random 20% chance of each element being `1` (otherwise `0`), so the expected sum is roughly `N × 0.2 ≈ 6,710,886`, though it varies run to run. + +## Output + +The program produces: + +1. **`sum =`** — the value computed on the GPU +2. **`expected =`** — the same sum computed on the CPU via LINQ's `Aggregate`, as a reference + +### Example Output + +\`\`\` +sum = 6710498 +expected = 6710498 +\`\`\` + +## How It Works + +`ReduceAdd` combines three standard CUDA reduction techniques in a single kernel: + +1. **Grid-stride accumulation** — each thread walks the array with a stride of `blockDim.x * gridDim.x`, accumulating a private partial sum (`tmp`) across however many elements fall to it. This means the kernel works regardless of how `N` compares to the total thread count. +2. **Shared-memory tree reduction** — each thread's partial sum is stored into a shared-memory array (`cache`, sized `blockDim.x`). The block then repeatedly halves the active range (`i = blockDim.x / 2`, then `i >>= 1` each round), with each active thread adding its "far" neighbor's value into its own slot, guarded by `__syncthreads()` between rounds. After `log2(blockDim.x)` rounds, `cache[0]` holds the full sum for that block. +3. **Atomic combine across blocks** — only `threadIdx.x == 0` in each block adds its block's total into the single global `result[0]`, via `Interlocked.Add`, so blocks never race on the same memory location. + +| Factor | Explanation | +|---|---| +| **Two-level reduction** | Per-thread partial sums first, then a per-block tree reduction, then a cross-block atomic combine — the standard way to reduce a huge array with a bounded number of blocks/threads | +| **Shared memory sized at launch** | `SetDistrib(..., BLOCK_DIM * sizeof(int))` reserves exactly `blockDim.x` ints of shared memory per block, matching `cache`'s size | +| **Single atomic per block, not per element** | Only one `Interlocked.Add` call per block (not per thread) keeps atomic contention low even with millions of elements | +| **Grid-stride loop** | Decouples the number of GPU threads launched (`16 × multiProcessorCount` blocks of `256` threads) from the size of the input array | + +### Hybridizer: C# philosophy on the GPU + +This kernel is, like the shared-memory matrix multiplication and N-body samples, **hand-written CUDA logic expressed directly in C#**: explicit shared memory (`SharedMemoryAllocator`), `__syncthreads()` barriers, bit-shift-based tree reduction, and an atomic add (`Interlocked.Add`, mapped straight onto CUDA's atomic operations) are all written as ordinary C# in the kernel body, with `[EntryPoint]` being the only attribute needed to turn it into a GPU-executable method. + +This shows that when an algorithm genuinely needs GPU-specific tricks — like a tree reduction relying on shared memory and synchronization — Hybridizer doesn't get in the way: the same constructs a CUDA C++ programmer would reach for (shared memory, sync barriers, atomics) are available as first-class C# from inside a `[EntryPoint]` method. \ No newline at end of file diff --git a/src/2.Imaging/Sobel/README.md b/src/2.Imaging/Sobel/README.md new file mode 100644 index 0000000..31449aa --- /dev/null +++ b/src/2.Imaging/Sobel/README.md @@ -0,0 +1,85 @@ +# Sobel — GPU-Accelerated Sobel Edge Detection + +This sample demonstrates **GPU-accelerated image processing** using Hybridizer: it loads a grayscale-converted image, runs a classic **Sobel edge detection** filter on the GPU, and saves the result to disk. + +The program loads `lena512.bmp` (512×512), converts it to grayscale, computes the Sobel gradient magnitude for every interior pixel on the GPU, and writes the edge map to `lena-sobel.bmp`. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- `SixLabors.ImageSharp` NuGet package (image loading/encoding) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) +- An input image named `lena512.bmp`, placed next to the built executable (`AppContext.BaseDirectory`) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +Make sure `lena512.bmp` is copied to the output directory (e.g. set it to **Copy to Output Directory** in the project, or copy it manually next to the built binary) — the program looks for it at `AppContext.BaseDirectory`. + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`ComputeSobel`) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — the input filename (`lena512.bmp`) and output filename (`lena-sobel.bmp`) are hardcoded. + +## Output + +The program produces: + +- **`lena-sobel.bmp`** — a grayscale image the same size as the input, where each pixel's brightness represents the strength of the detected edge at that location (bright = strong edge, dark = flat/uniform area) + +There is no console output, timing, or CPU/GPU comparison in this sample — it's a pure image-processing demo. + +## How It Works + +### Grayscale conversion (CPU) + +`ReadImage` converts the loaded RGBA image to a single-byte-per-pixel grayscale buffer using the standard luminance formula: + +\`\`\` +grey = 0.2126 * R + 0.7152 * G + 0.0722 * B +\`\`\` + +This runs on the CPU, before any GPU work — it's cheap relative to the convolution itself and only needs to happen once. + +### Sobel filter (GPU) + +`ComputeSobel` assigns each GPU thread one or more pixels via a **2D grid-stride loop** (rows via `threadIdx.y`/`blockIdx.y`, columns via `threadIdx.x`/`blockIdx.x`). For every interior pixel (borders are skipped, since the 3×3 neighborhood would go out of bounds), it: + +1. Reads the 8 surrounding pixels (`topl`, `top`, `topr`, `l`, `r`, `botl`, `bot`, `botr`) +2. Computes the horizontal gradient `sobelx` and vertical gradient `sobely` using the standard 3×3 Sobel kernels +3. Combines them into a gradient magnitude, `sqrt(sobelx² + sobely²)`, clamped to `[0, 255]` to fit back into a byte + +| Factor | Explanation | +|---|---| +| **Per-pixel parallelism** | Each output pixel's Sobel value is independent of every other — a textbook embarrassingly-parallel image filter | +| **Stencil pattern** | Each thread reads a fixed 3×3 neighborhood from `inputPixel`, a classic GPU stencil-computation pattern (related to the constant-memory stencil samples elsewhere in the codebase) | +| **Border handling** | Pixels on the image edge are left untouched (`output` stays `0`) rather than reading out of bounds | +| **2D grid-stride loop** | `SetDistrib(32, 32, 16, 16, 1, 0)` launches a fixed grid that covers images larger than one launch's worth of threads by looping | + +> **Note:** pixel indexing throughout (`i * height + j`, and `i * width + j` for the Sobel neighborhood) mixes `width` and `height` as the row stride. This is only correct because `lena512.bmp` is square (`width == height`) — using this code as-is on a non-square image would misalign rows. + +### Hybridizer: C# philosophy on the GPU + +`ComputeSobel` is a direct, idiomatic port of a CPU-style Sobel filter: explicit 2D thread indexing, but no shared memory or manual CUDA memory management — reads go straight to the input buffer, one 3×3 neighborhood per pixel. The image I/O (`ImageSharp` loading, grayscale conversion, PNG encoding) stays in ordinary, unattributed C# and runs on the CPU; only `ComputeSobel` is marked `[EntryPoint]` and offloaded to the GPU. + +This demonstrates Hybridizer's core value proposition applied to image processing: **the parts of the pipeline that benefit from massive parallelism (the per-pixel convolution) go to the GPU with a single attribute, while everything else — file I/O, color conversion, encoding — stays as plain, familiar C#**. \ No newline at end of file diff --git a/src/2.Imaging/Sobel_2D/Readme.md b/src/2.Imaging/Sobel_2D/Readme.md index d108adb..306aa39 100644 --- a/src/2.Imaging/Sobel_2D/Readme.md +++ b/src/2.Imaging/Sobel_2D/Readme.md @@ -1,3 +1,85 @@ -Example using 2D arrays +# Sobel2D — GPU-Accelerated Sobel Edge Detection (2D Arrays) -note the --use-hybrid-arrays in the hybridizer command line \ No newline at end of file +This sample demonstrates **GPU-accelerated image processing** using Hybridizer, and specifically that Hybridizer supports **true 2D arrays** (`byte[,]`) as kernel arguments — not just flattened 1D buffers with manual index arithmetic. + +The program loads `lena512.bmp` (512×512), converts it to grayscale, computes the Sobel gradient magnitude for every interior pixel on the GPU using natural `[i, j]` indexing, saves the result to `lena-sobel.bmp`, and opens it automatically. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- `SixLabors.ImageSharp` NuGet package (image loading/encoding) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) +- An input image named `lena512.bmp`, placed next to the built executable (`AppContext.BaseDirectory`) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +Make sure `lena512.bmp` is copied to the output directory (e.g. set it to **Copy to Output Directory** in the project, or copy it manually next to the built binary) — the program looks for it at `AppContext.BaseDirectory`. + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`ComputeSobel`), including the 2D array marshaling for `byte[,]` parameters +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — the input filename (`lena512.bmp`), output filename (`lena-sobel.bmp`), and image size (`512`) are hardcoded. + +## Output + +The program produces: + +- **`lena-sobel.bmp`** — a 512×512 grayscale image where each pixel's brightness represents the strength of the detected edge at that location (bright = strong edge, dark = flat/uniform area) +- The output image is **opened automatically** after processing (`Process.Start`), wrapped in a `try/catch` so the program doesn't fail on headless/non-interactive machines + +There is no console output, timing, or CPU/GPU comparison in this sample — it's a pure image-processing demo. + +## How It Works + +### Grayscale conversion (CPU) + +`ReadImage` converts the loaded RGBA image into a `byte[512, 512]` grayscale buffer using the standard luminance formula: + +\`\`\` +grey = 0.2126 * R + 0.7152 * G + 0.0722 * B +\`\`\` + +This runs on the CPU, before any GPU work — it's cheap relative to the convolution itself and only needs to happen once. + +### Sobel filter (GPU) + +`ComputeSobel` assigns each GPU thread one or more pixels via a **2D grid-stride loop** (rows via `threadIdx.y`/`blockIdx.y`, columns via `threadIdx.x`/`blockIdx.x`), reading the image size directly from the array with `inputPixel.GetLength(0)`. For every interior pixel (borders are skipped, since the 3×3 neighborhood would go out of bounds), it: + +1. Reads the 8 surrounding pixels using natural 2D indexing (`inputPixel[i - 1, j - 1]`, etc. — no manual `row * width + col` arithmetic) +2. Computes the horizontal gradient `sobelx` and vertical gradient `sobely` using the standard 3×3 Sobel kernels +3. Combines them into a gradient magnitude, `sqrt(sobelx² + sobely²)`, clamped to `[0, 255]` to fit back into a byte + +| Factor | Explanation | +|---|---| +| **True 2D arrays on the GPU** | `byte[,]` parameters are passed directly to the kernel — Hybridizer handles the underlying memory layout and index translation, so kernel code reads like ordinary 2D-array C# | +| **Per-pixel parallelism** | Each output pixel's Sobel value is independent of every other — a textbook embarrassingly-parallel image filter | +| **Stencil pattern** | Each thread reads a fixed 3×3 neighborhood from `inputPixel`, a classic GPU stencil-computation pattern | +| **Border handling** | Pixels on the image edge are left untouched (`output` stays `0`) rather than reading out of bounds | +| **2D grid-stride loop** | `SetDistrib(32, 32, 16, 16, 1, 0)` launches a fixed grid that covers images larger than one launch's worth of threads by looping | + +### Hybridizer: C# philosophy on the GPU + +Compared to a flat-buffer Sobel kernel, this version reads noticeably closer to plain, idiomatic C#: `inputPixel[i - 1, j - 1]` instead of `inputPixel[(i-1) * width + (j-1)]`. There is still no shared memory or manual CUDA memory management — reads go straight to the 2D input buffer, one 3×3 neighborhood per pixel — but the multi-dimensional array support removes an entire class of manual index-arithmetic bugs (like row/column stride mix-ups) that a flattened-array version has to get right by hand. + +This demonstrates Hybridizer's core value proposition applied to image processing: **you can keep using .NET's native multi-dimensional arrays on the GPU**, with the parts of the pipeline that benefit from massive parallelism (the per-pixel convolution) offloaded with a single attribute, while file I/O, color conversion, and encoding stay as plain, familiar C#. \ No newline at end of file diff --git a/src/3.Maths/ConjugateGradient/README.md b/src/3.Maths/ConjugateGradient/README.md new file mode 100644 index 0000000..a495a3c --- /dev/null +++ b/src/3.Maths/ConjugateGradient/README.md @@ -0,0 +1,90 @@ +# ConjugateGradient — GPU-Accelerated Iterative Linear Solver (CSR, 1D Laplacian) + +This sample demonstrates a **full iterative solver running almost entirely on the GPU** using Hybridizer: the **Conjugate Gradient** method, solving `A·X = B` for a sparse **1D Laplacian** matrix in CSR format, with every vector operation (SpMV, dot product, SAXPY, copy) offloaded as its own small GPU kernel. + +The program builds a `10,000 × 10,000` 1D Laplacian, solves `A·X = B` for a right-hand side of all `1.0`, and prints the residual norm every 10 iterations until convergence (or up to `1,000,000` iterations). + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from all five `[EntryPoint]` kernels (`ScalarProd`, `Copy`, `Saxpy`, `Fmsub`, `Multiply`) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — matrix size (`10,000`), max iterations (`1,000,000`), and convergence tolerance (`eps = 1e-8`) are hardcoded. As the code comments note, this problem size converges **very slowly** (or may not visibly converge within a reasonable time) without a preconditioner — this sample favors showing the algorithm's GPU structure over solver performance. + +## Output + +The program prints the residual norm (`√⟨R|R⟩`) every 10 iterations, until it drops below `eps`: + +\`\`\` +99.9812 +87.3541 +... +0.0000000089 +\`\`\` + +The final solution `X` is refreshed back to host memory (`X.RefreshHost()`) at the end but not printed. + +## How It Works + +### GPU-resident memory + +Unlike the other samples, vectors here are `FloatResidentArray`/`IntResidentArray` — buffers that **stay on the GPU across kernel calls**, only synchronized with the host explicitly via `RefreshDevice()`/`RefreshHost()`. Since the entire CG loop runs kernel after kernel without ever needing the intermediate vectors on the CPU, this avoids a host/device copy on every one of the potentially hundreds of thousands of iterations — only the inputs (`A`, `X`, `B`) are pushed to the GPU once at the start, and the result is pulled back once at the end. + +### The five GPU kernels + +Conjugate Gradient is built from a handful of basic linear-algebra operations, each implemented as its own kernel: + +| Kernel | Role | +|---|---| +| `Multiply` | Sparse matrix-vector product `A·v` (CSR row-wise dot product), used to compute `AP = A·P` each iteration | +| `Fmsub` | `res = A - m·v` — computes the initial residual `R = B - A·X` | +| `Saxpy` | `res = x + alpha·y` — updates `X`, `R`, and `P` each iteration | +| `Copy` | Simple vector copy, used to initialize `P = R` | +| `ScalarProd` | Dot product `⟨r1\|r2⟩`, using the same shared-memory tree-reduction + atomic-combine pattern as the standalone reduction sample | + +`Multiply`, `Fmsub`, `Saxpy`, and `Copy` are written as plain `Parallel.For` loops — Hybridizer maps each iteration to a GPU thread automatically. `ScalarProd` is the one hand-optimized kernel: shared memory, a tree reduction with `__syncthreads()`, and a single atomic add per block (`AtomicExpr.apply`) to combine each block's partial sum into the final result. + +### Conjugate Gradient algorithm + +The host-side loop (`ConjugateGradient`) implements the textbook CG iteration, with every vector operation dispatched to the GPU: + +1. `R = B - A·X`, `P = R` (initialization) +2. Each iteration: `AP = A·P`; `α = ⟨R|R⟩ / ⟨P|AP⟩`; `X += α·P`; `R -= α·AP`; check `⟨R|R⟩ < eps²` for convergence; `β = ⟨R_new|R_new⟩ / ⟨R_old|R_old⟩`; `P = R + β·P` + +Only a handful of scalars (`r`, `alpha`, `rr`, `beta`) ever touch the CPU each iteration — everything vector-sized stays resident on the GPU. + +| Factor | Explanation | +|---|---| +| **Minimizing host/device traffic** | `FloatResidentArray`/`IntResidentArray` + explicit `RefreshDevice`/`RefreshHost` calls mean the GPU is only synchronized with the host twice total, not once per iteration | +| **Kernel composition** | A non-trivial iterative algorithm is built by chaining several small, reusable GPU kernels, called from an ordinary C# host loop — no single monolithic kernel | +| **Reused reduction pattern** | `ScalarProd` reuses the same shared-memory tree-reduction technique as the standalone `Reduction` sample, applied here to compute dot products instead of a plain sum | + +### Hybridizer: C# philosophy on the GPU + +This sample shows Hybridizer used to build a **complete numerical algorithm**, not just a single kernel: the CG loop itself is ordinary, sequential C# on the host, orchestrating GPU kernels the same way it would orchestrate calls into any other library — `wrapper.Multiply(...)`, `wrapper.Saxpy(...)`, etc. Most of those kernels are simple `Parallel.For` ports; one (`ScalarProd`) needs genuine CUDA-level optimization (shared memory, reduction, atomics), and both styles coexist naturally in the same C# codebase. \ No newline at end of file diff --git a/src/3.Maths/Mandelbulb/README.md b/src/3.Maths/Mandelbulb/README.md new file mode 100644 index 0000000..694fdc3 --- /dev/null +++ b/src/3.Maths/Mandelbulb/README.md @@ -0,0 +1,71 @@ +# Mandelbulb — GPU-Accelerated Real-Time Raymarched Fractal + +This sample demonstrates a **real-time GPU raymarcher** using Hybridizer: for every pixel of an interactive OpenGL window, it sphere-traces a ray through a **Mandelbulb** fractal's distance field, shades the hit point (diffuse, specular, soft shadows), and writes the result directly into a CUDA/OpenGL-shared texture — a step up in complexity from the earlier N-body OpenGL interop sample, using **surface objects** instead of vertex buffers. + +The program opens an `800×600` animated window (camera and light slowly orbiting) and displays a live FPS / kernel-time counter in the console. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support, with OpenGL interop enabled (same GPU driving both the display and the simulation) — **make sure your preferred rendering GPU is CUDA-compatible, not an integrated Intel GPU** (the program prints a reminder about this at startup: check NVIDIA Control Panel → Manage 3D Settings → Preferred Graphics Processor) +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- `OpenTK` NuGet packages (`OpenTK.Graphics`, `OpenTK.Windowing.Desktop`) for the OpenGL window and rendering +- `vertex.glsl` and `fragment.glsl` shader files, present next to the built executable (loaded at runtime via `File.ReadAllText`) — a minimal textured-quad shader pair +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +Make sure `vertex.glsl` and `fragment.glsl` are copied to the output directory — the window fails to load (with a caught exception logged to console) without them. + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`Render`) and its helper functions (`Distance`, `Shadow`, and the `MathFunctions`/`TextureHelpers` intrinsics) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required. This opens an `800×600`, non-resizable window titled **"Mandelbulb"**, rendering the fractal live as the camera and light both slowly orbit around it. Every second, the console prints the average GPU kernel time and FPS. Closing the window ends the program. + +## How It Works + +### Distance-estimated raymarching (sphere tracing) + +Rather than intersecting explicit geometry, the Mandelbulb is defined by a **distance estimator** (`Distance`): given a 3D point, it iterates the Mandelbulb's power-8 formula in spherical coordinates (`theta`, `phi`, `r`) up to `iterations` times (or until the point escapes a radius of `DEPTH_OF_FIELD`), tracking a running derivative (`dr`) used to convert the escape behavior into an approximate **distance to the fractal surface**. + +For each pixel, `Render`: + +1. Computes a camera ray from the pixel's position on a near-field plane through the eye location +2. **Sphere-traces** the ray: repeatedly evaluates `Distance` at the current position and steps forward by that distance (guaranteed not to overshoot the surface), until the estimated distance is below half a pixel's width (a hit) or the ray leaves the depth of field (a miss) +3. On a hit, estimates the **surface normal** by evaluating `Distance` at six points offset slightly along each axis (a numerical gradient / finite-difference normal) +4. Shades the point with a **Blinn-Phong**-style model: diffuse light from `dotNL`, a sharp specular term (`s^35`), and a **soft shadow** (`Shadow`, itself another short raymarch toward the light source) — combined into an RGB color written to the pixel +5. On a miss, writes a simple sky-gradient color instead, based on how many raymarching steps were taken + +### CUDA / OpenGL surface interop + +Instead of a vertex buffer (as in the N-body sample), this sample shares a **2D texture**: `cuda.GraphicsGLRegisterImage` registers the OpenGL texture as a CUDA graphics resource, `GraphicsMapResources`/`GraphicsSubResourceGetMappedArray` retrieve the underlying CUDA array, and `cuda.CreateSurfaceObject` wraps it as a `cudaSurfaceObject_t`. Each frame, the `Render` kernel writes pixel colors directly into that surface via `surf2Dwrite` (an `[IntrinsicFunction]`), and the same texture is immediately drawn on a full-screen quad by OpenGL — no host round-trip for the rendered image at all. + +| Factor | Explanation | +|---|---| +| **Per-pixel parallelism** | Each of the `800 × 600` pixels independently sphere-traces and shades its own ray — an ideal one-thread-per-pixel GPU workload | +| **Distance-field rendering** | Sphere tracing needs far fewer steps than naive ray-sampling, since each step safely advances by the estimated distance to the nearest surface | +| **Zero-copy display** | The GPU writes shading results straight into the texture OpenGL renders from, via CUDA surface objects — the same zero-copy idea as the N-body sample's vertex-buffer interop, applied to a 2D image instead of point data | +| **Live performance readout** | `runner.LastKernelDuration` exposes the GPU kernel's execution time directly from C#, averaged and printed alongside FPS every second | + +### Hybridizer: C# philosophy on the GPU + +This is the most elaborate hand-written CUDA logic among the samples so far: a distance-estimator fractal formula, sphere tracing, finite-difference normals, a full shading model, and low-level CUDA surface/texture interop calls (`cudaCreateChannelDesc`, `CreateSurfaceObject`, `surf2Dwrite`) are all expressed as ordinary C# methods and structs (`float3`, `uchar4`), decorated with `[Kernel]`/`[EntryPoint]`/`[IntrinsicFunction]` where CUDA-specific behavior is needed. `MathFunctions` mirrors CUDA's math intrinsics (`sincosf`, `atan2f`, `rsqrtf`, etc.) as thin C# wrappers, and `TextureHelpers` does the same for the surface/texture API — keeping even fairly advanced CUDA features accessible from a single .NET codebase, with no separate `.cu` files to maintain. \ No newline at end of file diff --git a/src/3.Maths/MonteCarloHeatEquation/README.md b/src/3.Maths/MonteCarloHeatEquation/README.md new file mode 100644 index 0000000..bd3b2a0 --- /dev/null +++ b/src/3.Maths/MonteCarloHeatEquation/README.md @@ -0,0 +1,80 @@ +# MonteCarloHeatEquation — GPU-Accelerated Monte Carlo PDE Solver + +This sample demonstrates a **Monte Carlo random-walk solver for the (steady-state) heat equation** using Hybridizer, and — unlike the other samples — wraps an **instance method on a plain object** (`MonteCarloHeatSolver.Solve()`) rather than a static method on the entry-point class, showing that Hybridizer's `[EntryPoint]`/`Wrap(...)` mechanism works on ordinary object-oriented C# as well. + +The program solves the heat equation on a `128×128` grid using **512 Monte Carlo iterations per point**, then saves the resulting temperature field as a rainbow-colored heatmap, `result.png`. + +> **Note:** the problem/geometry classes referenced here (`I2DProblem`, `SquareProblem`, `SimpleWalker`, `SimpleBoundaryCondition`, and the commented-out `TetrisProblem`/`TetrisBoundaryCondition`) aren't included in the files provided, so the sections below describe what's inferable from `MonteCarloHeatSolver.cs` and `Program.cs` — the exact walker/boundary logic lives in those other files. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- `SixLabors.ImageSharp` NuGet package (PNG output) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`MonteCarloHeatSolver.Solve`), including whatever walker/boundary logic `SquareProblem.Solve(ii, jj)` calls into +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — grid size (`N = 128`) and Monte Carlo iteration count (`iterCount = 512`) are hardcoded. As the code comments note, this is fairly compute-intensive, so these constants may need adjusting depending on your GPU. + +## Output + +The program produces: + +1. **`CUDA time :`** — wall-clock time for the full GPU solve +2. **`result.png`** — a `128×128` heatmap image of the solved temperature field, colored from black (coldest) through red, orange, yellow, green, blue, indigo, violet, to white (warmest) + +### Example Output + +\`\`\` +CUDA time : 1834 ms +\`\`\` + +## How It Works + +### Monte Carlo random-walk PDE solving + +`MonteCarloHeatSolver.Solve()` loops (via `Parallel.For`, transpiled to a GPU thread grid) over every grid point of the problem (`_problem.MaxIndex()` points, mapped back to `(ii, jj)` coordinates via `_problem.Coordinates`), and calls `_problem.Solve(ii, jj)` for each one. + +This is a classic setup for the **"walk on grid" Monte Carlo method**: for a steady-state heat/Laplace problem, the temperature at an interior point equals the *average* boundary temperature reached by a large number of random walks started at that point. Each of the `iterCount` walks per grid point takes random steps (governed by the `SimpleWalker` strategy) until it reaches the domain's boundary (checked by `SimpleBoundaryCondition`), and the estimated temperature is the average of the boundary values hit. + +### Pluggable problem shape and walker strategy + +`SquareProblem` is generic over the walker and boundary-condition types, so the same Monte Carlo driver (`MonteCarloHeatSolver`) can be reused for different domain shapes and stepping rules — the commented-out `TetrisProblem` line in `Main` shows an alternate domain shape (a Tetris-like boundary) plugged into the same solver unchanged. + +### Heatmap coloring + +`GetColor` maps a normalized temperature `[0, 1]` onto a fixed 8-color rainbow gradient (black → red → orange → yellow → green → blue → indigo → violet → white), linearly interpolating between the two nearest colors for smooth shading. + +| Factor | Explanation | +|---|---| +| **Embarrassingly parallel** | Each grid point's Monte Carlo estimate is entirely independent of every other point's — ideal for one-thread-per-point GPU execution | +| **Stochastic method, deterministic parallelization** | Unlike the finite-difference/CG solvers elsewhere in the codebase, this approximates the PDE solution statistically — accuracy depends on `iterCount`, not on convergence tolerance | +| **Object-oriented `[EntryPoint]`** | The kernel is an **instance method** (`Solve()` on `MonteCarloHeatSolver`, closing over `_problem`), wrapped with `runner.Wrap(solver)` — not a static method on the top-level `Program` class like the other samples | + +### Hybridizer: C# philosophy on the GPU + +This sample pushes further than the others into idiomatic object-oriented design: an interface (`I2DProblem`), a generic problem class parameterized by strategy types, and an `[EntryPoint]` on an **instance method of a non-`Program` class** all compile down to a GPU kernel. There's no requirement for the GPU entry point to live on a particular class or be `static` on `Program` — Hybridizer wraps whatever object instance you give it (`runner.Wrap(solver)`) and generates CUDA from its `[EntryPoint]`-decorated method, letting the solver's design stay driven by ordinary C# abstractions (interfaces, generics) rather than by GPU-programming constraints. diff --git a/src/3.Maths/NBody/README.md b/src/3.Maths/NBody/README.md new file mode 100644 index 0000000..916e5bf --- /dev/null +++ b/src/3.Maths/NBody/README.md @@ -0,0 +1,80 @@ +# NBody — GPU-Accelerated N-Body Simulation with OpenGL Interop + +This sample demonstrates a **real-time GPU N-body gravity simulation** using Hybridizer, rendered live with OpenGL through **CUDA/OpenGL interop** (the physics buffers are shared between CUDA and the GPU renderer — no round-trip through the CPU each frame). + +The program initializes a two-cluster "galaxy-like" body distribution, then runs an interactive window where each frame integrates gravity on the GPU and immediately draws the resulting point cloud. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support, with OpenGL interop enabled (same GPU driving both the display and the simulation) +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- `OpenTK` NuGet packages (`OpenTK.Graphics`, `OpenTK.Windowing.Desktop`, `OpenTK.Mathematics`) for the OpenGL window and rendering +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`Solve`) and its `[Kernel]` device functions (`ComputeBodyAccel`, `BodyBodyInteraction`) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required. This opens an `800×600` window titled **"Hybridizer N body simulation"** showing the bodies as gold points on a black background, animated in real time. Closing the window ends the program — there is no console benchmark output in this sample. + +## How It Works + +### Physics: tile-based N-body gravity + +Each GPU thread owns one body. For every frame, `Solve`: + +1. Calls `ComputeBodyAccel` to sum the gravitational pull of every other body on this one +2. Updates velocity from acceleration (`velocity += accel * deltaT`), applies a damping factor (`damping = 0.9995`) to slowly dissipate energy +3. Updates position from the new velocity (`position += velocity * deltaT`) + +`ComputeBodyAccel` uses the classic **tile-based shared-memory algorithm** (from the *GPU Gems 3* N-body chapter): bodies are processed in blocks of `blockDim.x`, with each tile's positions staged into shared memory (`SharedMemoryAllocator`) so every thread in the block reuses the same loaded data instead of hitting global memory per pairwise interaction. `__syncthreads()` calls guard the shared buffer between the load and the compute phase. + +`BodyBodyInteraction` computes the force from one other body `bj` on body `bi`, using a **softened** inverse-cube law (`softeningSquared = 0.00125`, added to `distSqr` before the `rsqrtf`) to avoid numerical blow-up when two bodies get very close. + +### CUDA / OpenGL interop + +Instead of copying positions from GPU to CPU and back to the renderer every frame, `RenderingWindow`: + +1. Allocates two **OpenGL vertex buffers** (`_buffers`) and registers each one as a **CUDA graphics resource** (`cuda.GraphicsGLRegisterBuffer`) +2. Each frame, **maps** both buffers directly into CUDA's address space (`MapResources`) — one as the read-only previous positions, one as the write-only new positions +3. Runs the `Solve` kernel writing straight into the mapped OpenGL buffer +4. **Unmaps** the buffers (`UnMapResources`) and draws the just-written buffer directly with `GL.DrawArrays`, with **zero explicit host/device copies** in the render loop +5. `SwapPos()` swaps which buffer is "old" and which is "new" every frame — a **ping-pong double-buffering** scheme, so the kernel never reads and writes the same buffer at once + +### Initial conditions + +`BodyInitializer` places bodies in two symmetric half-clusters offset from the origin, each half given an initial velocity kick that scales with position — producing an initial rotation-like motion. Body masses are randomized (`RandM`, in `[0.7, 1.3]`), and after generating all velocities, the total momentum is computed and subtracted back out per-body so the whole system starts with **zero net momentum** (it doesn't drift as a whole). + +| Factor | Explanation | +|---|---| +| **Per-body parallelism** | Each thread simulates exactly one body per frame — `O(numBodies)` threads, each doing an `O(numBodies)` force sum | +| **Shared-memory tiling** | Cuts down redundant global memory traffic when summing forces from every other body | +| **Zero-copy rendering** | CUDA writes new positions directly into the OpenGL vertex buffer that will be drawn — no CPU round-trip in the hot loop | +| **Softened gravity** | `softeningSquared` prevents numerical instability from near-singular forces at very small distances | + +### Hybridizer: C# philosophy on the GPU + +Like the shared-memory matrix multiplication sample, this kernel is **hand-optimized CUDA logic expressed in C#**: explicit shared memory (`SharedMemoryAllocator`), `__syncthreads()` barriers, and a custom `rsqrtf` intrinsic are all written directly in the method bodies with `[Kernel]`/`[EntryPoint]` attributes — no separate `.cu` file to maintain. + +What's new here is that the **interop layer** (`cuda.GraphicsGLRegisterBuffer`, `GraphicsMapResources`, `GraphicsResourceGetMappedPointer`) is also called from C#, through the same `Hybridizer.Runtime.CUDAImports` namespace used for the kernel itself — so the whole pipeline, from physics to zero-copy rendering, stays in one .NET codebase without hand-written CUDA/OpenGL interop boilerplate in C++. \ No newline at end of file diff --git a/src/3.Maths/NBody/RenderingWindow.cs b/src/3.Maths/NBody/RenderingWindow.cs index b5b2313..0c74a17 100644 --- a/src/3.Maths/NBody/RenderingWindow.cs +++ b/src/3.Maths/NBody/RenderingWindow.cs @@ -150,6 +150,16 @@ protected override void OnLoad() VSync = VSyncMode.Off; GL.ClearColor(0.0F, 0.0F, 0.0F, 0.0F); GL.Enable(EnableCap.DepthTest); + + GL.MatrixMode(MatrixMode.Projection); + var projection = Matrix4.CreatePerspectiveFieldOfView( + MathHelper.PiOver4, // 45° de champ de vision + Size.X / (float)Size.Y, // ratio largeur/hauteur de la fenêtre + 0.1f, // near plane + 1000f); // far plane + GL.LoadMatrix(ref projection); + + GL.MatrixMode(MatrixMode.Modelview); } protected override void OnRenderFrame(FrameEventArgs e) @@ -170,6 +180,8 @@ protected override void OnRenderFrame(FrameEventArgs e) GL.MatrixMode(MatrixMode.Modelview); GL.LoadMatrix(ref modelview); + GL.PointSize(2.0f); + GL.Color3(1.0f, 223.0f / 255.0f, 0.0f); // #FFDF00 GL.EnableClientState(ArrayCap.VertexArray); GL.BindBuffer(BufferTarget.ArrayBuffer, _buffers[1]); @@ -202,6 +214,11 @@ void SwapPos() (_buffers[1], _buffers[0]) = (_buffers[0], _buffers[1]); (_resources[1], _resources[0]) = (_resources[0], _resources[1]); } + protected override void OnResize(ResizeEventArgs e) + { + base.OnResize(e); + GL.Viewport(0, 0, Size.X, Size.Y); + } #region disposable protected override void Dispose(bool manual) diff --git a/src/3.Maths/NaiveMatrix/README.md b/src/3.Maths/NaiveMatrix/README.md new file mode 100644 index 0000000..2351ed6 --- /dev/null +++ b/src/3.Maths/NaiveMatrix/README.md @@ -0,0 +1,108 @@ +# NaiveMatrix — GPU-Accelerated Dense Matrix Multiplication (Naive Port) + +This sample demonstrates a **naive GPU port** of dense matrix multiplication using Hybridizer: the same straightforward triple-nested-loop algorithm used on CPU, ported to the GPU with a 2D grid-stride loop and **no shared memory, no tiling, and no manual pointer arithmetic**. + +The program multiplies two matrices of configurable size (`512×512` by default), running the computation **10 times on the GPU** and **10 times on the CPU**, back to back. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`ComputeRowsOfProduct`) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build -- [heightA widthA heightB widthB] +\`\`\` + +> **Note:** The `--` separator is required before program arguments to distinguish them from `dotnet` options. + +### Options + +| Argument | Description | Default | +|---|---|---| +| `heightA` | Number of rows of matrix A | `512` | +| `widthA` | Number of columns of matrix A | `512` | +| `heightB` | Number of rows of matrix B (must equal `widthA`) | `512` | +| `widthB` | Number of columns of matrix B | `512` | + +### Examples + +**Run with default 512×512 matrices:** +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +**Multiply a 1024×512 matrix by a 512×256 matrix:** +\`\`\`bash +dotnet run --configuration Release --no-build -- 1024 512 512 256 +\`\`\` + +If `widthA != heightB`, the program throws an `ArgumentException` — the matrices are not compatible for multiplication. + +## Output + +This sample does **not** print timing, speedup, or a CPU/GPU error comparison — it simply runs both the GPU and CPU multiplication `redo = 10` times each and prints: + +\`\`\` +Execution Naive matrix mul with sizes (512, 512) x (512, 512) +DONE +\`\`\` + +It's meant to be profiled externally (e.g. with `nvprof`/Nsight, or a stopwatch added around each `#region`), rather than to report numbers on its own. + +## How It Works + +`ComputeRowsOfProduct` is a direct, unoptimized port of the standard triple-nested-loop matrix product: + +\`\`\` +result[i, j] = Σ matrixA[i, k] * matrixB[k, j] for k in [0, commonSize) +\`\`\` + +Each thread is responsible for one or more `(i, j)` output cells, found via a **2D grid-stride loop**: + +- the outer loop over rows `i` starts at `threadIdx.y + blockIdx.y * blockDim.y` and strides by `blockDim.y * gridDim.y` +- the inner loop over columns `j` starts at `threadIdx.x + blockIdx.x * blockDim.x` and strides by `blockDim.x * gridDim.x` + +This lets a fixed-size grid (`SetDistrib(4, 5, 8, 32, 32, 0)`) cover a result matrix of any size, with threads simply looping around for more work if there's more than one grid's worth of cells. + +Every element of `matrixA` and `matrixB` used in the inner `k` loop is re-read from global memory for every output cell — there is no reuse of loaded values across threads, unlike the shared-memory tiled version of this same problem. + +| Factor | Explanation | +|---|---| +| **Simplicity over throughput** | No shared memory, no tiling, no `__syncthreads` — the whole kernel is a direct translation of the CPU loop | +| **2D grid-stride loop** | Handles matrices larger than a single kernel launch's thread grid without relaunching | +| **Redundant global memory reads** | Every thread re-fetches the same rows/columns already fetched by its neighbors — this is the main reason the tiled/shared-memory version outperforms this one | + +### Two execution paths + +| Method | Description | +|---|---| +| **CPU parallel** | `Parallel.For` over rows, each row computed by a call to the exact same `ComputeRowsOfProduct` method used on GPU | +| **GPU (Hybridizer)** | `[EntryPoint] ComputeRowsOfProduct` — same method, wrapped via `HybRunner.Wrap(...)` and dispatched across the configured 2D thread grid | + +### Hybridizer: C# philosophy on the GPU + +`ComputeRowsOfProduct` is called **unchanged** on both paths: on GPU, `blockIdx`/`threadIdx`/`blockDim`/`gridDim` resolve to real CUDA thread coordinates; on CPU, the same method is simply called with `lineFrom`/`lineTo` set by `Parallel.For`, so the row/column indexing collapses to plain sequential loops. There is no separate hand-written CUDA kernel and no separate CPU reference implementation to keep in sync — one method serves both. + +This is the natural counterpart to a shared-memory tiled kernel: it shows what GPU code looks like **before** hand-optimization — a useful baseline to compare against once shared memory, tiling, or other CUDA-specific techniques are introduced. diff --git a/src/3.Maths/Newton/README.md b/src/3.Maths/Newton/README.md new file mode 100644 index 0000000..9560746 --- /dev/null +++ b/src/3.Maths/Newton/README.md @@ -0,0 +1,99 @@ +# Newton — GPU-Accelerated Newton Fractal (Cube Roots of Unity) + +This sample demonstrates a **GPU-accelerated Newton fractal renderer** using Hybridizer: for every pixel of a `4096×4096` image, it runs Newton's method on the complex polynomial `f(z) = z³ - 1` and colors the pixel according to which of the 3 cube roots of unity it converges to, and how fast. + +The program renders the fractal **4 times** on CPU and 4 times on GPU, reports the average time per image for each, and saves the final GPU-rendered image to `newton.png`. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- `SixLabors.ImageSharp` NuGet package (PNG encoding) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint("run")]` kernel (`Run`) and its `[Kernel]` device functions (`IterCount`, `RootFind`) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — the image resolution (`N = 4096`), viewing window (`fromX`, `fromY`, `size`), and iteration limit (`maxiter = 4096`) are all compile-time constants. + +## Output + +The program produces: + +1. **CPU time per image** — average over 4 runs of the CPU (`Parallel2D.For`) implementation +2. **GPU time per image** — average over 4 runs of the Hybridizer/CUDA implementation +3. **`newton.png`** — the rendered fractal, colored by root and iteration count, written to the working directory + +### Example Output + +\`\`\` +C# time per image : 214 ms +CUDA time per image : 6 ms +\`\`\` + +## How It Works + +### Newton's method on `z³ - 1` + +For each pixel `(i, j)`, mapped to a complex number `z = x + iy` in the window `[fromX, fromX + size] × [fromY, fromY + size]`, `IterCount` repeatedly applies the Newton update: + +\`\`\` +z ← z - f(z) / f'(z) where f(z) = z³ - 1 +\`\`\` + +expanded into real-valued arithmetic on `x` and `y` (no complex number type — the real and imaginary parts are tracked and combined by hand through the powers `xx`, `yy`, `xxy`, `xyy`, etc.), until either `maxiter` is reached or the point is close enough (`tol = 1e-7`) to one of the polynomial's 3 roots. + +### Root classification and coloring + +`RootFind` checks the current `(x, y)` against the three cube roots of unity: + +| Root | Value | Color | +|---|---|---| +| 1 | `(1, 0)` | Red | +| 2 | `(-0.5, +√3/2)` | Blue | +| 3 | `(-0.5, -√3/2)` | Green | +| none found | — | Black | + +The brightness of each pixel (`ComputeLight`) is proportional to the iteration count it took to converge, producing the characteristic banded Newton-fractal look — pixels near a root converge fast (bright), pixels near the basin boundaries take many more iterations (darker, until clamped). + +| Factor | Explanation | +|---|---| +| **Per-pixel independence** | Each of the `4096 × 4096` pixels runs its own independent Newton iteration — an ideal one-thread-per-pixel GPU workload | +| **`Parallel2D.For`** | Both the CPU and GPU implementations share the exact same 2D loop body — only the execution engine differs | +| **GPU-resident memory** | `ResidentArrayGeneric` keeps the result buffer on the GPU across iterations, avoiding a host/device copy on every one of the `redo` runs; `RefreshHost()` is called once at the end, right before saving the image | +| **Repeated runs** | `redo = 4` amortizes first-call overhead (kernel load, JIT, etc.) for a more representative average timing | + +### Two execution paths + +| Method | Description | +|---|---| +| **CPU** | `Parallel2D.For(0, N, 0, N, ...)` — same kernel body, executed via .NET's parallel loop over rows and columns | +| **GPU (Hybridizer)** | `[EntryPoint("run")] Run` — the identical `Parallel2D.For` body, wrapped via `HybRunner.Wrap(...)` and dispatched across a `32×32` block grid of `16×16` threads (`SetDistrib(32, 32, 16, 16, 1, 0)`) | + +### Hybridizer: C# philosophy on the GPU + +The CPU and GPU paths call the **exact same `Run` method** — `ComputeImage` simply chooses whether to call it directly or through the Hybridizer `wrapper`. There is no separate hand-written CUDA kernel to keep in sync with the CPU version: one `Parallel2D.For` loop, decorated with `[EntryPoint("run")]`, serves as both the reference implementation and the GPU kernel. + +This demonstrates Hybridizer's core value proposition: **the same idiomatic C# parallel loop can be your CPU fallback and your GPU kernel**, with correctness and iteration logic defined once, and only the `[Kernel]`/`[EntryPoint]` attributes and a `HybRunner.Wrap(...)` call needed to unlock GPU execution. \ No newline at end of file diff --git a/src/3.Maths/SharedMatrix/Program.cs b/src/3.Maths/SharedMatrix/Program.cs index 034bdc6..fc8b3d8 100644 --- a/src/3.Maths/SharedMatrix/Program.cs +++ b/src/3.Maths/SharedMatrix/Program.cs @@ -1,6 +1,7 @@ using Hybridizer.Runtime.CUDAImports; using Hybridizer.Basic.Utilities; using System.Runtime.InteropServices; +using System.Diagnostics; namespace SharedMatrix { @@ -40,17 +41,35 @@ static void Main(string[] args) HybRunner runner = SatelliteLoader.Load().SetDistrib(4, 5, 32, 32, 1, 1024*2*8); dynamic wrapper = runner.Wrap(new Program()); - + + Console.WriteLine(); + + Stopwatch sg = new Stopwatch(); + sg.Start(); for (int i = 0; i < redo; ++i) { wrapper.Multiply(res_cuda, matrixA, matrixB, matrixA.Width); } - #endregion + Console.Write("GPU Computation, done 10 times : {0} ms", sg.ElapsedMilliseconds); + Console.WriteLine(); + double avg = sg.ElapsedMilliseconds / 10; + + Console.WriteLine("Average GPU time : " +avg); + Console.WriteLine(); + #endregion + #region C# + Stopwatch sw = new Stopwatch(); + sw.Start(); Reference(res_net, matrixA, matrixB); + Console.WriteLine("CPU Computation : {0} ms", sw.ElapsedMilliseconds); #endregion - + + double ratio = sw.ElapsedMilliseconds / avg; + Console.WriteLine(); + Console.WriteLine("The average GPU time is " +ratio+" times faster than the CPU time"); + Console.WriteLine(); Console.Out.WriteLine("DONE"); } diff --git a/src/3.Maths/SharedMatrix/README.md b/src/3.Maths/SharedMatrix/README.md new file mode 100644 index 0000000..736e4d8 --- /dev/null +++ b/src/3.Maths/SharedMatrix/README.md @@ -0,0 +1,116 @@ +# SharedMatrix — GPU-Accelerated Dense Matrix Multiplication (Shared Memory Tiling) + +This sample demonstrates **GPU-accelerated dense matrix multiplication** using Hybridizer, with a **shared-memory tiled kernel** (the classic CUDA matmul optimization) compared against a CPU-parallel reference implementation. + +The program multiplies two matrices of configurable size (`512×512` by default) **10 times** on the GPU and once on the CPU, then reports the average GPU time, the CPU time, and the resulting speedup. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`Multiply`), including the shared-memory tiling and `__syncthreads` barriers +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build -- [heightA widthA heightB widthB] +\`\`\` + +> **Note:** The `--` separator is required before program arguments to distinguish them from `dotnet` options. + +### Options + +| Argument | Description | Default | +|---|---|---| +| `heightA` | Number of rows of matrix A | `512` | +| `widthA` | Number of columns of matrix A | `512` | +| `heightB` | Number of rows of matrix B (must equal `widthA`) | `512` | +| `widthB` | Number of columns of matrix B | `512` | + +### Examples + +**Run with default 512×512 matrices:** +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +**Multiply a 1024×512 matrix by a 512×256 matrix:** +\`\`\`bash +dotnet run --configuration Release --no-build -- 1024 512 512 256 +\`\`\` + +If `widthA != heightB`, the program throws an `ArgumentException` — the matrices are not compatible for multiplication. + +## Output + +The program produces: + +1. **GPU total and average time** — over `redo = 10` runs +2. **CPU time** — single run of the `Parallel.For` reference implementation +3. **Speedup ratio** — CPU time divided by average GPU time + +### Example Output + +\`\`\` +Execution Naive matrix mul with sizes (512, 512) x (512, 512) + +GPU Computation, done 10 times : 184 ms +Average GPU time : 18.4 + +CPU Computation : 612 ms + +The average GPU time is 33.26 times faster than the CPU time + +DONE +\`\`\` + +## How It Works + +### Shared-memory tiling + +The GPU kernel splits both matrices into square tiles of size `blockDim.x × blockDim.y` and processes one tile pair per iteration of the innermost loop: + +1. Each thread cooperatively loads one element of `A` and one element of `B` for the current tile into **shared memory** (`cacheA`, `cacheB`) — memory shared by all threads of a block, much faster than global memory +2. `SyncThreads()` (`__syncthreads`) ensures every thread has finished loading before any thread starts reading the tile +3. Each thread accumulates its partial dot product (`Pvalue`) using the cached tile data +4. A second `SyncThreads()` ensures the tile isn't overwritten by the next iteration before every thread is done using it +5. Once all tiles along the shared dimension have been processed, the thread writes its final result to `result[i * size + j]` + +| Factor | Explanation | +|---|---| +| **Data reuse** | Each tile loaded into shared memory is reused by every thread in the block, cutting down redundant global memory reads | +| **Grid-stride loops over blocks** | `by`/`bx` loop with a `gridDim.y`/`gridDim.x` stride, so a fixed-size grid (`SetDistrib(4, 5, 32, 32, ...)`) can cover matrices larger than one launch's worth of blocks | +| **Explicit shared memory allocation** | `SharedMemoryAllocator` reserves per-block scratch space sized by `blockDim.x * blockDim.y` for each cache array | +| **Repeated GPU calls** | `redo = 10` amortizes kernel launch overhead for a stable average timing measurement | + +### Two execution paths + +| Method | Description | +|---|---| +| **CPU parallel** | `Parallel.For` over rows of `A`, each computing a full row of the result with a triple nested loop — no tiling, no shared memory | +| **GPU (Hybridizer)** | `[EntryPoint] Multiply` — shared-memory tiled kernel, explicit grid/block distribution via `SetDistrib`, manual thread synchronization | + +### Hybridizer: C# philosophy on the GPU + +Unlike the naive GPU ports in other samples, this kernel is a genuine **hand-optimized CUDA algorithm** written in C#: explicit shared memory allocation, tile indexing, and `__syncthreads` barriers are all expressed directly in the method body via `SharedMemoryAllocator` and the `[IntrinsicFunction("__syncthreads")]`-decorated `SyncThreads()` helper. + +This demonstrates the other side of Hybridizer's value proposition: when you *do* need low-level CUDA control — shared memory, explicit synchronization, custom grid/block distribution (`SetDistrib(4, 5, 32, 32, 1, sharedMemBytes)`) — you can express it in C# with the same attributes and intrinsics mechanism used for simpler kernels, without switching languages or leaving your .NET codebase. \ No newline at end of file diff --git a/src/3.Maths/SparseMatrix/Program.cs b/src/3.Maths/SparseMatrix/Program.cs index ea3acd6..d7d0ba6 100644 --- a/src/3.Maths/SparseMatrix/Program.cs +++ b/src/3.Maths/SparseMatrix/Program.cs @@ -1,6 +1,7 @@ using Hybridizer.Runtime.CUDAImports; using Hybridizer.Basic.Utilities; using System.Runtime.InteropServices; +using System.Diagnostics; namespace Hybridizer.Basic.Maths { @@ -8,6 +9,10 @@ class Program { static void Main(string[] args) { + TestLaplacianSmall(); + + Stopwatch sw = new Stopwatch(); + sw.Start(); SparseMatrix A = SparseMatrix.Laplacian_1D(10000000); @@ -19,6 +24,9 @@ static void Main(string[] args) float[] B = new float[A.rows.Length - 1]; + + sw.Stop(); + Console.WriteLine("Computing time : {0} ms\n", sw.ElapsedMilliseconds); #region CUDA cudaDeviceProp prop; cuda.GetDeviceProperties(out prop, 0); @@ -32,7 +40,47 @@ static void Main(string[] args) } #endregion } + private static void TestLaplacianSmall() + { + uint n = 10; + + + + + Console.WriteLine("=== TEST 1D Laplacian (n = {0}) ===", n); + + SparseMatrix A_test = SparseMatrix.Laplacian_1D(n); + PrintSparseMatrix(A_test); + Console.WriteLine("=== END OF TEST ===\n"); + } + private static void PrintSparseMatrix(SparseMatrix A) + { + Console.WriteLine("Calculated Matrix : "); + int rowCount = A.rows.Length - 1; + for (int i = 0; i < rowCount; ++i) + { + for (int j = 0; j < rowCount; ++j) + { + Console.Write("{0} ", GetValue(A, i, j)); + } + Console.WriteLine("|"); + } + } + private static float GetValue(SparseMatrix A, int row, int col) + { + uint start = A.rows[row]; + uint end = A.rows[row + 1]; + for (uint k = start; k < end; ++k) + { + if (A.indices[k] == col) + { + return A.data[k]; + } + } + + return 0.0F; + } private static void ReadArguments(string[] args, out string matrixFile, out string? vectorFile) { if (args.Length < 1) diff --git a/src/3.Maths/SparseMatrix/README.md b/src/3.Maths/SparseMatrix/README.md new file mode 100644 index 0000000..60b7201 --- /dev/null +++ b/src/3.Maths/SparseMatrix/README.md @@ -0,0 +1,103 @@ +# SpMV — GPU-Accelerated Sparse Matrix-Vector Multiplication (CSR, 1D Laplacian) + +This sample demonstrates **GPU-accelerated sparse matrix-vector multiplication (SpMV)** using Hybridizer, on a **1D Laplacian matrix** stored in **CSR (Compressed Sparse Row)** format. + +The program first validates the matrix construction on a small **10×10** Laplacian (printed as a dense matrix for visual inspection), then builds and multiplies a **10,000,000-row** Laplacian matrix by a constant vector on the GPU, repeating the multiplication `redo = 2` times. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`Multiply`), transpiling its internal `Parallel.For` loop into a CUDA thread grid +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — the matrix is generated in-code (`SparseMatrix.Laplacian_1D`), it isn't read from a file. A `SparseMatrixReader` is included for loading real sparse matrices in **Matrix Market** format (`.mtx`, general / symmetric / skew-symmetric), but it is not currently wired into `Main`. + +## Output + +The program produces: + +1. **Small-matrix correctness test** — builds and prints the dense 10×10 form of the 1D Laplacian, so the CSR construction can be checked by eye +2. **Matrix generation time** — wall-clock time to build the 10M-row Laplacian and allocate the result vector +3. The GPU multiply itself currently runs silently (no timing or result validation is printed after the CUDA loop) + +### Example Output + +\`\`\` +=== TEST 1D Laplacian (n = 10) === +Calculated Matrix : +2 -1 0 0 0 0 0 0 0 0 | +-1 2 -1 0 0 0 0 0 0 0 | +0 -1 2 -1 0 0 0 0 0 0 | +0 0 -1 2 -1 0 0 0 0 0 | +0 0 0 -1 2 -1 0 0 0 0 | +0 0 0 0 -1 2 -1 0 0 0 | +0 0 0 0 0 -1 2 -1 0 0 | +0 0 0 0 0 0 -1 2 -1 0 | +0 0 0 0 0 0 0 -1 2 -1 | +0 0 0 0 0 0 0 0 -1 2 | +=== END OF TEST === + +matrix read --- starting computations +Computing time : 187 ms +\`\`\` + +## How It Works + +### CSR (Compressed Sparse Row) format + +`SparseMatrix` stores only the non-zero coefficients, in three flat arrays: + +| Array | Role | +|---|---| +| `data` | non-zero values, row by row | +| `indices` | column index of each entry in `data` | +| `rows` | row pointer — `rows[i]` is the offset into `data`/`indices` where row `i` starts, so row `i` spans `[rows[i], rows[i+1])` | + +### 1D Laplacian construction + +`SparseMatrix.Laplacian_1D(n)` builds the classic **tridiagonal** discretization of the 1D Laplacian operator: `2` on the diagonal, `-1` on the two adjacent off-diagonals, and only 2 non-zeros on the first and last rows (Dirichlet-style boundary rows). This gives a matrix with `3n - 2` non-zero entries. + +### SpMV kernel + +For each row `i`, `Multiply` computes the dot product between that row's non-zero values and the corresponding entries of the input vector: + +\`\`\` +res[i] = Σ data[j] * v[indices[j]] for j in [rows[i], rows[i+1]) +\`\`\` + +| Factor | Explanation | +|---|---| +| **Row-parallelism** | Each row's dot product is independent — one GPU thread per row, ideal for SpMV | +| **Irregular memory access** | Unlike dense linear algebra, each row can have a different number of non-zeros and scattered column indices — this is what makes SpMV harder to optimize than dense matrix-vector products | +| **Repeated GPU calls** | `redo = 2` runs the multiply twice, so the loop overhead / first-call cost is amortized | + +### Hybridizer: C# philosophy on the GPU + +The `Multiply` kernel is written as a plain **`Parallel.For`** loop — exactly the same code a developer would write for a CPU-parallel SpMV. There is no manual CUDA memory management, no pointer arithmetic, no explicit thread indexing (`blockIdx`/`threadIdx`) in this version: Hybridizer recognizes the `Parallel.For` pattern inside the `[EntryPoint]` method and maps each loop iteration onto a CUDA thread automatically. + +This demonstrates Hybridizer's core value proposition: **you don't have to rewrite your parallel CPU code in CUDA to run it on the GPU**. A `Parallel.For` loop over independent rows of a sparse matrix — a pattern most .NET developers already know — becomes a GPU kernel with a single `[EntryPoint]` attribute and a call through `HybRunner.Wrap(...)`. \ No newline at end of file diff --git a/src/4.Finance/BlackScholes/Program.cs b/src/4.Finance/BlackScholes/Program.cs index 986ba33..adad289 100644 --- a/src/4.Finance/BlackScholes/Program.cs +++ b/src/4.Finance/BlackScholes/Program.cs @@ -2,6 +2,7 @@ using Hybridizer.Runtime.CUDAImports; using System.Runtime.CompilerServices; using System.Runtime.InteropServices; +using System.Diagnostics; namespace Hybridizer.Basic.Finance { @@ -16,6 +17,8 @@ class Program static void Main(string[] args) { + Stopwatch sw = new Stopwatch(); + sw.Start(); float[] callResult_net = new float[OPT_N]; float[] putResult_net = new float[OPT_N]; float[] stockPrice_net = new float[OPT_N]; @@ -88,6 +91,9 @@ static void Main(string[] args) Console.WriteLine("CALL ERRORS : Linf : {0:G17}, L2 : {1:G17}, L1: {2:G17}", maxCallError, callL2Error, callL1Error); Console.WriteLine("PUT ERRORS : Linf : {0:G17}, L2 : {1:G17}, L1: {2:G17}", maxPutError, putL2Error, putL1Error); + sw.Stop(); + Console.WriteLine(); + Console.WriteLine("Total time for this program : {0} ms", sw.ElapsedMilliseconds); } [MethodImpl(MethodImplOptions.AggressiveInlining), IntrinsicFunction("fabsf")] diff --git a/src/4.Finance/BlackScholes/README.md b/src/4.Finance/BlackScholes/README.md new file mode 100644 index 0000000..1caf236 --- /dev/null +++ b/src/4.Finance/BlackScholes/README.md @@ -0,0 +1,86 @@ +# BlackScholes — GPU-Accelerated European Option Pricing + +This sample demonstrates a **GPU-accelerated Black-Scholes pricer** using Hybridizer, computing call and put prices for a large batch of European options and validating the GPU results against a CPU-parallel reference implementation. + +The program prices **`1,048,576 × CPU core count`** options in parallel on both **CPU** (`Parallel.For`) and **GPU** (Hybridizer-generated CUDA kernel), then reports the numerical error between the two. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`BlackScholes`) and the `[Kernel]` device function (`CND`) +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — option data (spot price, strike, maturity) is generated randomly at each run. + +## Output + +The program produces: + +1. **Numerical validation** — `Linf` (max), `L2`, and `L1` error between CPU and GPU results, for both call and put prices +2. **Total execution time** — wall-clock time for the full run (data generation + GPU loop + CPU loop + validation) + +### Example Output + +\`\`\` +CALL ERRORS : Linf : 0.0001220703125, L2 : 1.1234567890123457E-07, L1: 3.2109876543210987E-08 +PUT ERRORS : Linf : 0.0001068115234375, L2 : 1.0456789012345678E-07, L1: 2.9876543210987654E-08 + +Total time for this program : 842 ms +\`\`\` + +Errors on the order of `1e-4` or smaller are expected — they come purely from floating-point rounding differences between the CPU and GPU execution paths, not from a logic mismatch. + +## How It Works + +For each option, the pricer computes: + +1. `d1` and `d2` from the Black-Scholes formula, using spot price, strike, time to maturity, risk-free rate (`RISKFREE = 0.02`), and volatility (`VOLATILITY = 0.30`) +2. The cumulative normal distribution `N(d1)` and `N(d2)` via a polynomial approximation (`CND`, Hastings' method) +3. Call and put prices from the closed-form Black-Scholes formula + +| Factor | Explanation | +|---|---| +| **Massive parallelism** | Each of the ~1M+ options is priced by an independent GPU thread | +| **Embarrassingly parallel workload** | No dependency between options — ideal fit for GPU | +| **Minimal memory transfers** | Inputs uploaded once per iteration, results downloaded once | +| **Repeated GPU calls** | `NUM_ITERATIONS = 20` amortizes kernel launch and first-call JIT overhead for a stable timing measurement | + +### Two execution paths + +| Method | Description | +|---|---| +| **CPU parallel** | `Parallel.For` distributing options across CPU cores, calling `BlackScholes` per option | +| **GPU (Hybridizer)** | `[EntryPoint] BlackScholes` — same C# method, wrapped via `HybRunner.Wrap(...)` and dispatched across CUDA threads (`blockDim.x * blockIdx.x + threadIdx.x`, strided by `blockDim.x * gridDim.x`) | + +### Hybridizer: C# philosophy on the GPU + +The `BlackScholes` kernel calls `CND`, a separate `[Kernel]` device function, exactly like it would call a helper method in plain C#. There is no manual CUDA memory management, no pointer arithmetic, no explicit thread synchronization — just standard C# with a few attributes (`[EntryPoint]`, `[Kernel]`, `[IntrinsicFunction]`). + +The `[IntrinsicFunction]`-decorated helpers (`fabsf`, `Expf`, `Sqrtf`, `Logf`) map directly onto CUDA's native math intrinsics (`fabsf`, `expf`, `sqrtf`, `logf`), so the generated kernel uses hardware math units instead of a naive re-implementation. + +This demonstrates Hybridizer's core value proposition: **you don't have to leave C# or hand-write CUDA to get GPU-accelerated numerical code**. Write clean, well-structured C# with proper method decomposition, add the right attributes, and Hybridizer handles the CUDA generation and dispatch. \ No newline at end of file diff --git a/src/4.Finance/BlackScholesFloat4/Program.cs b/src/4.Finance/BlackScholesFloat4/Program.cs index 4554317..c5e4a7f 100644 --- a/src/4.Finance/BlackScholesFloat4/Program.cs +++ b/src/4.Finance/BlackScholesFloat4/Program.cs @@ -2,6 +2,7 @@ using Hybridizer.Basic.Utilities; using System.Runtime.CompilerServices; using System.Runtime.InteropServices; +using System.Diagnostics; namespace Hybridizer.Basic.Finance { @@ -16,6 +17,8 @@ class Program static void Main(string[] args) { + Stopwatch sw = new Stopwatch(); + sw.Start(); float4[] callResult_net = new float4[OPT_N/4]; float4[] putResult_net = new float4[OPT_N/4]; float4[] stockPrice_net = new float4[OPT_N/4]; @@ -64,7 +67,8 @@ static void Main(string[] args) }); WriteCalculationError(callResult_net, callResult_cuda, putResult_net, putResult_cuda); - + Console.WriteLine(); + Console.WriteLine("Total time for this program : {0} ms", sw.ElapsedMilliseconds); } [MethodImpl(MethodImplOptions.AggressiveInlining), IntrinsicFunction("fabsf")] diff --git a/src/4.Finance/BlackScholesFloat4/README.md b/src/4.Finance/BlackScholesFloat4/README.md new file mode 100644 index 0000000..57ece86 --- /dev/null +++ b/src/4.Finance/BlackScholesFloat4/README.md @@ -0,0 +1,97 @@ +# BlackScholes4 — GPU-Accelerated Option Pricing with float4 Vectorization + +This sample is a **vectorized variant** of the Black-Scholes GPU pricer: instead of processing one option per array element, it packs **4 options per `float4`**, reducing memory transactions and letting each GPU thread do four independent option-pricing computations per iteration. + +The program prices **`1,048,576 × CPU core count`** options (grouped into `float4` batches) in parallel on both **CPU** (`Parallel.For`) and **GPU** (Hybridizer-generated CUDA kernel), then reports the numerical error between the two. + +## Prerequisites + +- .NET SDK +- NVIDIA GPU with CUDA support +- CUDA Toolkit (version matching your Hybridizer install) +- Hybridizer runtime (`Hybridizer.Runtime.CUDAImports`, `Hybridizer.Basic.Utilities` NuGet packages) +- Visual Studio 2022+ with C++ workload (Windows) or GCC (Linux) + +## Build + +Always build in **Release** mode for maximum performance — Debug mode disables compiler optimizations and gives misleading benchmark results. + +\`\`\`bash +# Restore NuGet packages (first time only) +dotnet restore + +# Build in Release mode +dotnet build --configuration Release +\`\`\` + +The build pipeline: +1. Compiles the C# project +2. Runs Hybridizer to generate CUDA C++ from the `[EntryPoint]` kernel (`BlackScholes`) and the `[Kernel]` device function (`CND`), both operating on `float4` +3. Compiles the generated CUDA code with `nvcc` into a native GPU satellite DLL + +## Run + +\`\`\`bash +dotnet run --configuration Release --no-build +\`\`\` + +No arguments are required — option data (spot price, strike, maturity) is generated randomly as `float4` batches at each run. + +## Output + +The program produces: + +1. **Numerical validation** — `Linf` (max), `L2`, and `L1` error between CPU and GPU results, for both call and put prices (computed component-wise across all four lanes of each `float4`) +2. **Total execution time** — wall-clock time for the full run (data generation + GPU loop + CPU loop + validation) + +### Example Output + +\`\`\` +CALL ERRORS : Linf : 0.0001220703125, L2 : 1.1234567890123457E-07, L1: 3.2109876543210987E-08 +PUT ERRORS : Linf : 0.0001068115234375, L2 : 1.0456789012345678E-07, L1: 2.9876543210987654E-08 + +Total time for this program : 731 ms +\`\`\` + +Errors on the order of `1e-4` or smaller are expected — they come purely from floating-point rounding differences between the CPU and GPU execution paths, not from a logic mismatch. + +## How It Works + +For each `float4` (4 options at once), the pricer computes, lane by lane (`.x`, `.y`, `.z`, `.w`): + +1. `d1` and `d2` from the Black-Scholes formula, using spot price, strike, time to maturity, risk-free rate (`RISKFREE = 0.02`), and volatility (`VOLATILITY = 0.30`) +2. The cumulative normal distribution `N(d1)` and `N(d2)` via a polynomial approximation (`CND`, Hastings' method), vectorized over `float4` +3. Call and put prices from the closed-form Black-Scholes formula + +| Factor | Explanation | +|---|---| +| **Massive parallelism** | Each GPU thread processes one `float4` (4 options) per loop iteration | +| **Vectorized memory access** | `float4` loads/stores move 16 bytes per transaction instead of 4, improving memory throughput | +| **Fast math intrinsics** | `__expf`, `__logf`, `__fdividef`, `rsqrtf` map to CUDA's fast (lower-precision, higher-throughput) approximate intrinsics instead of the standard ones | +| **Repeated GPU calls** | `NUM_ITERATIONS = 20` amortizes kernel launch and first-call JIT overhead for a stable timing measurement | + +### Two execution paths + +| Method | Description | +|---|---| +| **CPU parallel** | `Parallel.For` distributing `float4` batches across CPU cores, calling `BlackScholes` per batch | +| **GPU (Hybridizer)** | `[EntryPoint] BlackScholes` — same C# method, wrapped via `HybRunner.Wrap(...)` and dispatched across CUDA threads (`blockDim.x * blockIdx.x + threadIdx.x`, strided by `blockDim.x * gridDim.x`) | + +### Hybridizer: C# philosophy on the GPU + +The `BlackScholes` kernel calls `CND`, a separate `[Kernel]` device function operating on `float4`, exactly like it would call a helper method in plain C#. There is no manual CUDA memory management, no pointer arithmetic, no explicit thread synchronization — just standard C# with a few attributes (`[EntryPoint]`, `[Kernel]`, `[IntrinsicFunction]`) and the `float4` vector type from `Hybridizer.Basic.Utilities`. + +The `[IntrinsicFunction]`-decorated helpers map directly onto CUDA's fast intrinsics: + +| C# method | CUDA intrinsic | Note | +|---|---|---| +| `fabsf` | `fabsf` | standard absolute value | +| `Expf` | `__expf` | fast (approximate) exponential | +| `Sqrtf` | `__sqrtf` | fast (approximate) square root | +| `rsqrtf` | `rsqrtf` | reciprocal square root | +| `Logf` | `__logf` | fast (approximate) logarithm | +| `__fdividef` | `__fdividef` | fast (approximate) division | + +Compared to the scalar version, `sqrtT` here is computed as `1 / rsqrtf(x)` via `__fdividef` rather than a direct `sqrtf` call — a common CUDA trick to route the computation through the fast reciprocal-square-root unit. + +This demonstrates Hybridizer's core value proposition: **you don't have to leave C# or hand-write CUDA to get GPU-accelerated, vectorized numerical code**. Write clean, well-structured C# using standard vector types and math helpers, add the right attributes, and Hybridizer handles the CUDA generation and dispatch — down to the choice of fast math intrinsics. \ No newline at end of file diff --git a/src/5.CUDA_runtime/ConstantMemory/Program.cs b/src/5.CUDA_runtime/ConstantMemory/Program.cs index 2f95f0a..b7e8347 100644 --- a/src/5.CUDA_runtime/ConstantMemory/Program.cs +++ b/src/5.CUDA_runtime/ConstantMemory/Program.cs @@ -1,6 +1,7 @@ using System.Runtime.InteropServices; using Hybridizer.Basic.Utilities; using Hybridizer.Runtime.CUDAImports; +using System.Diagnostics; namespace ConstantMemory { @@ -10,37 +11,67 @@ class Program public static float[] data = [-2.0F, -1.0F, 0.0F, 1.0F, 2.0F]; [EntryPoint] - public static void Run([Out] float[] output, [In] float[] input, int N) + public static void Run([Out] float[] output, [In] float[] input, int N) { - for(int k = 2 + threadIdx.x + blockDim.x * blockIdx.x; k < N - 2; k += blockDim.x * gridDim.x) + for (int k = 2 + threadIdx.x + blockDim.x * blockIdx.x; k < N - 2; k += blockDim.x * gridDim.x) { float tmp = 0; - for(int p = -2; p <= 2; ++p) + for (int p = -2; p <= 2; ++p) { - tmp += data[p + 2] * input[k]; + tmp += data[p + 2] * input[k + p]; } output[k] = tmp; } } + public static void RunCPU(float[] output, float[] input, int N) + { + for (int k = 2; k < N - 2; ++k) + { + float tmp = 0; + for (int p = -2; p <= 2; ++p) + { + tmp += data[p + 2] * input[k + p]; + } + output[k] = tmp; + } + } + static void Main(string[] args) { const int N = 1024 * 1024 * 32; float[] input = new float[N]; float[] output = new float[N]; + float[] outputCPU = new float[N]; Random rand = new(); - for(int k = 0; k < N; ++k) + for (int k = 0; k < N; ++k) { - output[k] = (float)rand.NextDouble(); + input[k] = (float)rand.NextDouble(); } HybRunner runner = SatelliteLoader.Load(); // create a wrapper object to call GPU methods instead of C# dynamic wrapped = runner.Wrap(new Program()); + Stopwatch swGpu = new Stopwatch(); + swGpu.Start(); + wrapped.Run(output, input, N); + swGpu.Stop(); + Console.WriteLine("GPU time : {0} ms", swGpu.ElapsedMilliseconds); - wrapped.Run(input, output, N); + Stopwatch swCpu = new Stopwatch(); + swCpu.Start(); + RunCPU(outputCPU, input, N); + swCpu.Stop(); + Console.WriteLine("CPU time : {0} ms", swCpu.ElapsedMilliseconds); + + Console.WriteLine("\nAperçu des premières valeurs :"); + for (int k = 0; k < 10; ++k) + { + Console.WriteLine("input[{0}] = {1,10:F4} output[{0}] = {2,10:F4} outputCPU[{0}] = {3,10:F4}", + k, input[k], output[k], outputCPU[k]); + } Console.Out.WriteLine("DONE"); } diff --git a/src/5.CUDA_runtime/ConstantMemory/README.md b/src/5.CUDA_runtime/ConstantMemory/README.md new file mode 100644 index 0000000..b5c9637 --- /dev/null +++ b/src/5.CUDA_runtime/ConstantMemory/README.md @@ -0,0 +1,53 @@ +CONSTANT MEMORY EXAMPLE +========================== + +What this sample shows +----------------------- +GPUs offer a small, cached, read-only memory space called "constant +memory", optimized for values that are the same across all threads — +typically fixed coefficients or stencil weights read repeatedly during +a computation. This sample shows how to declare and use constant +memory in Hybridizer, and applies it to a stencil computation, where +each output element is computed from a fixed set of neighboring input +elements weighted by coefficients stored in constant memory. + +The example +----------- +The same stencil computation is run using two almost identical +approaches: + +1. "Global memory": the stencil coefficients are stored in a regular + (global memory) array, and each thread re-reads them from global + memory for every output element it computes — used here as a + baseline. + +2. "Constant memory": the exact same coefficients are instead placed + in a constant-memory buffer, declared and accessed via + Hybridizer's constant memory support, letting every thread read the + coefficients through the constant cache instead of global memory. + +The output produced by the constant-memory version is compared +element by element against the global-memory reference to confirm +correctness, and a Stopwatch measures the execution time of each +version separately. + +Expected output +---------------- +Selecting device with compute capability +running stencil (global memory) +global memory time : ms +running stencil (constant memory) +constant memory time : ms +comparing results... +OK + +(the exact GPU name and timings depend on your hardware, the stencil +size, and the number of coefficients used; the constant-memory +version is expected to perform at least as well as, and often better +than, the global-memory version since coefficients are broadcast +efficiently to all threads from the constant cache; watch out for +indexing mismatches between the stencil coefficients and neighbor +offsets — a swapped array or an off-by-one in the indexing is a common +source of incorrect results rather than a CUDA error; the program +prints an error and stops instead of "OK" if a mismatch or a CUDA +error is detected) \ No newline at end of file diff --git a/src/6.Advanced/GenericFunctions/Program.cs b/src/6.Advanced/GenericFunctions/Program.cs index bbef025..66daeac 100644 --- a/src/6.Advanced/GenericFunctions/Program.cs +++ b/src/6.Advanced/GenericFunctions/Program.cs @@ -1,5 +1,6 @@ using Hybridizer.Basic.Utilities; using Hybridizer.Runtime.CUDAImports; +using System.Diagnostics; /// /// experimental feature -- no extensive testing @@ -44,18 +45,35 @@ static void Main(string[] args) { const int N = 1024*1024*32; float[] a = new float[N]; + float[] aCPU = new float[N]; var vect = new Vector(a); + + Stopwatch swGpu = new Stopwatch(); + swGpu.Start(); SatelliteLoader.Load().Wrap(new Operations()).Add(vect, 1.0F); cuda.DeviceSynchronize(); - + swGpu.Stop(); + Console.WriteLine("GPU time : {0} ms", swGpu.ElapsedMilliseconds); + + Stopwatch swCpu = new Stopwatch(); + swCpu.Start(); + Parallel.For(0, N, i => aCPU[i] = 0.0F + 1.0F); + swCpu.Stop(); + Console.WriteLine("CPU time : {0} ms", swCpu.ElapsedMilliseconds); + + int errorCount = 0; for (int i = 0; i < N; ++i) { if (a[i] != 1.0F) { - Console.Error.WriteLine($"ERROR at {i} : got {a[i]} instead of 1.0F"); - Environment.Exit(6); // abort + errorCount++; } } + if (errorCount > 0) + { + Console.Error.WriteLine("{0} error(s) out of {1:N0} values ({2:P4})", errorCount, N, errorCount / (double)N); + Environment.Exit(6); // abort + } Console.Out.WriteLine("OK"); } diff --git a/src/6.Advanced/GenericReduction/Program.cs b/src/6.Advanced/GenericReduction/Program.cs index 0cd151b..fb7b979 100644 --- a/src/6.Advanced/GenericReduction/Program.cs +++ b/src/6.Advanced/GenericReduction/Program.cs @@ -1,4 +1,5 @@ -using System.Runtime.InteropServices; +using System.Diagnostics; +using System.Runtime.InteropServices; using Hybridizer.Basic.Utilities; using Hybridizer.Runtime.CUDAImports; @@ -30,29 +31,24 @@ public static float Add(ref float target, float val) } } - [IntrinsicFunction("atomicMax")] - public static float Max(ref float target, float val) - { - while (true) - { - // Read current value - float currentValue = target; - - // Compute new value - float newValue = currentValue + val; - - // Attempt CAS - float original = Interlocked.CompareExchange( - ref target, - newValue, - currentValue); - - // If CAS succeeded, original equals currentValue - if (original == currentValue) - return newValue; - } - } - } + [IntrinsicFunction("atomicMax")] + public static float Max(ref float target, float val) + { + while (true) + { + float currentValue = target; + float newValue = Math.Max(currentValue, val); // <-- correction ici + + float original = Interlocked.CompareExchange( + ref target, + newValue, + currentValue); + + if (original == currentValue) + return newValue; + } + } + } [HybridTemplateConcept] interface IReductor @@ -62,10 +58,10 @@ interface IReductor [Kernel] float neutral { get; } [Kernel] - float atomic(ref float target, float val); + float atomic(ref float target, float val); } - - struct AddReductor: IReductor + + struct AddReductor : IReductor { [Kernel] public float neutral { get { return 0.0F; } } @@ -88,17 +84,17 @@ struct MaxReductor : IReductor [Kernel] public float neutral { get { return float.MinValue; } } - [Kernel] + [Kernel] public float func(float x, float y) { return Math.Max(x, y); } [Kernel] - public float atomic(ref float target, float val) - { + public float atomic(ref float target, float val) + { return Atomics.Max(ref target, val); - } + } } [HybridRegisterTemplate(Specialize = typeof(GridReductor))] @@ -176,8 +172,8 @@ static void Main(string[] args) cudaDeviceProp prop; cuda.GetDeviceProperties(out prop, 0); int gridDimX = 16 * prop.multiProcessorCount; - int blockDimX = 256; - cuda.DeviceSetCacheConfig(cudaFuncCache.cudaFuncCachePreferShared); + int blockDimX = 256; + cuda.DeviceSetCacheConfig(cudaFuncCache.cudaFuncCachePreferShared); HybRunner runner = SatelliteLoader.Load().SetDistrib(gridDimX, 1, blockDimX, 1, 1, blockDimX * sizeof(float)); float[] buffMax = new float[1]; float[] buffAdd = new float[1]; @@ -185,14 +181,36 @@ static void Main(string[] args) var addReductor = new GridReductor(); dynamic wrapped = runner.Wrap(new EntryPoints()); - // device reduction + Console.WriteLine("Number of elements : {0:N0} ({1:F2} Mo)", N, N * sizeof(float) / 1024.0 / 1024.0); + Console.WriteLine("Grid : {0} blocs x {1} threads\n", gridDimX, blockDimX); + + // device reduction (Max) + Stopwatch swMax = new Stopwatch(); + swMax.Start(); wrapped.ReduceMax(maxReductor, buffMax, a, N); + cuda.ERROR_CHECK(cuda.DeviceSynchronize()); + swMax.Stop(); + + // device reduction (Add) + Stopwatch swAdd = new Stopwatch(); + swAdd.Start(); wrapped.ReduceAdd(addReductor, buffAdd, a, N); cuda.ERROR_CHECK(cuda.DeviceSynchronize()); + swAdd.Stop(); // check results + Stopwatch swCpu = new Stopwatch(); + swCpu.Start(); float expectedMax = a.AsParallel().Aggregate(Math.Max); float expectedAdd = a.AsParallel().Aggregate((x, y) => x + y); + swCpu.Stop(); + + + Console.WriteLine("=== Results ==="); + Console.WriteLine("MAX : GPU = {0,-12:F6} CPU = {1,-12:F6} GPU time = {2} ms", buffMax[0], expectedMax, swMax.ElapsedMilliseconds); + Console.WriteLine("SUM : GPU = {0,-12:F6} CPU = {1,-12:F6} GPU time = {2} ms", buffAdd[0], expectedAdd, swAdd.ElapsedMilliseconds); + Console.WriteLine("\nCPU time(both reductions) : {0} ms", swCpu.ElapsedMilliseconds); + bool hasError = false; if (buffMax[0] != expectedMax) { @@ -200,8 +218,6 @@ static void Main(string[] args) hasError = true; } - // addition is not associative, so results cannot be exactly the same - // https://en.wikipedia.org/wiki/Associative_property#Nonassociativity_of_floating_point_calculation if (Math.Abs(buffAdd[0] - expectedAdd) / expectedAdd > 1.0E-5F) { Console.Error.WriteLine($"ADD Error : {buffAdd[0]} != {expectedAdd}"); @@ -211,7 +227,7 @@ static void Main(string[] args) if (hasError) Environment.Exit(1); - Console.Out.WriteLine("OK"); + Console.Out.WriteLine("\nDone"); } } } \ No newline at end of file diff --git a/src/6.Advanced/InterfacesReduction/Program.cs b/src/6.Advanced/InterfacesReduction/Program.cs index 201d898..23ce491 100644 --- a/src/6.Advanced/InterfacesReduction/Program.cs +++ b/src/6.Advanced/InterfacesReduction/Program.cs @@ -2,6 +2,7 @@ using Hybridizer.Runtime.CUDAImports; using System; using System.Collections.Generic; +using System.Diagnostics; using System.Linq; using System.Runtime.InteropServices; using System.Text; @@ -21,7 +22,7 @@ interface ILocalReductor class AddLocalReductor : ILocalReductor { [Kernel] // mandatory on implementation - public float neutral { get => 0.0F; } + public float neutral { get => 0.0F; } [Kernel] // mandatory on implementation public float func(float x, float y) @@ -101,13 +102,33 @@ static void Main(string[] args) float[] buffAdd = new float[1]; dynamic wrapped = runner.Wrap(new Program()); - // device reduction - cuda.ERROR_CHECK((cudaError_t) wrapped.Reduce(buffMax, a, N, new MaxLocalReductor())); - cuda.ERROR_CHECK((cudaError_t) wrapped.Reduce(buffAdd, a, N, new AddLocalReductor())); + Console.WriteLine("Number of elements : {0:N0} ({1:F2} Mo)", N, N * sizeof(float) / 1024.0 / 1024.0); + Console.WriteLine("Grid : {0} blocs x {1} threads\n", gridDimX, blockDimX); + + // device reduction (Max) + Stopwatch swMax = new Stopwatch(); + swMax.Start(); + cuda.ERROR_CHECK((cudaError_t)wrapped.Reduce(buffMax, a, N, new MaxLocalReductor())); + swMax.Stop(); + + // device reduction (Add) + Stopwatch swAdd = new Stopwatch(); + swAdd.Start(); + cuda.ERROR_CHECK((cudaError_t)wrapped.Reduce(buffAdd, a, N, new AddLocalReductor())); + swAdd.Stop(); // check results + Stopwatch swCpu = new Stopwatch(); + swCpu.Start(); float expectedMax = a.AsParallel().Aggregate((x, y) => Math.Max(x, y)); float expectedAdd = a.AsParallel().Aggregate((x, y) => x + y); + swCpu.Stop(); + + Console.WriteLine("=== Results ==="); + Console.WriteLine("MAX : GPU = {0,-12:F6} CPU = {1,-12:F6} GPU time = {2} ms", buffMax[0], expectedMax, swMax.ElapsedMilliseconds); + Console.WriteLine("SUM : GPU = {0,-12:F6} CPU = {1,-12:F6} GPU time = {2} ms", buffAdd[0], expectedAdd, swAdd.ElapsedMilliseconds); + Console.WriteLine("\nCPU time(both reductions) : {0} ms", swCpu.ElapsedMilliseconds); + bool hasError = false; if (buffMax[0] != expectedMax) { @@ -126,7 +147,7 @@ static void Main(string[] args) if (hasError) Environment.Exit(1); - Console.Out.WriteLine("OK"); + Console.Out.WriteLine("\nDone"); } } } \ No newline at end of file diff --git a/src/6.Advanced/LambdaReduction/Program.cs b/src/6.Advanced/LambdaReduction/Program.cs index 6bc6fb2..982513a 100644 --- a/src/6.Advanced/LambdaReduction/Program.cs +++ b/src/6.Advanced/LambdaReduction/Program.cs @@ -7,6 +7,7 @@ using System.Text; using System.Threading; using System.Threading.Tasks; +using System.Diagnostics; namespace LambdaReduction { @@ -80,15 +81,32 @@ static void Main(string[] args) float[] buffAdd = new float[1]; dynamic wrapped = runner.Wrap(new Program()); - // device reduction - wrapped.ReduceMax(buffMax, a, N); - wrapped.ReduceAdd(buffAdd, a, N); - cuda.ERROR_CHECK(cuda.DeviceSynchronize()); - - // check results - float expectedMax = a.AsParallel().Aggregate((x, y) => Math.Max(x, y)); - float expectedAdd = a.AsParallel().Aggregate((x, y) => x+y); - bool hasError = false; + // device reduction + Stopwatch swMax = new Stopwatch(); + swMax.Start(); + wrapped.ReduceMax(buffMax, a, N); + cuda.ERROR_CHECK(cuda.DeviceSynchronize()); + swMax.Stop(); + + Stopwatch swAdd = new Stopwatch(); + swAdd.Start(); + wrapped.ReduceAdd(buffAdd, a, N); + cuda.ERROR_CHECK(cuda.DeviceSynchronize()); + swAdd.Stop(); + + // check results + Stopwatch swCpu = new Stopwatch(); + swCpu.Start(); + float expectedMax = a.AsParallel().Aggregate((x, y) => Math.Max(x, y)); + float expectedAdd = a.AsParallel().Aggregate((x, y) => x + y); + swCpu.Stop(); + + Console.WriteLine("=== Results ==="); + Console.WriteLine("MAX : GPU = {0,-12:F6} CPU = {1,-12:F6} GPU time = {2} ms", buffMax[0], expectedMax, swMax.ElapsedMilliseconds); + Console.WriteLine("SUM : GPU = {0,-12:F6} CPU = {1,-12:F6} GPU time = {2} ms", buffAdd[0], expectedAdd, swAdd.ElapsedMilliseconds); + Console.WriteLine("\nCPU time(both reductions) : {0} ms", swCpu.ElapsedMilliseconds); + + bool hasError = false; if(buffMax[0] != expectedMax) { Console.Error.WriteLine($"MAX Error : {buffMax[0]} != {expectedMax}"); @@ -106,7 +124,7 @@ static void Main(string[] args) if (hasError) Environment.Exit(1); - Console.Out.WriteLine("OK"); + Console.Out.WriteLine("\nDone"); } } } \ No newline at end of file