feat: nvcuda::wmma bindings for tensor cores - #45
Merged
Merged
Conversation
Adds C# bindings for the nvcuda::wmma C++ API in src/CUDARuntime/wmma.cs,
exposing tensor-core matrix multiply-accumulate to Hybridizer-transcoded
kernels. Phase 1 scope: half input on (16,16,16)/(32,8,16)/(8,32,16) with
{half,float} accumulators, tf32 on (16,16,8) with float accumulator, and
double on (8,8,4) with double accumulator. Operations: fill_fragment,
load_matrix_sync, store_matrix_sync, mma_sync.
Adds IntrinsicFieldSubscriptAttribute in src/API.cs, used by the WMMA
fragment indexers to emit (frag)->x[i] in C++. Requires matching
transcoder support (separate PR in hybridizer-software-suite).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
JeWaVe
added a commit
that referenced
this pull request
May 19, 2026
* feat!: add gridimz to hybrunner BREAKING CHANGE (#18) * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * restore setdistrib * rebase development (#40) * feat: in and out attributes management (#21) * feat!: add gridimz to hybrunner BREAKING CHANGE (#18) * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * restore setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * feat: add NVRTC support for CUDA 13.x and GetCUBIN API (#23) * feat: add NVRTC support for CUDA 13.x and GetCUBIN API Add NVRTC bindings for CUDA 13.0 and 13.1 with Linux/Windows platform-specific implementations. Add GetCUBIN API for native GPU binary output (available since CUDA 11.1). * fix: use portable DllImport name for NVRTC 13.x Use "nvrtc" as the DLL name so the .NET runtime resolves to libnvrtc.so on Linux and nvrtc.dll on Windows, instead of maintaining separate platform-specific classes. * fix: use versioned platform-specific DLL names for NVRTC 13.x Use full paths on Linux (/usr/local/cuda-13.x/lib64/libnvrtc.so) and versioned DLL names on Windows (nvrtc64_130_0.dll) with runtime platform dispatch, matching the cuda runtime pattern. * feat: AlternateNameAttribute for type rename in C++ output (#192) (#24) C# permits a non-generic class Foo and a generic class Foo<T> to coexist in the same namespace; C++ does not. The transpiler currently emits both under the same identifier, producing invalid C++. Working around it by using different namespaces or intrinsic implementations is awkward. Add a new attribute, [AlternateName("X")], that overrides the emitted C++ identifier for the decorated class/struct while leaving the C# type name untouched. The transpiler still generates the type's full C++ definition; only the identifier differs (unlike IntrinsicTypeAttribute, which delegates the implementation to user code). Inherits IntrinsicAttribute so the existing Flavor scoping works for free: [AlternateName("FooCuda", flavor: 1)] applies only to CUDA. This is the user-facing half. The internal mirror in Transcoder.Disassembly and the rename plumbing in CTypeWriter live in the main repo PR; the GUID literal here must stay in lockstep with that mirror. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * feat: cuda 13.2 + ai found issues (#41) * fix: review-driven cleanup across cuda/nvrtc wrappers + cuda 13.2 support Closes: - #25 StreamAddCallback wrapper recursed infinitely in every cuda impl — routed through cudaStreamAddCallback (the PInvoke shim). - #26 cudaGetMipmappedArrayLevel was bound to EntryPoint "cudaFreeMipmappedArray" — fixed across all impls. - #27 EntryPoint "cudaOccupancyMaxActiveBlocksPerMultiprocessor " had a trailing space — stripped. - #28 DriverAPI hard-coded nvcuda.dll. The dllPath const becomes a synthetic name and a NativeLibrary.SetDllImportResolver (registered via reflection so netstandard2.0 still compiles) maps it to libcuda.so.1 on Linux. - #29 ADD_TO_PARAM_BUFFER had a duplicated typeof(byte) branch — removed. - #30 nvrtc was unmapped for CUDA 11.0 / 11.4 / 12.0 / 12.4 / 12.6 — new impls added (Linux + Windows variants), cloned from the 13.0 wrapper. - #31 dead `cudaVersion = "80"` fallback in nvrtc.GetCudaVersion — removed. - #32 SetCudaVersion was a no-op once the cuda type initializer had run — selection logic extracted into a SelectInstance helper that the setter calls to rebind both the cuda runtime instance and the nvrtc backend. - #33 StringArrayMarshal(empty[]) returned a non-null Ptr — empty arrays now mirror the null path and surface as IntPtr.Zero. - #34 StringArrayMarshal crashed on null entries; the partially-constructed finalizer would then crash the process — null entries map to IntPtr.Zero slots and the disposer guards every GCHandle with IsAllocated. - #35 nvrtc Get(ProgramLog|PTX|CUBIN) freed their pinned GCHandle outside any try/finally — wrapped. - #36 Linux DllImport paths were hard-coded to /usr/local/cuda-X.Y/lib64/... — switched to sonames (libcudart.so.<MAJOR>, libnvrtc.so.<MAJOR>) so ldconfig resolves them and the same wrapper rides ABI compat across minors. - #37 cudaDeviceProp_130/131.surfaceAlignment lacked `public`, so StructConvert silently dropped it — made public. - #38 unused CUDA_DLLS dictionary in cuda.cs — removed. - #39 JittedModule.GetEntryPoint(string) discarded the CUresult — added a TryGetEntryPoint(string, out CUfunction) overload that returns it. Also: full CUDA 13.2 support — new cuda-13.2.cs / nvrtc-13.2.cs wrappers, SetCudaVersion whitelist entry, dispatcher cases, and HYBRIDIZER_CUDA_VERSION_132 build constant. Tests: tests/BugRegressionTests.cs covers each case (reflection + native smoke calls); run against the local CUDA 13.2 install to validate soname resolution end-to-end. * test: skip CUDA-loading regression tests when libcudart is absent Catches DllNotFoundException so CI without the toolkit doesn't fail. * fix: better getcustomattribute * feat: nvcuda::wmma bindings for tensor cores (#45) Adds C# bindings for the nvcuda::wmma C++ API in src/CUDARuntime/wmma.cs, exposing tensor-core matrix multiply-accumulate to Hybridizer-transcoded kernels. Phase 1 scope: half input on (16,16,16)/(32,8,16)/(8,32,16) with {half,float} accumulators, tf32 on (16,16,8) with float accumulator, and double on (8,8,4) with double accumulator. Operations: fill_fragment, load_matrix_sync, store_matrix_sync, mma_sync. Adds IntrinsicFieldSubscriptAttribute in src/API.cs, used by the WMMA fragment indexers to emit (frag)->x[i] in C++. Requires matching transcoder support (separate PR in hybridizer-software-suite). Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
JeWaVe
added a commit
that referenced
this pull request
May 21, 2026
* feat!: add gridimz to hybrunner BREAKING CHANGE (#18) * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * restore setdistrib * rebase development (#40) * feat: in and out attributes management (#21) * feat!: add gridimz to hybrunner BREAKING CHANGE (#18) * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * Feat: in out attributes (#20) * initial attempt to handle in/out attributes : make memcpy optional * feat: in-out parameters. Modify Hybrunner * feat: in-out attribute. Add UT * feat: in/out attributes , add more UT * feat: in-ou attributes, fix UT * feat: in-out attributes, fix setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * fix : remove duplicated setdistrib * restore setdistrib --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> * feat: add NVRTC support for CUDA 13.x and GetCUBIN API (#23) * feat: add NVRTC support for CUDA 13.x and GetCUBIN API Add NVRTC bindings for CUDA 13.0 and 13.1 with Linux/Windows platform-specific implementations. Add GetCUBIN API for native GPU binary output (available since CUDA 11.1). * fix: use portable DllImport name for NVRTC 13.x Use "nvrtc" as the DLL name so the .NET runtime resolves to libnvrtc.so on Linux and nvrtc.dll on Windows, instead of maintaining separate platform-specific classes. * fix: use versioned platform-specific DLL names for NVRTC 13.x Use full paths on Linux (/usr/local/cuda-13.x/lib64/libnvrtc.so) and versioned DLL names on Windows (nvrtc64_130_0.dll) with runtime platform dispatch, matching the cuda runtime pattern. * feat: AlternateNameAttribute for type rename in C++ output (#192) (#24) C# permits a non-generic class Foo and a generic class Foo<T> to coexist in the same namespace; C++ does not. The transpiler currently emits both under the same identifier, producing invalid C++. Working around it by using different namespaces or intrinsic implementations is awkward. Add a new attribute, [AlternateName("X")], that overrides the emitted C++ identifier for the decorated class/struct while leaving the C# type name untouched. The transpiler still generates the type's full C++ definition; only the identifier differs (unlike IntrinsicTypeAttribute, which delegates the implementation to user code). Inherits IntrinsicAttribute so the existing Flavor scoping works for free: [AlternateName("FooCuda", flavor: 1)] applies only to CUDA. This is the user-facing half. The internal mirror in Transcoder.Disassembly and the rename plumbing in CTypeWriter live in the main repo PR; the GUID literal here must stay in lockstep with that mirror. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * feat: cuda 13.2 + ai found issues (#41) * fix: review-driven cleanup across cuda/nvrtc wrappers + cuda 13.2 support Closes: - #25 StreamAddCallback wrapper recursed infinitely in every cuda impl — routed through cudaStreamAddCallback (the PInvoke shim). - #26 cudaGetMipmappedArrayLevel was bound to EntryPoint "cudaFreeMipmappedArray" — fixed across all impls. - #27 EntryPoint "cudaOccupancyMaxActiveBlocksPerMultiprocessor " had a trailing space — stripped. - #28 DriverAPI hard-coded nvcuda.dll. The dllPath const becomes a synthetic name and a NativeLibrary.SetDllImportResolver (registered via reflection so netstandard2.0 still compiles) maps it to libcuda.so.1 on Linux. - #29 ADD_TO_PARAM_BUFFER had a duplicated typeof(byte) branch — removed. - #30 nvrtc was unmapped for CUDA 11.0 / 11.4 / 12.0 / 12.4 / 12.6 — new impls added (Linux + Windows variants), cloned from the 13.0 wrapper. - #31 dead `cudaVersion = "80"` fallback in nvrtc.GetCudaVersion — removed. - #32 SetCudaVersion was a no-op once the cuda type initializer had run — selection logic extracted into a SelectInstance helper that the setter calls to rebind both the cuda runtime instance and the nvrtc backend. - #33 StringArrayMarshal(empty[]) returned a non-null Ptr — empty arrays now mirror the null path and surface as IntPtr.Zero. - #34 StringArrayMarshal crashed on null entries; the partially-constructed finalizer would then crash the process — null entries map to IntPtr.Zero slots and the disposer guards every GCHandle with IsAllocated. - #35 nvrtc Get(ProgramLog|PTX|CUBIN) freed their pinned GCHandle outside any try/finally — wrapped. - #36 Linux DllImport paths were hard-coded to /usr/local/cuda-X.Y/lib64/... — switched to sonames (libcudart.so.<MAJOR>, libnvrtc.so.<MAJOR>) so ldconfig resolves them and the same wrapper rides ABI compat across minors. - #37 cudaDeviceProp_130/131.surfaceAlignment lacked `public`, so StructConvert silently dropped it — made public. - #38 unused CUDA_DLLS dictionary in cuda.cs — removed. - #39 JittedModule.GetEntryPoint(string) discarded the CUresult — added a TryGetEntryPoint(string, out CUfunction) overload that returns it. Also: full CUDA 13.2 support — new cuda-13.2.cs / nvrtc-13.2.cs wrappers, SetCudaVersion whitelist entry, dispatcher cases, and HYBRIDIZER_CUDA_VERSION_132 build constant. Tests: tests/BugRegressionTests.cs covers each case (reflection + native smoke calls); run against the local CUDA 13.2 install to validate soname resolution end-to-end. * test: skip CUDA-loading regression tests when libcudart is absent Catches DllNotFoundException so CI without the toolkit doesn't fail. * fix: better getcustomattribute * feat: nvcuda::wmma bindings for tensor cores (#45) Adds C# bindings for the nvcuda::wmma C++ API in src/CUDARuntime/wmma.cs, exposing tensor-core matrix multiply-accumulate to Hybridizer-transcoded kernels. Phase 1 scope: half input on (16,16,16)/(32,8,16)/(8,32,16) with {half,float} accumulators, tf32 on (16,16,8) with float accumulator, and double on (8,8,4) with double accumulator. Operations: fill_fragment, load_matrix_sync, store_matrix_sync, mma_sync. Adds IntrinsicFieldSubscriptAttribute in src/API.cs, used by the WMMA fragment indexers to emit (frag)->x[i] in C++. Requires matching transcoder support (separate PR in hybridizer-software-suite). Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * feat: pipeline, ldmatrix, and mma PTX bindings for Hybridizer (#47) Adds three CUDARuntime modules in the same shape as wmma.cs: * pipeline.cs — binds <cuda_pipeline.h>'s __pipeline_memcpy_async, __pipeline_commit, and __pipeline_wait_prior. Lets Hybridizer-transcoded kernels build producer/consumer cp.async pipelines for shmem tile staging (sm_80+; the header fallback is safe on older arches). * ldmatrix.cs — binds ldmatrix.sync.aligned.m8n8.x{1,2,4}[.trans].b16 plus a generic-to-shared address cvta helper. The underlying inline-PTX wrappers live in namespace hybridizer::ldmatrix inside the runtime header (see companion PR in hybridizer-software-suite). * mma.cs — binds the canonical fp16/fp32 tensor-core MMA mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 with C and D folded into a single ref-float operand set for in-place accumulate (CUTLASS Mma_sm80 pattern). Companion runtime helpers in namespace hybridizer::mma in the suite. All three follow the wmma.cs conventions: free-function `op` class, [IntrinsicFunction] on each binding, primitive-typed args only. The pipeline binding carries an [IntrinsicIncludeCUDA("<cuda_pipeline.h>")]; ldmatrix and mma reach for hybridizer::* helpers that ship in hybridizer.cuda.cuh and therefore need no extra include attribute. Phase 1 scope is the minimum needed to write a CUTLASS-shaped GEMM (cp.async double-buffered staging + ldmatrix.x4(.trans) loads + mma.m16n8k16). bf16/tf32/fp8 mma shapes and ldmatrix.{x1,x2,trans-x2} variants land in follow-up PRs once we have callers. Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Régis Portalez <regis.portalez@gmail.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Adds C# bindings for the nvcuda::wmma C++ API in src/CUDARuntime/wmma.cs, exposing tensor-core matrix multiply-accumulate to Hybridizer-transcoded kernels. Phase 1 scope: half input on (16,16,16)/(32,8,16)/(8,32,16) with {half,float} accumulators, tf32 on (16,16,8) with float accumulator, and double on (8,8,4) with double accumulator. Operations: fill_fragment, load_matrix_sync, store_matrix_sync, mma_sync.
Adds IntrinsicFieldSubscriptAttribute in src/API.cs, used by the WMMA fragment indexers to emit (frag)->x[i] in C++. Requires matching transcoder support (separate PR in hybridizer-software-suite).