feat: wmma - #46
Merged
Merged
Conversation
* 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>
* 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>
* 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>
* 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.
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>
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
1 participant
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.
implement #44