docs: improve CPU matmul and GPU buffer reuse - #720
Conversation
github-actions
Bot
changed the title
|
Note Other AI code review bot(s) detectedCodeRabbit has detected other AI code review bot(s) in this pull request and will avoid duplicating their findings in the review comments. This may lead to a less comprehensive review. Summary by CodeRabbit
✏️ Tip: You can customize this high-level summary in your review settings. WalkthroughAdds GPU‑resident benchmarking and execution paths, thread‑safe GPU buffer pooling across CUDA/HIP/OpenCL, AiDotNetEngine GPU context helpers, activation input→output refactor, zero‑copy tensor constructors, OpenBLAS packaging and native CPU BLAS loader, CPU/GPU benchmark suites and tooling, and multiple benchmark CLI/config enhancements. Changes
Sequence Diagram(s)sequenceDiagram
participant CLI as CLI
participant Harness as GpuResidentQuickHarness
participant Engine as AiDotNetEngine
participant Pool as GpuBufferPool
participant Backend as IDirectGpuBackend
participant Device as GPU
CLI->>Harness: Run(warmup, iterations)
Harness->>Engine: BeginGpuContext(options)
Engine-->>Harness: GpuExecutionContext
Harness->>Engine: Create data & Upload tensors
Engine->>Pool: TryRent(size)
alt pool hit
Pool-->>Engine: pooled buffer
Engine->>Backend: Copy host -> pooled buffer
else pool miss
Engine->>Backend: Allocate device buffer
end
Harness->>Backend: Enqueue MatMul/Add/Activation ops
Backend->>Device: Execute kernels
Device-->>Backend: Complete
Backend-->>Harness: Signal / timings
Harness->>Pool: Return buffers (Dispose -> Return)
Harness->>CLI: Print per-op averages
Estimated code review effort🎯 4 (Complex) | ⏱️ ~45 minutes Possibly related issues
Possibly related PRs
Poem
🚥 Pre-merge checks | ✅ 1 | ❌ 2❌ Failed checks (2 warnings)
✅ Passed checks (1 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing touches
🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
|
🤖 PR Title Auto-Fixed Your PR title was automatically updated to follow Conventional Commits format. Original title: New title: Detected type: Valid types and their effects:
If the detected type is incorrect, you can manually edit the PR title. |
There was a problem hiding this comment.
Pull request overview
This PR introduces CPU matrix multiplication optimizations via BLAS integration and blocked algorithms, adds GPU buffer pooling to reduce allocation overhead, and expands benchmark coverage with CPU/GPU comparisons against TorchSharp, TensorFlow.NET, and ML.NET.
Changes:
- Adds BLAS provider for high-performance CPU matrix multiplication
- Implements blocked matrix multiplication with cache-friendly access patterns
- Introduces GPU buffer pooling to reuse allocations and reduce memory fragmentation
- Adds Im2Col-based Conv2D optimization path for CPU tensors
- Expands benchmark suite with separate CPU/GPU comparison benchmarks and GPU-resident harness
Reviewed changes
Copilot reviewed 20 out of 20 changed files in this pull request and generated 6 comments.
Show a summary per file
| File | Description |
|---|---|
CompilerAttributePolyfills.cs |
Adds NotNullWhenAttribute polyfill for older frameworks |
MatrixBase.cs |
Integrates BLAS and blocked matmul paths in CPU matrix multiply |
MatrixMultiplyHelper.cs |
New helper providing BLAS integration, blocked matmul, and path selection |
BlasProvider.cs |
New BLAS library loader with P/Invoke to OpenBLAS/MKL |
DirectGpuTensorEngine.cs |
Refactors activation functions to support separate input/output buffers |
OpenClBackend.cs |
Adds buffer pooling and kernel work group size clamping |
LossKernels.cs |
Renames reduction kernels to avoid naming conflicts |
HipBackend.cs |
Adds buffer pooling support for HIP backend |
GpuBufferPool.cs |
New GPU buffer pool implementation for reuse |
CudaBackend.cs |
Adds buffer pooling support for CUDA backend |
CpuEngine.cs |
Routes matrix multiply through optimized paths and adds Im2Col Conv2D |
AiDotNetEngine.cs |
Adds GPU context management API |
BENCHMARKS.md |
Documents CPU/GPU benchmark split |
*ComparisonBenchmarks.cs |
New/updated benchmark files for CPU/GPU comparisons |
Program.cs |
Adds GPU harness command-line options |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
| var result = HipNativeBindings.hipMemcpy( | ||
| pooled.Handle, | ||
| handle.AddrOfPinnedObject(), | ||
| size, | ||
| HipMemcpyKind.HostToDevice); |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 6
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (5)
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
563-608: Fix GPU-memory leaks on allocation/init failure (pooled + non-pooled paths).
IfhipMemcpy/hipMemsetthrows, the rented buffer (or newlyhipMalloc’d pointer) can be leaked because nothing returns it to the pool or frees it. This is especially important for transient failures (OOM / device-lost / invalid params).Proposed fix (ensure cleanup on exceptions)
public IGpuBuffer AllocateBuffer(float[] data) { IntPtr devicePtr = IntPtr.Zero; var size = (UIntPtr)(data.Length * sizeof(float)); - if (_bufferPool.TryRent(data.Length, out var pooled) && pooled != null) - { - GCHandle handle = GCHandle.Alloc(data, GCHandleType.Pinned); - try - { - var result = HipNativeBindings.hipMemcpy( - pooled.Handle, - handle.AddrOfPinnedObject(), - size, - HipMemcpyKind.HostToDevice); - HipNativeBindings.CheckError(result, "hipMemcpy H2D"); - } - finally - { - handle.Free(); - } - - return pooled; - } + if (_bufferPool.TryRent(data.Length, out var pooled) && pooled != null) + { + try + { + GCHandle handle = GCHandle.Alloc(data, GCHandleType.Pinned); + try + { + var result = HipNativeBindings.hipMemcpy( + pooled.Handle, + handle.AddrOfPinnedObject(), + size, + HipMemcpyKind.HostToDevice); + HipNativeBindings.CheckError(result, "hipMemcpy H2D"); + } + finally + { + handle.Free(); + } + + return pooled; + } + catch + { + // Ensure the rented buffer isn't leaked if initialization fails. + pooled.Dispose(); // returns to pool or releases (depending on pool policy) + throw; + } + } var allocResult = HipNativeBindings.hipMalloc(ref devicePtr, size); HipNativeBindings.CheckError(allocResult, "hipMalloc"); // Copy data to device GCHandle allocHandle = GCHandle.Alloc(data, GCHandleType.Pinned); try { var copyResult = HipNativeBindings.hipMemcpy( devicePtr, allocHandle.AddrOfPinnedObject(), size, HipMemcpyKind.HostToDevice); HipNativeBindings.CheckError(copyResult, "hipMemcpy H2D"); } finally { allocHandle.Free(); } return new HipGpuBuffer(devicePtr, data.Length, _bufferPool.Return); } public IGpuBuffer AllocateBuffer(int size) { IntPtr devicePtr = IntPtr.Zero; var sizeBytes = (UIntPtr)(size * sizeof(float)); - if (_bufferPool.TryRent(size, out var pooled) && pooled != null) - { - var zeroResult = HipNativeBindings.hipMemset(pooled.Handle, 0, sizeBytes); - HipNativeBindings.CheckError(zeroResult, "hipMemset"); - return pooled; - } + if (_bufferPool.TryRent(size, out var pooled) && pooled != null) + { + try + { + var zeroResult = HipNativeBindings.hipMemset(pooled.Handle, 0, sizeBytes); + HipNativeBindings.CheckError(zeroResult, "hipMemset"); + return pooled; + } + catch + { + pooled.Dispose(); + throw; + } + } var allocResult = HipNativeBindings.hipMalloc(ref devicePtr, sizeBytes); HipNativeBindings.CheckError(allocResult, "hipMalloc"); - // Zero-initialize - var memsetResult = HipNativeBindings.hipMemset(devicePtr, 0, sizeBytes); - HipNativeBindings.CheckError(memsetResult, "hipMemset"); + // Zero-initialize + try + { + var memsetResult = HipNativeBindings.hipMemset(devicePtr, 0, sizeBytes); + HipNativeBindings.CheckError(memsetResult, "hipMemset"); + } + catch + { + // Avoid leaking the freshly allocated device pointer if initialization fails. + _ = HipNativeBindings.hipFree(devicePtr); + throw; + } return new HipGpuBuffer(devicePtr, size, _bufferPool.Return); }Also applies to: 610-630
src/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs (2)
1515-1536: DisposeoutputBufferon exceptions inActivationGpu
outputBufferis allocated on Line 1522, but ifCopy/ApplyGpuActivationthrows, it leaks. Wrap with exception-safe disposal.Proposed fix
public IGpuTensor<T> ActivationGpu<T>(IGpuTensor<T> input, FusedActivationType activation) { @@ // Allocate output buffer int size = input.ElementCount; var outputBuffer = backend.AllocateBuffer(size); - if (activation == FusedActivationType.None) - { - // Preserve previous behavior: output is a copy of input. - backend.Copy(input.Buffer, outputBuffer, size); - } - else - { - ApplyGpuActivation(backend, input.Buffer, outputBuffer, size, activation); - } - - // Return new GPU tensor - return new GpuTensor<T>(backend, outputBuffer, input.Shape, GpuTensorRole.Activation, ownsBuffer: true); + try + { + if (activation == FusedActivationType.None) + { + // Preserve previous behavior: output is a copy of input. + backend.Copy(input.Buffer, outputBuffer, size); + } + else + { + ApplyGpuActivation(backend, input.Buffer, outputBuffer, size, activation); + } + + // Return new GPU tensor + return new GpuTensor<T>(backend, outputBuffer, input.Shape, GpuTensorRole.Activation, ownsBuffer: true); + } + catch + { + outputBuffer.Dispose(); + throw; + } }
7003-7062:SumAxisGputranspose workaround looks right; makeoutputBufferexception-safeThe new axis==0 path (Line 7035+) correctly transposes to reuse
SumAxis, butoutputBuffer(Line 7033) will leak ifTranspose/SumAxisthrows.Proposed fix
- var outputBuffer = backend.AllocateBuffer(outputSize); + var outputBuffer = backend.AllocateBuffer(outputSize); - if (axis == 0) - { - if (innerSize == 1) - { - backend.SumAxis(input.Buffer, outputBuffer, 1, outerSize); - } - else - { - IGpuBuffer? transposedBuffer = null; - try - { - transposedBuffer = backend.AllocateBuffer(outerSize * innerSize); - backend.Transpose(input.Buffer, transposedBuffer, outerSize, innerSize); - backend.SumAxis(transposedBuffer, outputBuffer, innerSize, outerSize); - } - finally - { - transposedBuffer?.Dispose(); - } - } - } - else - { - backend.SumAxis(input.Buffer, outputBuffer, outerSize, innerSize); - } - - return new GpuTensor<T>(backend, outputBuffer, outputShape, GpuTensorRole.Activation, ownsBuffer: true); + try + { + if (axis == 0) + { + if (innerSize == 1) + { + backend.SumAxis(input.Buffer, outputBuffer, 1, outerSize); + } + else + { + IGpuBuffer? transposedBuffer = null; + try + { + transposedBuffer = backend.AllocateBuffer(outerSize * innerSize); + backend.Transpose(input.Buffer, transposedBuffer, outerSize, innerSize); + backend.SumAxis(transposedBuffer, outputBuffer, innerSize, outerSize); + } + finally + { + transposedBuffer?.Dispose(); + } + } + } + else + { + backend.SumAxis(input.Buffer, outputBuffer, outerSize, innerSize); + } + + return new GpuTensor<T>(backend, outputBuffer, outputShape, GpuTensorRole.Activation, ownsBuffer: true); + } + catch + { + outputBuffer.Dispose(); + throw; + }src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (1)
7664-7670:Dispose()should be exception-safe before setting_disposed(pool dispose can currently strand resources)
_disposed = trueis set before_bufferPool.Dispose(). If_bufferPool.Dispose()throws, the method exits early, other CUDA resources won’t be released, and the finalizer won’t retry because_disposedis already true.Consider guarding disposal so cleanup continues even if pool disposal fails.
Proposed fix
public void Dispose() { if (_disposed) return; - _disposed = true; - _bufferPool.Dispose(); + try + { + _bufferPool.Dispose(); + } + catch + { + // Best-effort cleanup; avoid leaking other CUDA resources if pool disposal fails. + } + _disposed = true; if (_cublasHandle != IntPtr.Zero) { CuBlasNative.cublasDestroy(_cublasHandle); _cublasHandle = IntPtr.Zero; }AiDotNetBenchmarkTests/TensorFlowComparisonBenchmarks.cs (1)
300-305: TensorFlow benchmark returns result but doesn't dispose it.The TensorFlow benchmark methods return the
Tensorflow.Tensorresult without disposing it, which could lead to memory pressure during repeated benchmark iterations. Consider disposing the result or using the consumer pattern used in TorchSharp benchmarks.🐛 Proposed fix: Dispose TensorFlow results
[Benchmark] [Arguments(256)] [Arguments(512)] - public Tensorflow.Tensor TensorFlow_MatMul(int size) + public void TensorFlow_MatMul(int size) { - var result = tf.matmul(_tfMatricesA[size], _tfMatricesB[size]); - result.numpy(); - return result; + using var result = tf.matmul(_tfMatricesA[size], _tfMatricesB[size]); + result.numpy(); }
🤖 Fix all issues with AI agents
In `@src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs`:
- Around line 182-198: Add a defensive null check for the fallback parameter in
the static WithGpuContext<TResult>(Func<GpuExecutionContext, TResult> func,
Func<TResult> fallback, GpuExecutionOptions? options = null) method: at the
start of AiDotNet.Tensors.Engines.AiDotNetEngine.WithGpuContext validate
fallback is not null and throw an ArgumentNullException(nameof(fallback)) if it
is, mirroring the validation style used in the Current setter and ensuring
clearer error messages when DirectGpuTensorEngine is not active.
In `@src/AiDotNet.Tensors/Engines/CpuEngine.cs`:
- Around line 3474-3607: TryConv2DIm2Col can OOM by allocating colMatrix sized
colCols*colRows per batch; add a guard in TryConv2DIm2Col before allocating
colMatrix (and optionally kernelMatrix) that estimates byte size (use sizeof 4
for float / 8 for double based on typeof(T)) and if the allocation would exceed
a safe threshold (e.g., 100MB or a configured cap) return false so the engine
falls back to the direct/safer convolution path; alternatively implement tiling
by iterating over output tiles (split along colRows or colCols), allocating a
smaller colTile buffer and calling MultiplyMatrixCore for each tile — reference
colMatrix, kernelMatrix, colRows, colCols, MultiplyMatrixCore and ensure any
early-return or tiling keeps batch semantics and uses allowParallel as before.
In `@src/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cs`:
- Around line 71-87: Dispose has a race with concurrent Return: _disposed may be
set after Return checked it, letting a buffer be added to a bucket after you've
drained that bucket causing a leak; to fix, either document that the pool must
not be used during disposal, or make Dispose defensive by setting _disposed
first, then repeatedly drain all _buckets until a full pass removes nothing (or
swap _buckets with an empty collection and drain the old one), and ensure
buffer.Release is called for every buffer removed; reference symbols to change:
Dispose, Return, _disposed, _buckets, bucket.Buffers.TryTake, and
buffer.Release.
In `@src/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs`:
- Around line 1323-1376: The out-of-place ApplyGpuActivation( IDirectGpuBackend
backend, IGpuBuffer input, IGpuBuffer output, int size, FusedActivationType
activation ) must not leave output uninitialized when activation ==
FusedActivationType.None; copy input to output (use the backend's buffer
copy/memcpy method or an elementwise copy routine) in that branch and add a
default: that throws an ArgumentOutOfRangeException for unknown
FusedActivationType values to avoid silent no-ops in future; leave the in-place
overload ApplyGpuActivation(backend, buffer, size, activation) unchanged.
In `@src/AiDotNet.Tensors/Helpers/BlasProvider.cs`:
- Around line 139-156: When an explicit AIDOTNET_BLAS_PATH is provided and
TryLoadNativeLibrary succeeds but TryLoadSymbols fails, the code currently
leaves _libraryHandle loaded; modify the explicit-path branch in the
BlasProvider (the explicitPath, TryLoadNativeLibrary, _libraryHandle,
TryLoadSymbols flow) to call FreeNativeLibrary(_libraryHandle) and set
_libraryHandle = IntPtr.Zero before returning/continuing (so the handle is
released just like in the candidate loop).
In `@src/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cs`:
- Around line 19-93: TryGemm<T> is allocating arrays (especially cArray) even
when the BLAS path won’t be used; add an early type guard at the start of
TryGemm<T> to return false unless typeof(T) is float or double (the only types
supported by TryGemmFromArray), so you skip all subsequent
TryGetArraySegment/ToArray/new T[] work; apply the same guard to the other
similar helper block (the second TryGemm occurrence referencing TryGemmFromArray
and TryGetArraySegment) to avoid wasted allocations.
🧹 Nitpick comments (16)
src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/Kernels/LossKernels.cs (1)
259-262: Unused variablefeatureIdx.The variable
featureIdxis declared but never used. It appears to be leftover from a different parallelization strategy. Consider removing it to avoid confusion.__kernel void triplet_loss_backward( ... { const int tripletIdx = get_global_id(0); - const int featureIdx = get_local_id(0); if (tripletIdx >= batchSize) return;src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs (1)
151-164: Consider documenting context lifecycle.If
GpuExecutionContextimplementsIDisposable, callers need to know they're responsible for disposing the returned context. Consider adding a remark to the XML documentation./// <remarks> /// The caller is responsible for disposing the returned context when finished. /// </remarks>src/AiDotNet.Tensors/Engines/CpuEngine.cs (1)
4385-4430: Batched matmul: consider skippingParallel.ForwhenbatchSize == 1.
Right nowallowParallel:falseis always used insideParallel.For, which can leave performance on the table for single-batch workloads (and still paysParallel.Foroverhead).Optional perf tweak for batchSize==1
int matrixSizeA = m * n; int matrixSizeResult = m * p; + if (batchSize == 1) + { + MultiplyMatrixCore(numOps, aData, 0, bData, 0, resultData, 0, m, n, p, allowParallel: true); + return result; + } + Parallel.For(0, batchSize, batch => { int aOffset = batch * matrixSizeA; int resultOffset = batch * matrixSizeResult; MultiplyMatrixCore(numOps, aData, aOffset, bData, 0, resultData, resultOffset, m, n, p, allowParallel: false); });int matrixSizeA = m * n; int matrixSizeB = n * p; int matrixSizeResult = m * p; + if (batchSize == 1) + { + MultiplyMatrixCore(numOps, aData, 0, bData, 0, resultData, 0, m, n, p, allowParallel: true); + return result; + } + Parallel.For(0, batchSize, batch => { int aOffset = batch * matrixSizeA; int bOffset = batch * matrixSizeB; int resultOffset = batch * matrixSizeResult; MultiplyMatrixCore(numOps, aData, aOffset, bData, bOffset, resultData, resultOffset, m, n, p, allowParallel: false); });Also applies to: 4436-4488
src/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cs (1)
31-49: Minor: Redundant null assignment.Lines 47-48 are redundant since
bufferis already set tonullat line 33 and isn't modified on the false path.♻️ Suggested simplification
if (_buckets.TryGetValue(size, out var bucket) && bucket.Buffers.TryTake(out var candidate)) { Interlocked.Decrement(ref bucket.Count); candidate.MarkRented(); buffer = candidate; return true; } - buffer = null; return false; }src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs (1)
9212-9263:DirectOpenClGpuBufferpooling state: avoid magic numbers + make Dispose resilient if pool is disposed.
Today,Dispose()calls_returnToPool(this)directly; if the backend/pool is already disposed (or decides to throw),Dispose()can throw and/or leak the underlying CL buffer. Consider a defensive fallback toRelease()and make the states self-documenting.Proposed hardening (non-breaking)
internal sealed class DirectOpenClGpuBuffer : IGpuBuffer, IPoolableGpuBuffer { internal readonly DirectOpenClBuffer Buffer; private readonly Action<DirectOpenClGpuBuffer>? _returnToPool; private int _poolState; + private const int PoolStateRented = 0; + private const int PoolStateReturned = 1; + private const int PoolStateReleased = 2; @@ public void MarkRented() { - Interlocked.Exchange(ref _poolState, 0); + Interlocked.Exchange(ref _poolState, PoolStateRented); } public void Release() { - if (Interlocked.Exchange(ref _poolState, 2) == 2) + if (Interlocked.Exchange(ref _poolState, PoolStateReleased) == PoolStateReleased) return; Buffer.Dispose(); } public void Dispose() { if (_returnToPool == null) { Release(); return; } - if (Interlocked.CompareExchange(ref _poolState, 1, 0) != 0) + if (Interlocked.CompareExchange(ref _poolState, PoolStateReturned, PoolStateRented) != PoolStateRented) return; - _returnToPool(this); + try + { + _returnToPool(this); + } + catch + { + // Pool might be disposed or rejecting; don't leak and avoid throwing from Dispose(). + Release(); + } } }src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (2)
7-10: Potentially redundantusing AiDotNet.Tensors.Engines.DirectGpu;with repo-wide global usingsIf
AiDotNet.Tensors.*is already globally imported (per repo learnings), thisusingcan likely be dropped to avoid drift in style.using System.Threading;makes sense forInterlocked.
Based on learnings, ...
41-45: Make pool limits configurable (or at least document the memory rationale)Hard-coding
MaxPooledBufferElements/MaxPooledBuffersPerSizeis fine as a first pass, but consider making them tunable (env var / options) or adding a short comment explaining the intended memory cap and why those defaults are safe across GPUs.AiDotNetBenchmarkTests/TensorFlowComparisonBenchmarks.cs (1)
121-178: Cleanup is thorough but consider consolidating disposal loops.The cleanup properly disposes all GPU resources. For maintainability, you could consolidate the repetitive disposal patterns into a helper.
♻️ Optional: Consolidate disposal loops
+ private static void DisposeAll<T>(IEnumerable<T> items) where T : IDisposable + { + foreach (var item in items) item.Dispose(); + } + [GlobalCleanup] public void Cleanup() { - foreach (var tensor in _aiGpuMatricesA.Values) - { - tensor.Dispose(); - } - - foreach (var tensor in _aiGpuMatricesB.Values) - { - tensor.Dispose(); - } - // ... etc + DisposeAll(_aiGpuMatricesA.Values); + DisposeAll(_aiGpuMatricesB.Values); + DisposeAll(_aiGpuVectorsA.Values); + DisposeAll(_aiGpuVectorsB.Values); + DisposeAll(_aiGpuAddOutputs.Values); + DisposeAll(_aiGpuMultiplyOutputs.Values); + _aiGpuConvInput?.Dispose(); + DisposeAll(_tfMatricesA.Values); + // ... etc }AiDotNetBenchmarkTests/TorchSharpCpuComparisonBenchmarks.cs (1)
228-234: AiDotNet benchmarks return results while TorchSharp benchmarks are void.The AiDotNet benchmarks return
Tensor<float>which BenchmarkDotNet will consume to prevent dead-code elimination, while TorchSharp benchmarks use theConsumerpattern. This inconsistency is minor but worth noting for consistency.AiDotNetBenchmarkTests/MlNetCpuComparisonBenchmarks.cs (2)
128-136: VBuffer allocation on every iteration adds measurement overhead.Line 133 creates a new
VBuffer<float>on each benchmark iteration. WhileVBufferis a struct (no heap allocation for the struct itself), this could still introduce slight overhead compared to the AiDotNet path.For a fairer comparison, consider pre-allocating the VBuffer in setup similar to
_mlVBufferA.♻️ Suggested improvement
Add a pre-allocated VBuffer dictionary for multiply destinations:
private readonly Dictionary<int, VBuffer<float>> _mlVBufferA = new(); +private readonly Dictionary<int, VBuffer<float>> _mlVBufferB = new();In Setup, after creating
_mlVBufferA:_mlVBufferA[size] = new VBuffer<float>(size, dataA); +_mlVBufferB[size] = new VBuffer<float>(size, _mlMultiplyTargets[size]);Then in benchmark:
public void MlNet_Multiply(int size) { var dstValues = _mlMultiplyTargets[size]; Array.Copy(_mlVectorsB[size], dstValues, size); var src = _mlVBufferA[size]; - var dst = new VBuffer<float>(size, dstValues); + var dst = _mlVBufferB[size]; MulElementWise(ref src, ref dst); _consumer.Consume(dstValues); }
138-160: Sum and Mean benchmarks lack size parameterization.
AiDotNet_TensorSum,MlNet_Sum,AiDotNet_TensorMean, andMlNet_Meanonly test withVectorSizes[1](1M elements). This is inconsistent with the Add/Multiply benchmarks which test both 100k and 1M sizes.Consider adding
[Arguments]attributes to match the other benchmarks for consistent coverage.♻️ Suggested improvement
[Benchmark] -public float AiDotNet_TensorSum() +[Arguments(100_000)] +[Arguments(1_000_000)] +public float AiDotNet_TensorSum(int size) { - return AiDotNetEngine.Current.TensorSum(_aiVectorsA[VectorSizes[1]]); + return AiDotNetEngine.Current.TensorSum(_aiVectorsA[size]); } [Benchmark] -public float MlNet_Sum() +[Arguments(100_000)] +[Arguments(1_000_000)] +public float MlNet_Sum(int size) { - return SumVector(_mlVectorsA[VectorSizes[1]]); + return SumVector(_mlVectorsA[size]); }Apply similar changes to
AiDotNet_TensorMeanandMlNet_Mean.AiDotNetBenchmarkTests/TensorFlowCpuComparisonBenchmarks.cs (1)
253-307: ReLU, Sigmoid, Sum, and Mean benchmarks lack size parameterization.Similar to the ML.NET benchmarks, these operations only test with
VectorSizes[1](1M elements), while Add/Multiply test both sizes. Consider adding[Arguments]for consistency across the benchmark suite.AiDotNetBenchmarkTests/Program.cs (1)
83-148: Consider usingManualConfig.Create(baseConfig)to inherit the base configuration and eliminate manual copying.The manual copying of all config elements is thorough and correct. However, BenchmarkDotNet's
ManualConfig.Create(baseConfig)constructor accepts a base config and automatically inherits all its elements (exporters, loggers, diagnosers, etc.), which would simplify this code by eliminating the foreach loops entirely:var config = ManualConfig.Create(baseConfig) .WithUnionRule(ConfigUnionRule.AlwaysUseLocal) .WithOptions(ConfigOptions.DisableOptimizationsValidator) .AddJob(Job.ShortRun.WithId("ShortRun"));That said, the explicit approach ensures full control and works reliably across BenchmarkDotNet versions.
src/AiDotNet.Tensors/Helpers/BlasProvider.cs (1)
109-127: Make the double-checked init safe (volatile orLazyInitializer).With the current pattern, another thread can observe
_initialized == truewithout a matching acquire fence.Low-impact fix
- private static bool _initialized; + private static volatile bool _initialized;src/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cs (1)
691-697: Optional: consider poolingresultDatafor large fallback multiplies.If the recursive fallback is hit for large
M*N, renting fromArrayPool<T>(with atry/finallyreturn) can reduce LOH/GC churn.src/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cs (1)
106-170: Blocked kernel looks correct and parallel-safe for the row-block partitioning.Each task owns a disjoint row range, so
cwrites shouldn’t overlap. Consider adding cheap argument validation (orDebug.Assert) for offsets/strides vs span lengths since this method is callable from multiple sites.
📜 Review details
Configuration used: Organization UI
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (20)
AiDotNetBenchmarkTests/GpuResidentQuickHarness.csAiDotNetBenchmarkTests/MlNetCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/Program.csAiDotNetBenchmarkTests/TensorFlowComparisonBenchmarks.csAiDotNetBenchmarkTests/TensorFlowCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/TorchSharpComparisonBenchmarks.csAiDotNetBenchmarkTests/TorchSharpCpuComparisonBenchmarks.csBENCHMARKS.mdsrc/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Engines/CpuEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/Kernels/LossKernels.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cssrc/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cssrc/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cssrc/Polyfills/CompilerAttributePolyfills.cs
🧰 Additional context used
🧠 Learnings (3)
📚 Learning: 2025-12-18T08:49:25.295Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 444
File: src/Interfaces/IPruningMask.cs:1-102
Timestamp: 2025-12-18T08:49:25.295Z
Learning: In the AiDotNet repository, the project-level global using includes AiDotNet.Tensors.LinearAlgebra via AiDotNet.csproj. Therefore, Vector<T>, Matrix<T>, and Tensor<T> are available without per-file using directives. Do not flag missing using directives for these types in any C# files within this project. Apply this guideline broadly to all C# files (not just a single file) to avoid false positives. If a file uses a type from a different namespace not covered by the global using, flag as usual.
Applied to files:
src/Polyfills/CompilerAttributePolyfills.cssrc/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cssrc/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cssrc/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cssrc/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/Kernels/LossKernels.csAiDotNetBenchmarkTests/MlNetCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/TorchSharpCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/GpuResidentQuickHarness.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.csAiDotNetBenchmarkTests/TensorFlowCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/TensorFlowComparisonBenchmarks.csAiDotNetBenchmarkTests/Program.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Engines/CpuEngine.csAiDotNetBenchmarkTests/TorchSharpComparisonBenchmarks.cs
📚 Learning: 2025-12-18T08:49:53.103Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 444
File: src/Interfaces/IPruningStrategy.cs:1-4
Timestamp: 2025-12-18T08:49:53.103Z
Learning: In this repository, global using directives are declared in AiDotNet.csproj for core namespaces (AiDotNet.Tensors.* and AiDotNet.*) and common system types. When reviewing C# files, assume these global usings are in effect; avoid adding duplicate using statements for these namespaces and for types like Vector<T>, Matrix<T>, Tensor<T>, etc. If a type is not found, verify the global usings or consider adding a file-scoped using if needed. Prefer relying on global usings to reduce boilerplate.
Applied to files:
src/Polyfills/CompilerAttributePolyfills.cssrc/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cssrc/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cssrc/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cssrc/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/Kernels/LossKernels.csAiDotNetBenchmarkTests/MlNetCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/TorchSharpCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/GpuResidentQuickHarness.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.csAiDotNetBenchmarkTests/TensorFlowCpuComparisonBenchmarks.csAiDotNetBenchmarkTests/TensorFlowComparisonBenchmarks.csAiDotNetBenchmarkTests/Program.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Engines/CpuEngine.csAiDotNetBenchmarkTests/TorchSharpComparisonBenchmarks.cs
📚 Learning: 2025-11-19T04:08:26.895Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 0
File: :0-0
Timestamp: 2025-11-19T04:08:26.895Z
Learning: For ILGPU GPU operations in GpuEngine.cs, use standard .NET exception types (InvalidOperationException, ArgumentException, OutOfMemoryException) instead of ILGPU-specific exception types, as ILGPU exception types may be version-specific. Combine with message-based filtering using ex.Message.Contains("device") or ex.Message.Contains("accelerator") as a fallback for GPU-specific errors.
Applied to files:
src/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.csAiDotNetBenchmarkTests/TorchSharpComparisonBenchmarks.cs
⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (4)
- GitHub Check: CodeQL analysis (csharp)
- GitHub Check: Agent
- GitHub Check: Build (Windows)
- GitHub Check: CodeQL Analysis
🔇 Additional comments (52)
src/Polyfills/CompilerAttributePolyfills.cs (1)
73-86: LGTM! Correct polyfill implementation forNotNullWhenAttribute.The implementation correctly matches the BCL signature with proper
AttributeUsage, constructor, and property. Thepublicvisibility is appropriate here since this attribute is applied to parameters in public API signatures (e.g.,bool TryGetValue([NotNullWhen(true)] out T? value)), unlike the compiler-internal attributes which can remaininternal.src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/Kernels/LossKernels.cs (2)
991-991: LGTM!The kernel name array correctly reflects the renamed kernels in the source string.
372-372: Verify that all callers have been updated to reference the new kernel names.The
loss_prefix aligns with other loss-related kernels (e.g.,loss_mse,loss_bce) and helps avoid naming collisions. However, confirm that no references to the old kernel names (reduce_sumandreduce_mean) remain in the codebase before approving. Search the C# source files for any string references to the old kernel names that may need updating toloss_reduce_sumandloss_reduce_mean.src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs (2)
1-2: LGTM!The using directive is required for the new GPU-related types (
GpuExecutionContext,GpuExecutionOptions,DirectGpuTensorEngine).
166-180: LGTM!The pattern matching approach is thread-safe as it captures the engine reference locally. The delegation to the underlying GPU engine keeps this method simple.
src/AiDotNet.Tensors/Engines/CpuEngine.cs (5)
1403-1417: MatrixMultiply delegation toMatrix<T>.Multiplyis clean.
Nice simplification, and it should automatically pick up the optimized matmul paths behindMatrix<T>.Multiply.
3426-3432: Conv2D early-exit fast-path is well placed.
Falling back to the existing direct convolution keeps behavior intact when the heuristic says “no”.
4367-4378: TensorMatMul2D correctly funnels through the shared core.
This should ensure the 2D path benefits from GEMM/blocked implementations consistently.
4508-4511: Conv2D(int[] ...) symmetric fast-path looks good.
Keeps the “fully general” path while letting the common symmetric case reuse the optimized scalar-params implementation.
4321-4363: Verify GEMM/blocked paths honor offsets (batched correctness depends on it).
MultiplyMatrixCoreis now called with non-zeroaOffset/bOffset/cOffsetfrom batched matmul and the Conv2D im2col path. Please double-check thatMatrixMultiplyHelper.TryGemm(...)andMatrixMultiplyHelper.MultiplyBlocked(...)interpret those offsets/leading dimensions correctly (otherwise results will silently corrupt for batched inputs).src/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cs (4)
6-10: LGTM!Clean interface design for poolable GPU buffers with the essential lifecycle methods.
14-18: LGTM!Maintaining a separate
Countfield withInterlockedoperations is a reasonable optimization to avoid the O(n) cost ofConcurrentBag.Count. The slight count drift during concurrent operations is acceptable for pool capacity management.
20-29: LGTM!Constructor defensively handles invalid parameters with sensible defaults. The use of
intfor the_disposedflag is appropriate for atomic operations.
51-69: LGTM!The return logic correctly handles pool capacity limits. The increment-before-add pattern may cause brief count drift, but
TryRenthandles empty bags gracefully by returning false.src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (3)
8-10: Newusingdirectives are justified (pool +Interlocked).
No issues here.
9025-9033: Disposing_bufferPoolinHipBackend.Dispose()is the right direction.
Order (pool dispose beforehipStreamDestroy) looks sensible.
45-48: [rewritten review comment]
[exactly one classification tag]src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs (2)
9-10: Imports look fine for the new pooling/state code.
System.Threadingis needed forInterlocked, andAiDotNet.Tensors.Engines.DirectGpuis needed for the pool types.
40-43: Buffer pooling: verify queue/stream safety for multi-queue scenarios.Because
SupportsMultiStream => true, buffers returned to_bufferPoolfrom work enqueued on queue A can be re-rented and used on queue B while queue A is still in-flight, causing data corruption. This risk only disappears if:
- The pool is queue/event-aware and enforces reuse constraints per queue, or
- Completion events gate all rentals until prior work completes, or
- Pooling is intentionally disabled for multi-stream execution.
Confirm the
GpuBufferPoolimplementation guards against cross-queue aliasing, properly releases oversize buffers, and behaves safely after disposal. If not queue-aware, consider pool-per-queue isolation or disabling multi-stream pooling.Also applies to: 721-741, 9184-9205
src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (1)
370-389: Good: zeroing pooled buffers on rent; verify pool buckets by exact size
cuMemsetD32(pooled.Handle, 0, (ulong)size)is correct if the pool only returns buffers of the exact requested element count (per-size buckets). If the pool allows “>= size” reuse, this will leave tail data dirty andIGpuBuffer.Sizemay not match expectations.Not necessarily a change needed here—just ensure the pool keys by exact
Size(or adjust clearing / metadata accordingly).AiDotNetBenchmarkTests/TensorFlowComparisonBenchmarks.cs (2)
49-119: LGTM! Proper GPU initialization and validation.The setup correctly validates GPU availability for both AiDotNet and TensorFlow backends before proceeding. The null checks and exception throwing ensure benchmarks fail fast if GPU is unavailable.
287-295: LGTM! GPU-resident benchmark with proper synchronization.The benchmark correctly uses GPU-resident tensors, disposes the result, and synchronizes to ensure accurate timing measurements.
AiDotNetBenchmarkTests/TorchSharpCpuComparisonBenchmarks.cs (2)
43-88: LGTM! Clean CPU-only setup with proper TorchSharp configuration.The setup correctly initializes the CPU engine, disables gradient computation for TorchSharp, and validates tensor initialization. Good defensive programming with null checks.
90-115: Cleanup disposes TorchSharp tensors but not AiDotNet tensors.The cleanup correctly disposes TorchSharp tensors. AiDotNet CPU tensors (
Tensor<float>) don't appear to require explicit disposal based on the learnings about this repository, so this is acceptable.AiDotNetBenchmarkTests/TorchSharpComparisonBenchmarks.cs (3)
56-129: LGTM! Comprehensive GPU initialization with proper validation.The setup validates both AiDotNet and TorchSharp GPU availability before proceeding. The sequential checks and early-exit exceptions ensure the benchmark suite fails fast if either GPU backend is unavailable.
296-304: Good synchronization pattern for TorchSharp CUDA results.The
ConsumeTorchResultmethod properly synchronizes CUDA before consuming, ensuring accurate timing measurements and preventing measurement artifacts from asynchronous GPU execution.
448-455: Mean calculation creates two GPU tensors but only synchronizes the second.The mean calculation creates both
sumandmeanGPU tensors. While themean.Synchronize()call should implicitly wait for thesumcomputation, it would be more explicit to ensure both are synchronized. This is likely fine in practice since they're sequential operations.BENCHMARKS.md (3)
18-26: LGTM! Clear documentation of GPU vs CPU benchmark paths.The updated run commands with comments (
# GPU vs GPU,# CPU vs CPU) make it easy to understand which benchmark suite to run for each scenario.
28-29: Good historical context for existing benchmark results.The note explaining that historical results predate the CPU/GPU split helps users understand the benchmark comparison context.
89-99: LGTM! New ML.NET CPU comparison section is well-formatted.The new section follows the established documentation pattern with location and run instructions.
AiDotNetBenchmarkTests/GpuResidentQuickHarness.cs (5)
16-73: LGTM! Robust GPU availability checking with graceful fallback.The multiple validation steps (AutoDetectAndConfigureGpu, engine type check, backend availability, context creation) provide comprehensive error handling. The early returns with console messages make debugging easier.
78-86: Conv input upload failure should not prevent other benchmarks.Good design decision to allow other benchmarks to proceed when Conv2D upload fails, with the appropriate skip message.
166-186: LGTM! Accurate timing methodology with proper synchronization.The
RunTimedmethod correctly:
- Synchronizes before warmup
- Synchronizes after each warmup iteration
- Synchronizes before measurement
- Synchronizes after each measured iteration
- Computes average correctly
This ensures GPU operations complete before timing boundaries.
208-224: CompositeDisposable disposes in reverse order (LIFO), which is correct.The disposal order (second then first) follows the standard LIFO pattern for dependent resources, which is appropriate for the mean calculation where
meandepends onsum.
87-142: Resource cleanup in finally block may be incomplete but cannot be verified without access to codebase.The finally block does not dispose
convKernel(created at line 79). However, verification requires checking theTensor<float>class definition to confirm whetherIDisposableis implemented and whetherconvKernelrequires explicit cleanup. This cannot be determined from available sources.AiDotNetBenchmarkTests/Program.cs (3)
11-43: LGTM!The argument preprocessing logic is well-structured. The GPU harness integration with conditional compilation is appropriate, and the short-run configuration path is cleanly handled. The early return after GPU harness execution prevents unintended benchmark runs.
45-55: LGTM!Clean and correct implementation of flag removal.
57-81: LGTM!The option parsing handles edge cases well: missing option, option without value, and invalid/non-positive values all fall back to defaults. The removal order (value before option) is correct to maintain index validity.
AiDotNetBenchmarkTests/MlNetCpuComparisonBenchmarks.cs (4)
34-55: LGTM!The setup correctly initializes both AiDotNet and ML.NET data structures with deterministic seed values for reproducibility. Pre-allocating target arrays and VBuffers is appropriate for benchmark accuracy.
57-76: LGTM!Comprehensive warmup covering all benchmarked operations to prime JIT compilation and CPU caches.
78-96: LGTM!Deterministic data generation using well-known LCG constants ensures reproducible benchmark results across runs.
162-209: LGTM!The reflection-based delegate creation is well-implemented with proper error handling. Using
MakeByRefType()for theMulElementWisedelegate correctly handles therefparameters. Static initialization ensures the reflection cost is paid once.AiDotNetBenchmarkTests/TensorFlowCpuComparisonBenchmarks.cs (6)
13-36: LGTM!Clean separation of AiDotNet and TensorFlow tensor dictionaries. The nullable fields for Conv2D tensors are appropriate given the deferred initialization pattern.
37-80: LGTM!Comprehensive setup with proper CPU-forcing via
CUDA_VISIBLE_DEVICES=-1. The validation checks after Conv2D initialization are good defensive coding.
82-107: LGTM!Proper cleanup of TensorFlow tensors which wrap unmanaged resources. The null-conditional disposal for conv tensors is appropriate.
109-148: LGTM!Thorough warmup with
.numpy()calls to force TensorFlow eager execution. All benchmarked operations are covered.
150-177: LGTM!Conv2D initialization correctly handles the different data layouts between AiDotNet (NCHW) and TensorFlow (NHWC). While the actual data values differ due to the layout transformation, this is acceptable for benchmarking compute performance.
207-329: LGTM!The benchmark methods correctly return results to prevent dead code elimination. The
.numpy()calls on TensorFlow operations ensure eager execution for accurate timing comparison.src/AiDotNet.Tensors/Helpers/BlasProvider.cs (1)
49-107: Good: bounds/stride validation before unsafe BLAS call.
HasEnoughData+ thefixed/offset math keep the native call reasonably well-guarded, andalpha=1/beta=0matches the “C = A*B” usage.Also applies to: 184-198
src/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cs (1)
675-689: Nice fast-path ordering: GEMM → blocked → recursive fallback.This keeps the public API unchanged while avoiding the recursive-path allocation when GEMM/blocked succeed.
src/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cs (2)
209-265: Good: block-size heuristic + env overrides with clamping.The override +
[16,128]clamp keeps tuning sane while allowing experimentation.
276-296: Good: env parsing rejects non-positive values.Keeps “unset/invalid” distinct from “0” and avoids surprising thresholds.
✏️ Tip: You can disable this entire section by setting review_details to false in your review settings.
There was a problem hiding this comment.
Actionable comments posted: 8
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (3)
src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs (1)
4131-4165: CLBlast library GEMM: add minimal error checking aroundFinishand confirmSgemmsignature/ABI.
GemmWithClBlastcallsOpenClNativeBindings.Finish(...)but ignores the return code. Also, since this is a P/Invoke boundary, please confirm theClBlastNative.Sgemmsignature matches the platform ABI (esp.ref queueand the lasteventarg).Proposed fix (check Finish return)
--- a/src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs +++ b/src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs @@ var sw = Stopwatch.StartNew(); bool ok = TryExecuteClBlastLibraryGemm(A, B, C, M, N, K, alpha, beta); - OpenClNativeBindings.Finish(_context.CommandQueue); + int finishErr = OpenClNativeBindings.Finish(_context.CommandQueue); + if (finishErr != OpenClNativeBindings.CL_SUCCESS) + throw new InvalidOperationException($"clFinish failed: {finishErr}"); sw.Stop();Also applies to: 4167-4181
src/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs (1)
1526-1542: DisposeoutputBufferon exception inActivationGputo avoid GPU memory leaks.
IfApplyGpuActivationthrows (e.g., unsupported activation),outputBufferis currently leaked.Proposed fix
public IGpuTensor<T> ActivationGpu<T>(IGpuTensor<T> input, FusedActivationType activation) { if (!TryGetBackend(out var backend)) throw new InvalidOperationException("No GPU backend available for ActivationGpu"); // Allocate output buffer int size = input.ElementCount; var outputBuffer = backend.AllocateBuffer(size); - if (activation == FusedActivationType.None) - { - // Preserve previous behavior: output is a copy of input. - backend.Copy(input.Buffer, outputBuffer, size); - } - else - { - ApplyGpuActivation(backend, input.Buffer, outputBuffer, size, activation); - } - - // Return new GPU tensor - return new GpuTensor<T>(backend, outputBuffer, input.Shape, GpuTensorRole.Activation, ownsBuffer: true); + try + { + if (activation == FusedActivationType.None) + { + // Preserve previous behavior: output is a copy of input. + backend.Copy(input.Buffer, outputBuffer, size); + } + else + { + ApplyGpuActivation(backend, input.Buffer, outputBuffer, size, activation); + } + + // Return new GPU tensor + return new GpuTensor<T>(backend, outputBuffer, input.Shape, GpuTensorRole.Activation, ownsBuffer: true); + } + catch + { + outputBuffer.Dispose(); + throw; + } }src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
9819-9866:HipGpuBuffer.SizeInBytescan overflow (int math) + consider hardening pool state transitions.
Size * sizeof(float)multiplies asintfirst; cast tolongshould happen before multiply. Also,MarkRented()currently allows resurrecting a buffer afterRelease()(state2) if misused.Proposed fix (SizeInBytes)
internal sealed class HipGpuBuffer : IGpuBuffer, IPoolableGpuBuffer { public IntPtr Handle { get; } public int Size { get; } - public long SizeInBytes => Size * sizeof(float); + public long SizeInBytes => (long)Size * sizeof(float);
🤖 Fix all issues with AI agents
In `@AiDotNetBenchmarkTests/GpuResidentQuickHarness.cs`:
- Around line 219-235: CompositeDisposable currently calls _second.Dispose() and
_first.Dispose() unconditionally, which can throw if the constructor was passed
null; make the class null-safe by allowing the backing fields to be nullable
(IDisposable?) and guarding disposal calls in Dispose() (e.g., if (_second !=
null) _second.Dispose(); if (_first != null) _first.Dispose();), and update the
constructor signature to accept IDisposable? for first and second so null inputs
are handled without throwing.
In `@src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs`:
- Around line 172-180: The Action<GpuExecutionContext> overload of
WithGpuContext lacks a null check for the action parameter; add validation at
the start of AiDotNet.Tensors.Engines.AiDotNetEngine.WithGpuContext to throw
ArgumentNullException(nameof(action)) if action is null, mirroring the generic
overload's behavior, then proceed to cast Current to DirectGpuTensorEngine and
call gpuEngine.WithGpuContext(action, options) as before.
In `@src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs`:
- Line 368: The pooled dispose currently calls _returnToPool(this) directly
(e.g., CudaGpuBuffer.Dispose()), which can throw or be invalid if the
pool/backend is already disposed; change Dispose to call _returnToPool(this)
inside a try and on any exception or if _returnToPool is null fall back to
calling Release() to free the device memory, swallowing/logging secondary errors
but ensuring device memory is always freed; apply the same pattern wherever
buffers are created with _bufferPool.Return (the ctor call sites that pass
_bufferPool.Return and corresponding Dispose implementations around the
codebase).
In `@src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs`:
- Around line 310-325: The multiplication in ShouldUseVendorGemm uses (long)m *
n * k which can overflow; change the work calculation to use unsigned 64-bit
arithmetic (e.g., compute work as (ulong)m * (ulong)n * (ulong)k) and compare
against an unsigned threshold, updating the helper call to a GetEnvULong (or
cast GetEnvLong/DefaultVendorGemmThreshold to ulong) so the comparison is done
safely without signed overflow; adjust references to GemmVendorThresholdEnvVar
and DefaultVendorGemmThreshold accordingly.
In `@src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cs`:
- Around line 11-13: Mark the field _checkedAvailability as volatile to fix the
double-checked locking memory-ordering issue so threads cannot see
_checkedAvailability == true while reading a stale _isAvailable; update the
declaration "private static bool _checkedAvailability;" to "private static
volatile bool _checkedAvailability;" (or alternatively replace the pattern by
using Volatile.Read/Volatile.Write around reads/writes to _checkedAvailability
and _isAvailable, or switch to Lazy<bool>), keeping AvailabilityLock and the
existing double-check logic intact.
In `@src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs`:
- Around line 52-56: Overflow risk: the expression long work = (long)m * n * k
can overflow for large m,n,k and must be saturated. Replace each raw
cast-multiplication that computes work (the occurrences that reference m, n, k
used for GEMM selection, e.g., the line with "long work = (long)m * n * k" and
the similar block later) with a checked multiply sequence that catches
OverflowException and sets work = long.MaxValue (or another saturated cap), e.g.
perform the multiplies inside a try { work = checked((long)m * n * k); } catch
(OverflowException) { work = long.MaxValue; } so GEMM routing never receives a
wrapped/negative value; apply the same change to all occurrences (including the
second region noted around the later GEMM-selection code).
♻️ Duplicate comments (3)
src/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cs (1)
50-74: Defensive disposal check is good, but a subtle race remains.The re-check of
_disposedafter incrementingbucket.Count(lines 60-65) is a good defensive pattern. However, as noted in previous reviews, there's still a small window where a buffer could be added to the bag (line 73) afterDisposehas already iterated past that bucket.This is generally acceptable if
Disposeis called only during application shutdown when no concurrent operations are expected. If stricter guarantees are needed, consider documenting this assumption or implementing a multi-pass drain inDispose.src/AiDotNet.Tensors/Engines/CpuEngine.cs (1)
3474-3629: Harden Im2Col size guards againstlongoverflow (OOM bypass risk).
work,kernelElems, andperBatchElemsare computed with unchecked multiplications; overflow can wrap negative and slip past themaxScratchBytesgate.Proposed fix (overflow-safe arithmetic)
// Avoid Im2Col packing overhead for small workloads. const long im2ColWorkThreshold = 1_000_000; - long work = colRowsLong * colColsLong * outChannels; - if (work < im2ColWorkThreshold) + long work; + try + { + work = checked(colRowsLong * colColsLong * outChannels); + } + catch (OverflowException) + { + return false; + } + if (work < im2ColWorkThreshold) { return false; } @@ - long kernelElems = (long)outChannels * colColsLong; - long perBatchElems = colSizeLong + ((long)outChannels * colRowsLong); + long kernelElems; + long perBatchElems; + try + { + kernelElems = checked((long)outChannels * colColsLong); + perBatchElems = checked(colSizeLong + checked((long)outChannels * colRowsLong)); + } + catch (OverflowException) + { + return false; + } long approxScratchBytes; try { approxScratchBytes = checked((kernelElems + (perBatchElems * parallelBatches)) * elementSize); } catch (OverflowException) { return false; }Im2Col threshold constant is still magic.
im2ColWorkThreshold = 1_000_000should be explained or surfaced as a named/configurable knob (matches prior review feedback).Optional perf improvement: consider pooling
colMatrix/outputMatrix(e.g.,ArrayPool<T>) to avoid per-batch GC churn.src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
555-560: Unmanaged HIP calls: comments don’t address code-scanning finding by themselves.
If GitHub Advanced Security continues flagging these P/Invoke sites, you may need a suppression/justification mechanism (or a centralized interop wrapper) rather than inline comments.Also applies to: 577-582, 625-632
🧹 Nitpick comments (10)
AiDotNetBenchmarkTests/Program.cs (1)
83-148: LGTM! Comprehensive config builder.The method thoroughly transfers all configuration components from the base config while applying the
Job.ShortRunprofile. TheAlwaysUseLocalunion rule ensures the short-run job takes precedence.Minor observation: Line 145 passes
baseConfig.SummaryStyletoWithSummaryStylewithout a null check (unlikeOrdereron line 140). Consider adding a null guard ifSummaryStylecan be null, though BenchmarkDotNet typically provides a default.🔧 Optional: Add null check for SummaryStyle
- config.WithSummaryStyle(baseConfig.SummaryStyle); + if (baseConfig.SummaryStyle != null) + { + config.WithSummaryStyle(baseConfig.SummaryStyle); + }src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs (1)
9290-9341: Pool state machine is mostly sound; consider guarding against “rent after Release” and clarifyingRelease()semantics.
TodayRelease()permanently frees the underlying CL buffer even for pool-backed buffers (if someone calls it directly). If that’s intended, ok; otherwise consider making itinternal(ifIPoolableGpuBufferallows) or documenting “Dispose returns to pool; Release permanently frees”. Also consider detectingMarkRented()called after_poolState==2(released) to fail fast.src/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cs (1)
701-728: Redundant allocation in the recursive fallback path.When falling back to the recursive path,
resultDatais allocated as a new array (line 722), then immediately copied toresult.AsWritableSpan()(line 726). Theresultmatrix was already allocated at line 696. Consider writing directly to the result's backing memory to avoid this extra allocation and copy.♻️ Suggested optimization
// Use cache-oblivious recursive algorithm MatrixMultiplyHelper.TraceMatmul("RECURSIVE", M, N, K); - var resultData = new T[M * N]; - MultiplyRecursive(_memory.ToArray(), other._memory.ToArray(), resultData, 0, 0, 0, 0, 0, 0, M, K, N, K, N, N); - - // Copy result back to the result matrix - resultData.AsSpan().CopyTo(result.AsWritableSpan()); + var resultData = result._memory.ToArray(); + MultiplyRecursive(_memory.ToArray(), other._memory.ToArray(), resultData, 0, 0, 0, 0, 0, 0, M, K, N, K, N, N); + resultData.AsSpan().CopyTo(result.AsWritableSpan()); return result;Note: This still requires a copy back because
MultiplyRecursiveworks on arrays whileresult._memorymay not be array-backed. Consider refactoringMultiplyRecursiveto work withMemory<T>directly for true zero-copy in a future iteration.src/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cs (2)
423-455: Consider lock contention under high-throughput scenarios.The
GetOrCreatemethod uses a global lock (CacheLock) for all cache operations. In highly concurrent matmul workloads with different matrices, this could become a bottleneck. Consider using aConcurrentDictionarywith per-matrix locking orReaderWriterLockSlimfor the read-heavy case where cached entries are hit frequently.For now, this is acceptable if most workloads don't involve many different matrices being multiplied concurrently. The version check ensures correctness.
362-370: Replace the customClampmethod withMath.Clamp.The project targets .NET 8.0, which has
Math.Clampavailable as a built-in method since .NET 6.0. Replace the custom implementation withMath.Clamp(value, min, max).Current implementation (lines 362-370)
private static int Clamp(int value, int min, int max) { if (value < min) { return min; } return value > max ? max : value; }src/AiDotNet.Tensors/Helpers/BlasProvider.cs (2)
121-139: Consider marking_initializedasvolatilefor double-checked locking correctness.The double-checked locking pattern reads
_initializedoutside the lock (line 123). Withoutvolatile, this read could observe a stale value on weakly-ordered architectures (e.g., ARM), or reordering could cause_availableto be read before it's fully written.Option 1: Add volatile
- private static bool _initialized; + private static volatile bool _initialized;Option 2: Use Lazy<T> for simpler thread-safe initialization
private static readonly Lazy<bool> _lazyAvailable = new Lazy<bool>(TryLoadLibrary); private static bool EnsureInitialized() => _lazyAvailable.Value;
164-174: Avoid side effects inside LINQ predicates.The
Whereclause modifies_libraryHandleas a side effect. While this works due to lazy evaluation stopping at the first match, side effects in LINQ predicates make the code harder to reason about and can lead to subtle bugs if the query is modified.Proposed refactor using explicit loop
- foreach (var candidate in GetCandidateLibraryNames() - .Where(name => TryLoadNativeLibrary(name, out _libraryHandle))) + foreach (var candidate in GetCandidateLibraryNames()) { + if (!TryLoadNativeLibrary(candidate, out _libraryHandle)) + { + continue; + } + if (TryLoadSymbols()) { return true; } FreeNativeLibrary(_libraryHandle); _libraryHandle = IntPtr.Zero; }src/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs (1)
7039-7060: DisposeoutputBufferon exception inSumAxisGpu(new transpose path).
The transpose-based reduction looks right, butoutputBufferwill leak ifTranspose/SumAxisthrows before the owningGpuTensor<T>is created.Proposed fix
var outputBuffer = backend.AllocateBuffer(outputSize); - - if (axis == 0) - { - if (innerSize == 1) - { - backend.SumAxis(input.Buffer, outputBuffer, 1, outerSize); - } - else - { - using var transposedBuffer = backend.AllocateBuffer(outerSize * innerSize); - backend.Transpose(input.Buffer, transposedBuffer, outerSize, innerSize); - backend.SumAxis(transposedBuffer, outputBuffer, innerSize, outerSize); - } - } - else - { - backend.SumAxis(input.Buffer, outputBuffer, outerSize, innerSize); - } - - return new GpuTensor<T>(backend, outputBuffer, outputShape, GpuTensorRole.Activation, ownsBuffer: true); + try + { + if (axis == 0) + { + if (innerSize == 1) + { + backend.SumAxis(input.Buffer, outputBuffer, 1, outerSize); + } + else + { + using var transposedBuffer = backend.AllocateBuffer(outerSize * innerSize); + backend.Transpose(input.Buffer, transposedBuffer, outerSize, innerSize); + backend.SumAxis(transposedBuffer, outputBuffer, innerSize, outerSize); + } + } + else + { + backend.SumAxis(input.Buffer, outputBuffer, outerSize, innerSize); + } + + return new GpuTensor<T>(backend, outputBuffer, outputShape, GpuTensorRole.Activation, ownsBuffer: true); + } + catch + { + outputBuffer.Dispose(); + throw; + }src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (2)
41-45: Consider making pool sizing configurable (or at least centrally documented).
These constants effectively define memory retention behavior and may need tuning across devices/workloads.
7820-7845:Release()should check CUDA return codes (try/catch won’t catch non-throwing error returns).
Right now failures fromcuCtxPushCurrent/cuMemFree/cuCtxPopCurrentcan be silently ignored because these APIs typically return error codes rather than throwing.Proposed fix (still finalizer-safe)
try { if (_context != IntPtr.Zero) { - CuBlasNative.cuCtxPushCurrent(_context); - CuBlasNative.cuMemFree(_devicePtr); - CuBlasNative.cuCtxPopCurrent(out _); + CuBlasNative.CheckCudaResult(CuBlasNative.cuCtxPushCurrent(_context), "cuCtxPushCurrent"); + CuBlasNative.CheckCudaResult(CuBlasNative.cuMemFree(_devicePtr), "cuMemFree"); + CuBlasNative.CheckCudaResult(CuBlasNative.cuCtxPopCurrent(out _), "cuCtxPopCurrent"); } } catch { // Suppress disposal errors to avoid crashing finalizers. }Also applies to: 7861-7864
📜 Review details
Configuration used: Organization UI
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (15)
AiDotNetBenchmarkTests/GpuResidentQuickHarness.csAiDotNetBenchmarkTests/Program.cssrc/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Engines/CpuEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cssrc/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cssrc/AiDotNet.Tensors/LinearAlgebra/Matrix.cssrc/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cstests/AiDotNet.Tensors.Benchmarks/Program.cs
🧰 Additional context used
🧠 Learnings (4)
📚 Learning: 2025-12-18T08:49:25.295Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 444
File: src/Interfaces/IPruningMask.cs:1-102
Timestamp: 2025-12-18T08:49:25.295Z
Learning: In the AiDotNet repository, the project-level global using includes AiDotNet.Tensors.LinearAlgebra via AiDotNet.csproj. Therefore, Vector<T>, Matrix<T>, and Tensor<T> are available without per-file using directives. Do not flag missing using directives for these types in any C# files within this project. Apply this guideline broadly to all C# files (not just a single file) to avoid false positives. If a file uses a type from a different namespace not covered by the global using, flag as usual.
Applied to files:
tests/AiDotNet.Tensors.Benchmarks/Program.cssrc/AiDotNet.Tensors/Engines/AiDotNetEngine.csAiDotNetBenchmarkTests/GpuResidentQuickHarness.csAiDotNetBenchmarkTests/Program.cssrc/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cssrc/AiDotNet.Tensors/LinearAlgebra/Matrix.cssrc/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cssrc/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/CpuEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs
📚 Learning: 2025-12-18T08:49:53.103Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 444
File: src/Interfaces/IPruningStrategy.cs:1-4
Timestamp: 2025-12-18T08:49:53.103Z
Learning: In this repository, global using directives are declared in AiDotNet.csproj for core namespaces (AiDotNet.Tensors.* and AiDotNet.*) and common system types. When reviewing C# files, assume these global usings are in effect; avoid adding duplicate using statements for these namespaces and for types like Vector<T>, Matrix<T>, Tensor<T>, etc. If a type is not found, verify the global usings or consider adding a file-scoped using if needed. Prefer relying on global usings to reduce boilerplate.
Applied to files:
tests/AiDotNet.Tensors.Benchmarks/Program.cssrc/AiDotNet.Tensors/Engines/AiDotNetEngine.csAiDotNetBenchmarkTests/GpuResidentQuickHarness.csAiDotNetBenchmarkTests/Program.cssrc/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cssrc/AiDotNet.Tensors/LinearAlgebra/Matrix.cssrc/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cssrc/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/CpuEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs
📚 Learning: 2025-11-19T04:08:26.895Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 0
File: :0-0
Timestamp: 2025-11-19T04:08:26.895Z
Learning: For ILGPU GPU operations in GpuEngine.cs, use standard .NET exception types (InvalidOperationException, ArgumentException, OutOfMemoryException) instead of ILGPU-specific exception types, as ILGPU exception types may be version-specific. Combine with message-based filtering using ex.Message.Contains("device") or ex.Message.Contains("accelerator") as a fallback for GPU-specific errors.
Applied to files:
src/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs
📚 Learning: 2025-12-19T19:05:13.598Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 445
File: src/Interfaces/IPredictionModelBuilder.cs:7-8
Timestamp: 2025-12-19T19:05:13.598Z
Learning: In ooples/AiDotNet, IAggregationStrategy<TModel> is declared in AiDotNet.Interfaces (src/Interfaces/IAggregationStrategy.cs). For src/Interfaces/IPredictionModelBuilder.cs (also in AiDotNet.Interfaces), referencing IAggregationStrategy does not require an additional using. Avoid flagging missing using directives for this case in future reviews.
Applied to files:
AiDotNetBenchmarkTests/Program.cs
⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)
- GitHub Check: Build (Windows)
- GitHub Check: CodeQL Analysis
🔇 Additional comments (49)
src/AiDotNet.Tensors/LinearAlgebra/Matrix.cs (1)
1055-1059: LGTM! Correct mutation tracking placement.The
MarkDirty()call is appropriately placed before the in-place division operation, ensuring any cached views are invalidated before data mutation. This aligns with the mutation-tracking pattern established inMatrixBase.cs.tests/AiDotNet.Tensors.Benchmarks/Program.cs (2)
72-72: LGTM! Usage documentation is complete.The help text correctly documents the new
--cpu-matmuloption alongside existing options.
32-36: VerifyCpuMatMulDiagnostics.Run()exists and is accessible.The
--cpu-matmuloption correctly follows the established CLI pattern for command-line handling. However, confirm that theCpuMatMulDiagnosticsclass and itsRun()method are properly implemented and accessible from this context.AiDotNetBenchmarkTests/Program.cs (5)
1-5: LGTM! Using directives are correctly added.Both
System.Collections.Generic(forList<string>) andBenchmarkDotNet.Jobs(forJob.ShortRun) are required by the new functionality.
22-32: LGTM! GPU harness integration is well-structured.The implementation correctly:
- Uses
NET8_0_OR_GREATERguard for GPU-dependent code- Provides sensible defaults for warmup (1) and iterations (3)
- Gracefully handles unsupported runtimes with an informative message
- Returns early after handling the flag
37-42: LGTM! Short-run configuration handling.The dual flag support (
--short-runand--gpu-short) provides flexibility, and the filtered argument list is correctly passed to the benchmark switcher.
45-55: LGTM! Clean flag removal helper.The implementation correctly handles flag detection and removal with appropriate return semantics.
57-81: LGTM! Robust option parsing with edge case handling.Good handling of:
- Missing option (returns default)
- Option at end without value (removes option, returns default)
- Invalid/non-positive values (returns default)
The removal order (value first, then option) correctly avoids index shift issues.
src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs (5)
9-10: New usings are fine; just ensure they’re actually required in this file.
System.Threadingis now required forInterlockedinDirectOpenClGpuBuffer, andAiDotNet.Tensors.Engines.DirectGpuappears needed forGpuBufferPool<>/IPoolableGpuBuffer.
545-573:ClampLocalSizeForKernelis a solid safety improvement; verify the native binding behavior on error paths.
IfGetKernelWorkGroupInfoSizeTcan throw (vs returningUIntPtr.Zero), this needs guarding to avoid reductions taking down the process.
761-786: Pooling integration inAllocateBufferlooks correct; verify pool rent semantics (MarkRented + correct length).
This code assumes a rentedDirectOpenClGpuBufferis in a “rented” state and sized to at least the requested element count. Please double-checkGpuBufferPool<DirectOpenClGpuBuffer>.TryRent(...)callsIPoolableGpuBuffer.MarkRented()and only returns size-compatible buffers.
1809-1827: Vendor CLBlast fast path: good fallback behavior; please verify env-var modes and thresholds behave as intended.
The routing is now: CLBlast library (if enabled by env/threshold) → CLBlast baseline kernels → dynamic tuned → built-in. That seems reasonable; just confirm expected overrides forAIDOTNET_GPU_GEMM_IMPL(“vendor/clblast” vs “builtin/internal/...”).
2646-2650: Nice: kernel-aware local-size clamping + explanatory comment.
This addresses kernel-specific work-group constraints forreduce_sum/reduce_max.Also applies to: 2676-2679
src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs (3)
1-2: LGTM!The new using directive for
AiDotNet.Tensors.Engines.Gpuis appropriate to support the new GPU context types.
156-164: LGTM!The
BeginGpuContextmethod correctly delegates to the GPU engine when available and returnsnullotherwise. The nullable return type clearly communicates the conditional availability.
190-203: LGTM!The generic overload properly validates both
funcandfallbackparameters withArgumentNullException. The fallback is correctly invoked when no GPU engine is available.src/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cs (2)
47-58: LGTM!The version tracking implementation uses
Interlockedoperations correctly for thread-safe reads and increments. Theinternalvisibility ofVersionappropriately limits access to within the assembly whileMarkDirty()isprotectedfor derived class usage.
240-245: LGTM!The
MarkDirty()call is correctly placed before the mutation in the indexer setter. This pattern is consistently applied across all mutation points.src/AiDotNet.Tensors/Helpers/MatrixMultiplyHelper.cs (4)
26-31: LGTM!The early type guard for
float/doubleis correctly placed at the start ofTryGemm<T>, avoiding wasted allocations for unsupported types. This addresses the previous review feedback.
107-143: LGTM!The packed multiplication path correctly leverages the column-major packed format for improved cache locality. The early exit when packing isn't beneficial and the use of SIMD-accelerated dot products are good optimizations.
188-252: LGTM!The blocked multiplication correctly handles the constraint that
Span<T>cannot be captured in closures by capturingMemory<T>and accessing.Spaninside the delegate. The i-k-j loop order is optimal for cache usage, and parallel execution is appropriately gated by the threshold and processor count.
285-310: LGTM!The block size calculation correctly uses the null-coalescing operator for
PlatformDetector.Capabilities?.L1CacheSizewith a sensible 32KB fallback. The formulasqrt(L1 / (2 * elementSize))ensures two blocks fit in L1 cache.src/AiDotNet.Tensors/Engines/DirectGpu/GpuBufferPool.cs (3)
6-10: LGTM!The
IPoolableGpuBufferinterface provides a clean contract for poolable GPU buffers withMarkRented()andRelease()lifecycle methods.
31-48: LGTM!The
TryRentmethod correctly usesVolatile.Readfor the disposal check andInterlocked.Decrementfor thread-safe count management. The compound conditional efficiently combines bucket lookup and buffer extraction.
76-92: LGTM!The
Disposemethod correctly usesInterlocked.Exchangefor idempotent disposal and properly releases all pooled buffers. The single-pass drain is appropriate for shutdown scenarios.src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cs (3)
84-113: LGTM!The
TryLoadLibraryimplementation correctly probes for the hipBLAS library across platforms. Loading and immediately freeing the handle is an appropriate pattern for availability checking without holding resources.
115-121: LGTM!The NETFRAMEWORK-specific kernel32 imports are correctly declared with appropriate attributes for Windows library loading.
15-33: HipBlasOperation enum values verified as correct.Both enums are correct. The
HipBlasOperationenum values (111, 112, 113) match the hipBLAS API definitions forHIPBLAS_OP_N,HIPBLAS_OP_T, andHIPBLAS_OP_Crespectively. The difference from CUDA's values (0, 1, 2) is expected since they are separate libraries with their own API specifications.AiDotNetBenchmarkTests/GpuResidentQuickHarness.cs (4)
17-74: LGTM! Well-structured entry point with defensive checks.The
Runmethod validates inputs, performs multiple GPU availability checks before proceeding, and logs diagnostic info. Resource upload validation is properly handled.
177-197: LGTM! Sound GPU timing methodology.Proper synchronization before warmup, between warmup and measurement, and after each iteration ensures accurate GPU timing without overlapping kernel executions.
199-217: LGTM! Deterministic data generation.The LCG constants (1664525, 1013904223) are well-known Numerical Recipes values, and the normalization to
[0, 1)is correct.
60-86: Verify whether CPU tensors require disposal.The CPU tensors (
matrixA,matrixB,vectorA,vectorB,convInput,convKernel) are created but never disposed, while their GPU counterparts are properly cleaned up in thefinallyblock.src/AiDotNet.Tensors/Helpers/BlasProvider.cs (5)
10-22: LGTM! Clean static field organization.Thread-safe initialization fields, delegate pointers, and environment configuration are well-organized. Reading env vars into static readonly fields ensures consistent behavior across the lifetime.
61-119: LGTM! Solid GEMM implementations with proper validation.Both overloads correctly validate bounds before unsafe operations, use proper pinning, and pass standard CBLAS parameters for
C = A * Bmultiplication.
179-185: LGTM! Reasonable partial success handling.Returning
trueif eithersgemmordgemmis loaded allows partial functionality when a library only exposes one precision.
233-247: LGTM! Correct overflow-safe bounds checking.Using
longforlastIndexcalculation prevents integer overflow when computingoffset + (rows-1)*stride + (cols-1)with large matrices.
310-339: LGTM! Clean platform abstraction for native library loading.The conditional compilation provides unified API over
NativeLibrary(.NET Core+) andkernel32P/Invoke (.NET Framework).src/AiDotNet.Tensors/Engines/CpuEngine.cs (7)
3428-3431: Conv2D early-exit via Im2Col path looks good.
Returning the preallocatedresultkeeps the API behavior consistent and enables the optimized path cleanly.
4398-4400: 2D matmul now routes through shared core — good.
This should make CPU matmul behavior consistent across call sites.
4436-4452: Batched matmul now uses shared core and avoids nested parallelism — good.
Parallel.Forover batches +allowParallel: falseinside is the right shape to prevent oversubscription.
4492-4510: Full batched matmul routing through shared core looks correct.
Same nested-parallelism avoidance applies here and should improve maintainability.
4530-4533: Nice fast-path: symmetric stride/pad/dilation delegates to optimized scalar overload.
This ensures the Im2Col path is reachable for the common square-parameter case.
4343-4385: Verify helper method parameter contracts. TheMultiplyMatrixCoreconsolidation is clean and row-major indexing is consistent, but the calls toMatrixMultiplyHelper.TryGemm(...)andMultiplyBlocked(...)require verification against their actual method signatures to ensure parameter ordering and types match correctly.
1404-1417: Potential recursion risk:MatrixMultiply<T>delegates toMatrix<T>.Multiply.Verify that
Matrix<T>.Multiply(...)does not call back intoIEngine.MatrixMultiply(...). If it does, this creates infinite recursion (StackOverflow) and changes which optimized code path executes.src/AiDotNet.Tensors/Engines/DirectGpuTensorEngine.cs (1)
1323-1382: Activation dispatcher refactor is consistent (out-of-place + in-place wrapper).
The overload split keeps call sites explicit about in-place vs out-of-place and thedefault:guard prevents silent future enum drift.src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (2)
7-10: New usings look appropriate for pooling (Interlocked) and shared pool types.
7671-7672: Disposing the pool before destroying the CUDA context is the right ordering.
This keeps_contextvalid for pooled buffer frees during pool teardown.src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (2)
8-9: Imports update looks correct (Threading + DirectGpu).
Needed forInterlockedusage inHipGpuBufferand for the pooling types.
649-667: [Rewritten review comment]
[Classification tag]
✏️ Tip: You can disable this entire section by setting review_details to false in your review settings.
| internal static extern HipBlasStatus hipblasSetStream(IntPtr handle, IntPtr stream); | ||
|
|
||
| [DllImport(HipBlasLibrary, EntryPoint = "hipblasSgemm")] | ||
| internal static extern HipBlasStatus hipblasSgemm( |
Check notice
Code scanning / CodeQL
Unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
| var status = HipBlasNative.hipblasSgemm( | ||
| _hipblasHandle, | ||
| HipBlasNative.HipBlasOperation.None, | ||
| HipBlasNative.HipBlasOperation.None, | ||
| N, M, K, | ||
| ref alphaVal, | ||
| bufferB.Handle, N, | ||
| bufferA.Handle, K, | ||
| ref betaVal, | ||
| bufferC.Handle, N); |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
|
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Fix all issues with AI agents
In `@AiDotNetBenchmarkTests/GpuResidentQuickHarness.cs`:
- Around line 65-74: If any of the UploadToContext calls (gpuMatrixA,
gpuMatrixB, gpuVectorA, gpuVectorB) returns null the method currently returns
early and leaks the previously-uploaded GPU buffers; before returning when a
null is detected, Dispose or otherwise release any non-null uploads (gpuMatrixA,
gpuMatrixB, gpuVectorA, gpuVectorB) to free GPU resources, or refactor the
uploads into a try/finally (or using) pattern so that the finally block always
disposes any non-null uploaded buffers.
♻️ Duplicate comments (5)
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cs (1)
11-54: Make the double-checked locking safe withvolatile(orVolatile.Read/Write).Without proper memory barriers, a thread can observe
_checkedAvailability == truewhile reading a stale_isAvailablevalue.🔒 Suggested fix
- private static bool _isAvailable; - private static bool _checkedAvailability; + private static volatile bool _isAvailable; + private static volatile bool _checkedAvailability;src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (1)
7851-7863: Make pooled Dispose resilient to pool teardown.
Dispose()directly calls_returnToPool(this); if the pool/backend is already disposed (orReturnthrows), device memory can leak or finalizers can fault. Add a try/catch fallback toRelease().🐛 Proposed fix
public void Dispose() { if (_returnToPool == null) { Release(); return; } if (Interlocked.CompareExchange(ref _poolState, 1, 0) != 0) return; - _returnToPool(this); + try + { + _returnToPool(this); + } + catch + { + // Pool already disposed or rejected; ensure device memory is freed. + Release(); + } }src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (3)
322-343: Avoid overflow in vendor GEMM work estimate.Line 341 multiplies into a signed
long; large dimensions can overflow and flip negative, incorrectly disabling vendor GEMM.♻️ Proposed fix
- long work = (long)m * n * k; - return work >= GetEnvLong(GemmVendorThresholdEnvVar, DefaultVendorGemmThreshold); + ulong work = (ulong)m * (ulong)n * (ulong)k; + ulong threshold = (ulong)GetEnvLong(GemmVendorThresholdEnvVar, DefaultVendorGemmThreshold); + return work >= threshold;
877-919: Verify hipBLAS SGEMM argument ordering/layout.Line 901+ swaps
(M, N)and passesBbeforeA. That can be correct for row‑major adaptation, but if the layout/lda/ldb/ldc assumptions don’t match the kernel path, results are silently wrong only on the vendor path.#!/bin/bash set -euo pipefail fd -t f -e cs 'HipBlasNative|hipblas' src || true rg -n --type=cs -C2 'hipblasSgemm|HipBlasOperation|hipblas(SetStream|Create|Destroy)' src rg -n --type=cs -C2 'Gemm\\(|MatMul\\(|\\blda\\b|\\bldb\\b|\\bldc\\b|row-major|col-major' src
9182-9186: Guard pool returns after backend disposal.
_bufferPool.Dispose()can run while buffers are still alive;HipGpuBuffer.Dispose()always returns to the pool when a callback exists. That can call into a disposed pool and crash. Consider a return‑or‑free callback and set_disposedbefore pool disposal.🛠️ Suggested hardening
public sealed class HipBackend : IAsyncGpuBackend { + private void ReturnBufferToPool(HipGpuBuffer buffer) + { + if (_disposed) + { + buffer.Release(); + return; + } + _bufferPool.Return(buffer); + } ... public IGpuBuffer AllocateBuffer(float[] data) { ... - return new HipGpuBuffer(devicePtr, data.Length, _bufferPool.Return); + return new HipGpuBuffer(devicePtr, data.Length, ReturnBufferToPool); } public IGpuBuffer AllocateBuffer(int size) { ... - return new HipGpuBuffer(devicePtr, size, _bufferPool.Return); + return new HipGpuBuffer(devicePtr, size, ReturnBufferToPool); } public void Dispose() { if (_disposed) return; + _disposed = true; // Dispose the default stream wrapper (does not destroy underlying stream) _defaultStream?.Dispose(); _defaultStream = null; _bufferPool.Dispose(); ... - _kernelCache.Clear(); - _disposed = true; + _kernelCache.Clear(); } }Also applies to: 9878-9890
🧹 Nitpick comments (3)
scripts/snapshot-cpu-benchmarks.ps1 (1)
16-24: Consider recursion if reports can land in subfolders.
If BenchmarkDotNet outputs can be nested in this repo, the current non-recursive search will miss them. Please verify the output layout; if nested, add-Recurseand exclude the snapshots path to avoid re-copying.♻️ Possible update if nested outputs exist
-$files = foreach ($pattern in $patterns) { - Get-ChildItem -Path $Source -Filter $pattern -File -ErrorAction SilentlyContinue -} +$files = foreach ($pattern in $patterns) { + Get-ChildItem -Path $Source -Filter $pattern -File -Recurse -ErrorAction SilentlyContinue | + Where-Object { $_.FullName -notlike "$DestinationRoot\*" } +}src/AiDotNet.Tensors/Helpers/BlasProvider.cs (1)
121-139: Consider usingvolatileorLazy<T>for robust double-checked locking.The double-checked locking pattern here can have subtle memory ordering issues on certain runtimes or architectures. While it typically works on x86/x64 with modern .NET, marking
_initializedasvolatileor usingLazy<T>provides stronger guarantees and clearer intent.Option 1: Add volatile
- private static bool _initialized; - private static bool _available; + private static volatile bool _initialized; + private static bool _available;Option 2: Use Lazy<T> for simpler thread-safe initialization
private static readonly Lazy<bool> _lazyAvailable = new Lazy<bool>(TryLoadLibrary); private static bool EnsureInitialized() => _lazyAvailable.Value;src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs (1)
192-205: Consider using consistent null validation style.The null checks here use explicit
if (x is null) throwpattern, while theActionoverload at line 174 usesArgumentNullException.ThrowIfNull(). Both are functionally equivalent, but consistency improves readability.♻️ Optional: Align with ThrowIfNull pattern
public static TResult WithGpuContext<TResult>(Func<GpuExecutionContext, TResult> func, Func<TResult> fallback, GpuExecutionOptions? options = null) { - if (func is null) - throw new ArgumentNullException(nameof(func)); - if (fallback is null) - throw new ArgumentNullException(nameof(fallback)); + ArgumentNullException.ThrowIfNull(func); + ArgumentNullException.ThrowIfNull(fallback); if (Current is DirectGpuTensorEngine gpuEngine) {
📜 Review details
Configuration used: Organization UI
Review profile: CHILL
Plan: Pro
📒 Files selected for processing (8)
AiDotNetBenchmarkTests/GpuResidentQuickHarness.csscripts/snapshot-cpu-benchmarks.ps1src/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.cs
🧰 Additional context used
🧠 Learnings (3)
📚 Learning: 2025-11-19T04:08:26.895Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 0
File: :0-0
Timestamp: 2025-11-19T04:08:26.895Z
Learning: For ILGPU GPU operations in GpuEngine.cs, use standard .NET exception types (InvalidOperationException, ArgumentException, OutOfMemoryException) instead of ILGPU-specific exception types, as ILGPU exception types may be version-specific. Combine with message-based filtering using ex.Message.Contains("device") or ex.Message.Contains("accelerator") as a fallback for GPU-specific errors.
Applied to files:
src/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs
📚 Learning: 2025-12-18T08:49:25.295Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 444
File: src/Interfaces/IPruningMask.cs:1-102
Timestamp: 2025-12-18T08:49:25.295Z
Learning: In the AiDotNet repository, the project-level global using includes AiDotNet.Tensors.LinearAlgebra via AiDotNet.csproj. Therefore, Vector<T>, Matrix<T>, and Tensor<T> are available without per-file using directives. Do not flag missing using directives for these types in any C# files within this project. Apply this guideline broadly to all C# files (not just a single file) to avoid false positives. If a file uses a type from a different namespace not covered by the global using, flag as usual.
Applied to files:
src/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.csAiDotNetBenchmarkTests/GpuResidentQuickHarness.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs
📚 Learning: 2025-12-18T08:49:53.103Z
Learnt from: ooples
Repo: ooples/AiDotNet PR: 444
File: src/Interfaces/IPruningStrategy.cs:1-4
Timestamp: 2025-12-18T08:49:53.103Z
Learning: In this repository, global using directives are declared in AiDotNet.csproj for core namespaces (AiDotNet.Tensors.* and AiDotNet.*) and common system types. When reviewing C# files, assume these global usings are in effect; avoid adding duplicate using statements for these namespaces and for types like Vector<T>, Matrix<T>, Tensor<T>, etc. If a type is not found, verify the global usings or consider adding a file-scoped using if needed. Prefer relying on global usings to reduce boilerplate.
Applied to files:
src/AiDotNet.Tensors/Engines/AiDotNetEngine.cssrc/AiDotNet.Tensors/Helpers/BlasProvider.csAiDotNetBenchmarkTests/GpuResidentQuickHarness.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cssrc/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cssrc/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs
⏰ Context from checks skipped due to timeout of 90000ms. You can increase the timeout in your CodeRabbit configuration to a maximum of 15 minutes (900000ms). (2)
- GitHub Check: CodeQL Analysis
- GitHub Check: Build (Windows)
🔇 Additional comments (42)
scripts/snapshot-cpu-benchmarks.ps1 (4)
1-5: Clean parameterization and sensible defaults.
7-10: Good early exit on missing source path.
12-15: Destination setup is straightforward and safe.
26-32: Nice guard for empty results and clear output of snapshot path.src/AiDotNet.Tensors/Helpers/BlasProvider.cs (10)
8-22: LGTM!Clean static field setup with appropriate types for lazy initialization and native interop. Reading configuration from environment variables at static initialization time is a reasonable approach for optional BLAS support.
35-59: LGTM!Delegate signatures correctly match the CBLAS interface with
CallingConvention.Cdecl. The unsafe pointer parameters are appropriate for high-performance GEMM operations.
61-89: LGTM!Proper validation via
HasEnoughDatabefore the unsafe block. The GEMM call withalpha=1.0fandbeta=0.0fcorrectly performsC = A * B. Pointer arithmetic with offsets is handled correctly.
91-119: LGTM!Consistent implementation with the single-precision version. Bounds validation and pointer handling are correct.
141-181: LGTM!The library loading logic correctly handles the explicit path case and iterates through candidates. The handle is properly freed when symbol loading fails, addressing the previous review feedback.
183-189: Verify intent: provider is available if either precision works.The condition
_sgemm != null || _dgemm != nullmeans the provider is marked available even if only one precision loaded successfully. This is fine since eachTryGemmoverload checks its specific delegate, but callers should be aware that availability doesn't guarantee both precisions work.
199-232: LGTM!Good exception handling with specific catch blocks for expected P/Invoke failure modes. Silently ignoring failures is appropriate here since BLAS threading configuration is best-effort and non-critical.
249-263: LGTM!Bounds validation is correct. The use of
longarithmetic prevents integer overflow for large matrix dimensions, and the formula correctly computes the last accessed index for row-major layout.
265-300: LGTM!Comprehensive library name candidates covering common naming conventions across Windows, Linux, and macOS. The MKL preference logic is configurable via environment variable with a reasonable platform-based default.
302-355: LGTM!Clean helper functions for environment variable parsing and platform-specific native library operations. The
#if NETFRAMEWORKseparation keeps the code maintainable for both legacy and modern .NET targets.src/AiDotNet.Tensors/Engines/AiDotNetEngine.cs (3)
1-2: LGTM!The import is required for the new GPU context types (
GpuExecutionContext,GpuExecutionOptions,DirectGpuTensorEngine) used in the methods below.
151-164: LGTM!The method correctly delegates to
DirectGpuTensorEnginewhen available and returnsnullotherwise. The nullable return type properly signals GPU unavailability to callers.
166-182: LGTM!The null validation for
actionhas been added using the modernArgumentNullException.ThrowIfNullpattern. The method correctly returnsfalsewhen GPU context is unavailable, allowing callers to handle fallback logic.src/AiDotNet.Tensors/Engines/DirectGpu/OpenCL/OpenClBackend.cs (13)
9-10: Imports for Interlocked/pooling are appropriate.
40-43: Buffer pool wiring looks fine.
52-55: Vendor GEMM env flags and defaults are clear.
545-573: Kernel-aware local-size clamping is a solid safeguard.
614-662: Vendor GEMM gating logic is reasonable.
1817-1836: CLBlast vendor path gating looks good.
2653-2658: Reduction local-size clamp order is correct.
2684-2687: Consistent clamp for reduce_max is good.
4139-4173: CLBlast SGEMM error handling is solid.
4175-4186: CLBlast benchmark wrapper behavior is clear.
9275-9275: Pool disposal on backend dispose is correct.
9298-9349: Poolable buffer lifecycle changes look correct.
769-794: Verify pooled rent resets state and size matches request.
IfGpuBufferPool.TryRentdoesn't callMarkRented()(or if it returns a larger buffer than requested),Dispose()can skip returning the buffer or downstream ops may observe an unexpectedSize. Please confirm pool behavior; if it doesn't reset state, add a reset here.🔧 Optional safeguard (only if TryRent doesn't already MarkRented)
- if (_bufferPool.TryRent(data.Length, out var pooled) && pooled != null) + if (_bufferPool.TryRent(data.Length, out var pooled) && pooled != null) { + pooled.MarkRented(); pooled.Buffer.CopyFromHost(data); return pooled; }- if (_bufferPool.TryRent(size, out var pooled) && pooled != null) - return pooled; + if (_bufferPool.TryRent(size, out var pooled) && pooled != null) + { + pooled.MarkRented(); + return pooled; + }AiDotNetBenchmarkTests/GpuResidentQuickHarness.cs (6)
17-59: GPU availability gating and context setup look solid.Clamping inputs and early exits keep the harness safe to run across environments.
88-143: Good benchmark mix and reuse coverage.The sequence exercises both resident ops and reused outputs, which is helpful for perf comparisons.
158-175: Deterministic Conv2D input/kernel creation is clean.Straightforward shapes and seeded data make results stable across runs.
177-197: Timing loop is clean and synchronized.The synchronization around warmup and measured loops keeps GPU timings consistent.
199-216: Deterministic data generation looks good.The seeded LCG provides stable inputs without extra dependencies.
219-234: Null-safe composite disposal is well handled.Disposing in reverse order with null guards avoids leaks and exceptions.
src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (1)
332-348: VerifyTryRentcallsMarkRentedfor pooled buffers.Both allocation overloads return pooled buffers without explicitly resetting
_poolState. IfGpuBufferPool.TryRentdoes not callMarkRented,Dispose()will early-exit (state ≠ 0) and the buffer won't return to the pool or be freed.Also applies to: 380-389
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (5)
8-9: No issues noted in these additions.Also applies to: 45-48, 68-75, 211-212, 271-320
573-576: No issues noted in these comments.Also applies to: 595-598, 643-649
770-773: No issues noted in these changes.Also applies to: 9255-9257, 9270-9275
9843-9876: No issues noted in these poolable buffer additions.
667-733: Verify pooled buffers are reset/marked as rented on reuse.Please confirm
GpuBufferPool.TryRent(...)callsMarkRented()(or otherwise resets pool state) and enforces size invariants; otherwise pooled buffers can be double-returned or carry an incorrectSize.
✏️ Tip: You can disable this entire section by setting review_details to false in your review settings.
| fixed (float* src = data) | ||
| { | ||
| CuBlasNative.CheckCudaResult( | ||
| CuBlasNative.cuMemcpyHtoD(pooled.Handle, (IntPtr)src, byteSize), // lgtm[cs/call-to-unmanaged-code] CUDA interop requires native driver calls. |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
| if (_bufferPool.TryRent(size, out var pooled) && pooled != null) | ||
| { | ||
| CuBlasNative.CheckCudaResult( | ||
| CuBlasNative.cuMemsetD32(pooled.Handle, 0, (ulong)size), // lgtm[cs/call-to-unmanaged-code] CUDA interop requires native driver calls. |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
|
|
||
| try | ||
| { | ||
| var status = HipBlasNative.hipblasCreate(ref _hipblasHandle); // lgtm[cs/call-to-unmanaged-code] HIP BLAS uses native bindings. |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
| return; | ||
| } | ||
|
|
||
| status = HipBlasNative.hipblasSetStream(_hipblasHandle, _stream); // lgtm[cs/call-to-unmanaged-code] HIP BLAS uses native bindings. |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
|
|
||
| if (_hipblasHandle != IntPtr.Zero) | ||
| { | ||
| HipBlasNative.hipblasDestroy(_hipblasHandle); // lgtm[cs/call-to-unmanaged-code] HIP BLAS uses native bindings. |
Check notice
Code scanning / CodeQL
Calls to unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
| } | ||
|
|
||
| [DllImport(HipBlasLibrary, EntryPoint = "hipblasCreate")] | ||
| internal static extern HipBlasStatus hipblasCreate(ref IntPtr handle); // lgtm[cs/unmanaged-code] HIP BLAS requires native bindings. |
Check notice
Code scanning / CodeQL
Unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
| internal static extern HipBlasStatus hipblasCreate(ref IntPtr handle); // lgtm[cs/unmanaged-code] HIP BLAS requires native bindings. | ||
|
|
||
| [DllImport(HipBlasLibrary, EntryPoint = "hipblasDestroy")] | ||
| internal static extern HipBlasStatus hipblasDestroy(IntPtr handle); // lgtm[cs/unmanaged-code] HIP BLAS requires native bindings. |
Check notice
Code scanning / CodeQL
Unmanaged code Note
Copilot Autofix
AI 9 months ago
Copilot could not generate an autofix suggestion
Copilot could not generate an autofix suggestion for this alert. Try pushing a new commit or if the problem persists contact support.
| internal static extern HipBlasStatus hipblasDestroy(IntPtr handle); // lgtm[cs/unmanaged-code] HIP BLAS requires native bindings. | ||
|
|
||
| [DllImport(HipBlasLibrary, EntryPoint = "hipblasSetStream")] | ||
| internal static extern HipBlasStatus hipblasSetStream(IntPtr handle, IntPtr stream); // lgtm[cs/unmanaged-code] HIP BLAS requires native bindings. |
Check notice
Code scanning / CodeQL
Unmanaged code Note
There was a problem hiding this comment.
Actionable comments posted: 2
🤖 Fix all issues with AI agents
In `@AiDotNetBenchmarkTests/GpuResidentQuickHarness.cs`:
- Around line 60-157: The GPU uploads (gpuEngine.UploadToContext calls producing
gpuMatrixA/gpuMatrixB/gpuVectorA/gpuVectorB and gpuConvInput) must be protected
from exceptions so they don't leak; move the try block to start immediately
after the UploadToContext successes and perform subsequent AllocateBuffer calls
(addOutput, multiplyOutput) and conv setup inside that try, or preferably create
those buffers with using (or ensure they are disposed in the same finally) so
any exception from AllocateBuffer, conv setup, or later code triggers cleanup;
ensure gpuConvInput upload is also guarded so its exception is handled and that
all uploaded tensors (gpuMatrixA,gpuMatrixB,gpuVectorA,gpuVectorB,gpuConvInput)
are disposed in the finally if they were created.
In `@src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cs`:
- Around line 59-83: The DllImport declarations for hipBLAS methods
(hipblasCreate, hipblasDestroy, hipblasSetStream, hipblasSgemm) must explicitly
set the calling convention to cdecl to avoid a StdCall/default mismatch on
32-bit Windows; update each DllImport attribute to include CallingConvention =
CallingConvention.Cdecl so the native C API is invoked correctly across
platforms and the signatures remain unchanged.
♻️ Duplicate comments (1)
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
780-783: Re‑verify hipBLAS SGEMM argument order/layout.This is the same concern previously raised: the row‑major adaptation (A/B swap and M/N swap) needs validation to ensure results match the kernel path for all shapes/strides.
Also applies to: 887-930
🧹 Nitpick comments (3)
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBlasNative.cs (1)
95-95: PreferAny(CanLoadLibrary)for clarity and slight efficiency.This avoids creating a filtered enumerable and is a bit more idiomatic.
Refactor
- return candidates.Where(CanLoadLibrary).Any(); + return candidates.Any(CanLoadLibrary);src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
8-9: Drop redundant DirectGpu using to align with global usings.
AiDotNet.Tensors.*is globally imported in this repo, so this using is likely redundant. Keeping onlySystem.Threadingshould be sufficient.💡 Suggested cleanup
-using AiDotNet.Tensors.Engines.DirectGpu;Based on learnings, please confirm the global usings still cover this namespace.
src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (1)
7797-7873: Make_poolStatevalues self‑documenting.
The 0/1/2 state values are easy to misread. Consider named constants (or an enum) to clarify state transitions.♻️ Optional refactor
- private int _poolState; + private const int PoolStateInUse = 0; + private const int PoolStateReturned = 1; + private const int PoolStateReleased = 2; + private int _poolState = PoolStateInUse; ... - Interlocked.Exchange(ref _poolState, 0); + Interlocked.Exchange(ref _poolState, PoolStateInUse); ... - if (Interlocked.Exchange(ref _poolState, 2) == 2) + if (Interlocked.Exchange(ref _poolState, PoolStateReleased) == PoolStateReleased) return; ... - if (Interlocked.CompareExchange(ref _poolState, 1, 0) != 0) + if (Interlocked.CompareExchange(ref _poolState, PoolStateReturned, PoolStateInUse) != PoolStateInUse) return;
- Replace MemoryMarshal.Cast<T, TTarget> with MemoryMarshal.CreateSpan pattern using Unsafe.As to work around struct constraint requirement on unconstrained generic T - Wrap SIMD-optimized ReLU paths in #if NET5_0_OR_GREATER since MemoryMarshal.CreateSpan is not available in net471 - Replace ArgumentNullException.ThrowIfNull with standard null check for net471 compatibility Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Make MultiplyMatrixCore fallback loop use Parallel.For when allowParallel is true - This matches master's behavior which always parallelizes TensorMatMul2D - Increase MatMulTransposeMaxProduct from 262,144 to 1,048,576 (handles up to 1024x1024) - Increase MatMulTransposeMaxElements from 2M to 4M for better SIMD utilization - Decrease MatMulTransposeParallelThreshold from 32,768 to 16,384 for earlier parallelization The regression was caused by the fallback path running sequentially while master always used Parallel.For for row iteration. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Fix all issues with AI agents
In `@src/AiDotNet.Tensors/Engines/CpuEngine.cs`:
- Around line 3775-3776: Comment in CpuEngine.cs near the
MathHelper.GetNumericOperations<T>() call contains corrupted UTF-8 characters;
replace the garbled sequence with the intended multiplication symbol (e.g.,
change "3-6Ã..." to "3–6× speedup for float") so the comment reads clearly
(e.g., "Use SIMD-optimized Tanh (3–6× speedup for float)"), update any other
occurrences of the same garbled text in the file, and ensure the source file is
saved with UTF-8 encoding to prevent reintroduction of the double-encoding
issue.
🧹 Nitpick comments (2)
src/AiDotNet.Tensors/Engines/CpuEngine.cs (2)
1849-1871: Repeated parallel chunking pattern — consider extracting a helper.This exact pattern (threshold check, chunking,
Parallel.Forwith bounds validation) is repeated in Add, Subtract, Multiply, Tanh, Sigmoid, ReLU, and Sum. Extracting a generic helper likeParallelChunked(int length, Action<int start, int count> body)would reduce ~30 lines of boilerplate per operation and centralize the parallelization logic.♻️ Example helper extraction
private static void ParallelChunked(int length, int threshold, Action<int, int> processChunk) { if (length >= threshold && Environment.ProcessorCount > 1) { int workerCount = Environment.ProcessorCount; int chunkSize = (length + workerCount - 1) / workerCount; Parallel.For(0, workerCount, worker => { int start = worker * chunkSize; if (start >= length) return; int count = Math.Min(chunkSize, length - start); processChunk(start, count); }); } else { processChunk(0, length); } }Usage:
ParallelChunked(length, ElementwiseParallelThreshold, (start, count) => numOps.Add(aData.Span.Slice(start, count), bData.Span.Slice(start, count), resultData.Span.Slice(start, count)));
3920-4019: Significant code duplication in Tensor ReLU — consider refactoring.The float and double SIMD paths (lines 3926-3966 and 3968-4008) are nearly identical, differing only in the type cast. This ~80-line block could be reduced to ~30 lines by extracting a generic SIMD dispatch helper or using a local function with generics.
♻️ Sketch of refactored approach
private static void ApplyReLUSimd<TElement>( ReadOnlySpan<TElement> input, Span<TElement> output, int length, int threshold) where TElement : unmanaged { if (length >= threshold && Environment.ProcessorCount > 1) { int workerCount = Environment.ProcessorCount; int chunkSize = (length + workerCount - 1) / workerCount; Parallel.For(0, workerCount, worker => { int start = worker * chunkSize; if (start >= length) return; int count = Math.Min(chunkSize, length - start); SimdKernels.ReLU(input.Slice(start, count), output.Slice(start, count)); }); } else { SimdKernels.ReLU(input, output); } }Then in Tensor ReLU:
if (typeof(T) == typeof(float)) { ApplyReLUSimd( MemoryMarshal.Cast<T, float>(inputData.Span), MemoryMarshal.Cast<T, float>(resultData.Span), length, ElementwiseParallelThreshold); return result; }
…ation - Use GC.AllocateUninitializedArray<T> for faster large array allocation (~30-50% faster for large arrays by avoiding zeroing) - Add zero-copy TensorBase constructor from Vector<T> - Add Tensor<T>.FromMemory() static method for zero-copy tensor creation - These optimizations reduce allocation overhead vs native libraries Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Fix all issues with AI agents
In `@src/AiDotNet.Tensors/LinearAlgebra/VectorBase.cs`:
- Around line 64-70: The constructor in VectorBase currently uses
GC.AllocateUninitializedArray<T>(length) which breaks the documented
default-zero semantics for the _memory field; change the default constructor in
VectorBase to allocate with new T[length] to preserve zero-initialization and
XML contract, and introduce an explicit opt-in path for performance (either an
internal/static factory like VectorBase.CreateUninitialized or an
internal/optional constructor parameter e.g. bool uninitialized) that uses
GC.AllocateUninitializedArray<T> for callers that can guarantee full
initialization; update the XML docs on VectorBase and the new opt-in API to note
the unsafe/uninitialized behavior and audit callers to use the new opt-in API if
they fully initialize before exposure.
🧹 Nitpick comments (1)
src/AiDotNet.Tensors/LinearAlgebra/TensorBase.cs (1)
142-150: Consider adding null validation for thedataparameter.If
datais null, the access to_data.Lengthon line 146 will throw aNullReferenceException, which is less informative than anArgumentNullException. While this is a protected constructor intended for internal/derived class use, explicit validation improves debuggability.Suggested improvement
protected TensorBase(Vector<T> data, int[] shape) { + ArgumentNullException.ThrowIfNull(data); Shape = shape; _data = data; if (_data.Length != shape.Aggregate(1, (acc, dim) => acc * dim)) { throw new ArgumentException("The number of values does not match the specified shape."); } }
- Add AiDotNet.Native.OpenBLAS project with bundled OpenBLAS DLL (48MB) - Modify BlasProvider to search assembly directory, AppContext.BaseDirectory, and current directory for BLAS libraries - Add OpenBLAS reference to Tensors project and benchmark project - Results: 512x512 MatMul: 1886us -> 799us (58% faster, now only 1.8x slower than TorchSharp) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Fix corrupted UTF-8 encoding in CpuEngine.cs comments (x speedup for float) - Add CallingConvention.Cdecl to HipBlasNative.cs DllImport attributes for 32-bit Windows compatibility - Update VectorBase.cs documentation to clarify uninitialized allocation behavior on .NET 5+ Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Use System.Numerics.Tensors.TensorPrimitives directly for double type operations on .NET 8+ instead of going through TensorPrimitivesCore with struct-based dispatch overhead. Operations optimized: - Add, Subtract, Multiply, Divide (elementwise) - Dot product and Sum (reductions) - Exp, Log, Log2, Tanh (transcendentals) - Sigmoid (activation function) - Max, Min (comparisons) - SoftMax, CosineSimilarity (specialized) This matches the existing FloatOperations pattern and eliminates the 18-20x performance gap for double-precision operations. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Merge origin/master into fix/issue-682-gpu-perf to bring the following: - oneDNN integration for Conv2D with brgconv:avx2 algorithm - SIMD-optimized matrix multiply helpers (TryGemm, blocked multiply) - New Conv2D helpers (Im2Col, DirectConv2D, Winograd, FusedConv) - TensorPrimitives optimizations for float operations - LSTM/GRU kernel implementations for CUDA, HIP, and OpenCL - InProcess benchmark mode for native DLL testing Conflict resolutions: - CudaBackend: merged buffer pooling with cooperative launch - HipBackend: merged hipblas with LSTM/GRU module cleanup - Program.cs: merged GPU harness with InProcess benchmark mode - AiDotNetBenchmarkTests.csproj: kept both OpenBLAS and oneDNN Additional fix: - Removed incomplete TryMultiplyPacked call from MatrixBase.cs Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Fix all issues with AI agents
In `@src/AiDotNet.Tensors/NumericOperations/DoubleOperations.cs`:
- Around line 803-855: Update the misleading remark in the XML doc for the Add
method (and similarly worded remarks for Subtract/Multiply/Divide) to state that
the NET8_0_OR_GREATER guard enables hardware-accelerated TensorPrimitives and
that the fallback via TensorPrimitivesCore.InvokeSpanSpanIntoSpan is used for
all pre-.NET 8 targets (e.g., .NET 6/7 and .NET Framework), referencing the
methods Add/Subtract/Multiply/Divide and the
TensorPrimitives/TensorPrimitivesCore calls to ensure the documentation
accurately reflects behavior across all target frameworks.
♻️ Duplicate comments (1)
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
10306-10324: Don’t swallow pool-return failures silently.The generic catch in
HipGpuBuffer.Dispose()hides pool/interop issues; consider at least logging the exception.
🧹 Nitpick comments (4)
src/AiDotNet.Tensors/LinearAlgebra/MatrixBase.cs (1)
701-720: Good algorithm selection hierarchy; consider reducing allocations in recursive fallback.The BLAS → blocked → recursive hierarchy is well-structured. However, the recursive fallback path allocates three arrays:
resultData(line 716), plus two.ToArray()calls (line 717). Sinceresult._memoryis already allocated at line 696, consider refactoringMultiplyRecursiveto acceptMemory<T>orSpan<T>to avoid the extra allocation and copy.♻️ Suggested optimization for recursive path
// Use cache-oblivious recursive algorithm MatrixMultiplyHelper.TraceMatmul("RECURSIVE", M, N, K); - var resultData = new T[M * N]; - MultiplyRecursive(_memory.ToArray(), other._memory.ToArray(), resultData, 0, 0, 0, 0, 0, 0, M, K, N, K, N, N); - - // Copy result back to the result matrix - resultData.AsSpan().CopyTo(result.AsWritableSpan()); + // Consider refactoring MultiplyRecursive to work with Memory<T> directly + // to avoid these allocations and the final copy + var aData = _memory.ToArray(); + var bData = other._memory.ToArray(); + var cData = result._memory.ToArray(); + MultiplyRecursive(aData, bData, cData, 0, 0, 0, 0, 0, 0, M, K, N, K, N, N); + cData.AsSpan().CopyTo(result.AsWritableSpan());Note: A more substantial refactor would modify
MultiplyRecursiveto work withMemory<T>directly, eliminating the need for array conversions altogether.AiDotNetBenchmarkTests/Program.cs (1)
93-117: Minor: Silent fallback for invalid values.The parsing logic is correct, but invalid or non-positive values (e.g.,
--iterations 0or--iterations abc) silently fall back to defaults. Consider logging a warning for visibility during debugging.🔧 Optional: Add warning for invalid values
var valueText = args[index + 1]; args.RemoveAt(index + 1); args.RemoveAt(index); if (int.TryParse(valueText, out var value) && value > 0) { return value; } + Console.WriteLine($"Warning: Invalid value '{valueText}' for {option}, using default {defaultValue}"); return defaultValue;src/AiDotNet.Tensors/NumericOperations/DoubleOperations.cs (1)
1191-1249: Consider a shared guard for span-length checks
The new validations are great; to reduce duplication, consider a small helper (and optionally include parameter names for clearer exceptions).♻️ Possible refactor
+ private static void EnsureSameLength(ReadOnlySpan<double> x, Span<double> destination) + { + if (x.Length != destination.Length) + throw new ArgumentException("Spans must have the same length", nameof(destination)); + } public void LeakyReLU(ReadOnlySpan<double> x, double alpha, Span<double> destination) { - if (x.Length != destination.Length) - throw new ArgumentException("Spans must have the same length"); + EnsureSameLength(x, destination); Engines.Simd.SimdKernels.LeakyReLU(x, alpha, destination); }(Apply the same guard to GELU/Mish/Swish/ELU/ReLU.)
src/AiDotNet.Tensors/Engines/DirectGpu/CUDA/CudaBackend.cs (1)
43-46: Consider making pool limits configurable.
Hard-coded pool sizes may be suboptimal across GPUs; a ctor option or settings hook would let users tune memory/perf trade-offs.
Add parallel processing using CpuParallelSettings.ParallelForChunks for: - TensorAddInPlace: ~2.7x faster for 1M elements - TensorMultiplyInPlace: parallel SIMD for large tensors - SigmoidInPlace: parallel TensorPrimitives.Sigmoid (NET8+) - ReLUInPlace: parallel ReLU processing (NET8+) Operations use ParallelThreshold (65536 elements) to avoid overhead on small tensors. Double support uses NET5_0_OR_GREATER for TensorPrimitives compatibility. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Add Conv2DInto method to CpuEngine for zero-allocation Conv2D - Add AiDotNet_Conv2D_ZeroAlloc benchmark - Pre-allocate conv output tensor in benchmark setup Results: - AiDotNet Conv2D (zero-alloc): 178.5 µs, 0 B allocated - TorchSharp Conv2D: 245.4 µs, 48 B allocated - WE WIN by 1.37x on Conv2D with zero allocations! Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
There was a problem hiding this comment.
CodeQL found more than 20 potential problems in the proposed changes. Check the Files changed tab for more details.
- Fix generic catch clauses in CudaBackend.cs and HipBackend.cs (now catch Exception explicitly) - Fix GPU resource leak guard in GpuResidentQuickHarness.cs (moved try block earlier) - Fix DoubleOperations.cs docstring (pre-.NET 8 instead of .NET Framework) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Fix all issues with AI agents
In `@AiDotNetBenchmarkTests/Benchmarking/FixedProjectFileToolchain.cs`:
- Around line 56-78: The code assumes a Windows-only executable name
"dotnet.exe", breaking Linux/macOS; update FixedProjectFileToolchain to choose
the executable name based on platform (e.g., use
System.Runtime.InteropServices.RuntimeInformation and OSPlatform) and replace
all hardcoded "dotnet.exe" occurrences with a variable like executableName =
RuntimeInformation.IsOSPlatform(OSPlatform.Windows) ? "dotnet.exe" : "dotnet",
then use Path.GetFileName(processPath) == executableName and
Path.Combine(dotNetRoot, executableName) / Path.Combine(programFiles, "dotnet",
executableName) respectively so path resolution works cross-platform.
♻️ Duplicate comments (1)
src/AiDotNet.Tensors/Engines/DirectGpu/HIP/HipBackend.cs (1)
981-992: Validate hipBLAS SGEMM argument order for row‑major adaptation.
The swap of A/B and M/N is easy to get wrong and will silently transpose or mis‑stride results. Please re‑verify the layout/lda/ldb/ldc mapping against the hipBLAS signature.
🧹 Nitpick comments (14)
.deps/openblas/include/lapacke_mangling.h (1)
1-2: Prefer a more specific include guard name.This guard is a bit generic and could collide with other LAPACK headers in larger builds. Consider a filename‑specific guard.
♻️ Proposed change
-#ifndef LAPACK_HEADER_INCLUDED -#define LAPACK_HEADER_INCLUDED +#ifndef LAPACKE_MANGLING_H_INCLUDED +#define LAPACKE_MANGLING_H_INCLUDEDsrc/AiDotNet.Tensors/Engines/CpuEngine.cs (1)
4306-4311: Misleading comment and missed SIMD opportunity.The comment states "TensorPrimitives.Max with zeros is fastest" but the implementation uses a scalar loop that won't auto-vectorize well. Consider using
Vector<float>operations orTensorPrimitives.Maxwith a pre-allocated zero span for consistent SIMD performance.♻️ Suggested vectorized approach using MathF.Max
var floatSpan = ((Memory<float>)(object)data).Span.Slice(start, count); - // ReLU = max(x, 0), TensorPrimitives.Max with zeros is fastest - for (int i = 0; i < count; i++) - { - if (floatSpan[i] < 0) floatSpan[i] = 0; - } + // ReLU = max(x, 0) + for (int i = 0; i < count; i++) + { + floatSpan[i] = MathF.Max(floatSpan[i], 0f); + }Alternatively, for better SIMD, consider
System.Runtime.Intrinsicsor vectorized primitives if available.src/AiDotNet.Tensors/Engines/Optimized/IOptimizedTensor.cs (1)
42-53: Consider read‑only shape exposure and a 64‑bit element count.
int[] Shapecan allow external mutation, andint ElementCountcan overflow for large tensors. ConsiderIReadOnlyList<int>/ReadOnlyMemory<int>for shape andlongfor element counts to make the API safer for large models.Also applies to: 115-121
src/AiDotNet.Tensors/Engines/Optimized/ICpuOptimizedTensor.cs (1)
42-82: Consider SafeHandle wrapper and enum for format tags per oneDNN ownership semantics.Exposing raw
IntPtrhandles violates oneDNN's explicit ownership contract: memory descriptors returned from the C API must be destroyed by the caller viadnnl_memory_desc_destroy. Without a documented ownership/lifetime contract or SafeHandle-like wrapper, callers risk resource leaks or use-after-free bugs. Additionally, oneDNN definesdnnl_format_tag_tas a proper C enum; wrappingint OneDnnFormatTagas a strongly typed enum eliminates magic values and improves type safety.AiDotNetBenchmarkTests/TorchSharpCpuComparisonBenchmarks.cs (1)
299-303: Zero-allocation benchmark result is not consumed.The benchmark returns
voidand doesn't consume or return the output tensor, which may allow the JIT/runtime to optimize away the computation in some scenarios. For consistency with other benchmarks (which return results), consider returning_aiConvOutputor using the consumer pattern.Proposed fix
[Benchmark] - public void AiDotNet_Conv2D_ZeroAlloc() + public Tensor<float> AiDotNet_Conv2D_ZeroAlloc() { _cpuEngine.Conv2DInto(_aiConvOutput!, _aiConvInput!, _aiConvKernel!, _convStride, _convPadding, _convDilation); + return _aiConvOutput!; }AiDotNetBenchmarkTests/CpuProfilingHarness.cs (2)
95-105: Same disposal consideration applies to MatMul results.Similar to Conv2D, the MatMul results are not disposed. Apply
usingifTensor<float>is disposable.Proposed fix (if Tensor is disposable)
if (options.Mode is CpuProfilingMode.All or CpuProfilingMode.MatMul) { for (int iter = 0; iter < options.Iterations; iter++) { foreach (var size in MatMulSizes) { - var result = AiDotNetEngine.Current.TensorMatMul(matA[size], matB[size]); - _sink = result.GetFlat(0); + using var result = AiDotNetEngine.Current.TensorMatMul(matA[size], matB[size]); + _sink = result.GetFlat(0); } } }
167-174: Minor inconsistency with other benchmark files.This uses
intand(float)0x01000000whileGpuResidentQuickHarness.csandTorchSharpCpuComparisonBenchmarks.csuseuintand16777216f. Both are functionally equivalent, but consider aligning for consistency.AiDotNetBenchmarkTests/Benchmarking/FixedProjectFileToolchain.cs (2)
47-54: Empty catch block silently swallows exceptions.The generic
catchwithout exception type or logging makes debugging difficult ifMainModuleaccess fails unexpectedly (e.g., due to permissions). Consider catching the specific exception types (InvalidOperationException,NotSupportedException) or at minimum, avoid swallowing silently.♻️ Suggested improvement
try { processPath = Process.GetCurrentProcess().MainModule?.FileName; } - catch + catch (InvalidOperationException) + { + // Process has exited or MainModule is unavailable + processPath = null; + } + catch (NotSupportedException) { + // Platform does not support MainModule access processPath = null; }
83-95: Reflection-based override is fragile; consider documenting the risk.Accessing
<Generator>k__BackingFieldrelies on compiler-generated backing field naming and BenchmarkDotNet internals. This could silently break on library upgrades. Consider adding a comment explaining why this approach is necessary and noting the BenchmarkDotNet version it was tested against.+ /// <summary> + /// Overrides the generator field via reflection. This is fragile and depends on + /// BenchmarkDotNet internals (tested with v0.15.8). If BenchmarkDotNet changes its + /// toolchain implementation, this may need adjustment. + /// </summary> private static void OverrideGenerator(IToolchain toolchain, IGenerator generator)src/AiDotNet.Native.OneDNN/AiDotNet.Native.OneDNN.csproj (1)
3-19: Centralize package version metadata.Hard‑coding
<Version>here can drift from repo‑wide versioning. Consider moving it toDirectory.Build.propsso all packages stay aligned. Based on learnings, prefer a single version source of truth.src/AiDotNet.Tensors/Engines/CpuNativeBlas.cs (2)
33-79: Guard unsafe GEMM calls with argument validation (or verify caller contracts).These overloads pin arrays and pass raw pointers + offsets into native BLAS. If any offset/leading-dimension is wrong, the native call can read/write past array bounds and crash the process. Please confirm callers fully validate sizes/offsets, or add a fast guard/Debug.Assert here (e.g.,
aOffset + (m-1)*lda + k <= a.Length, etc.) before invoking.Also applies to: 82-129
167-239: Avoid forcing BLAS thread count when no env var is set.
ReadThreadCountfalls back toEnvironment.ProcessorCountwhenAIDOTNET_CPU_BLAS_THREADSis unset (and no fallback BLAS envs exist), soTryConfigureThreadsalways overrides library defaults. That can cause oversubscription in parallel workloads. Consider returningnullunless an explicit env var is present so BLAS keeps its default unless the user opts in.♻️ Suggested change to keep BLAS defaults unless explicitly set
- return Environment.ProcessorCount; + return null; @@ - return Environment.ProcessorCount; + return null;scripts/compare-cpu-benchmarks.ps1 (2)
23-51: Handle BenchmarkDotNet CSV unit variants.
Convert-ToNanosecondsassumes the unit is embedded in the value. If BenchmarkDotNet is configured to print units in headers (e.g.,Mean [ns]), values become numeric and this throws. Consider detecting the unit from the header and allowing numeric-only values via a unit hint.♻️ Suggested update to support header units
-function Convert-ToNanoseconds { - param([string]$value) +function Convert-ToNanoseconds { + param( + [string]$value, + [string]$unitHint + ) ... - if ($clean -notmatch '^([0-9.,]+)\s*([\p{L}]+)$') { + if ($clean -notmatch '^([0-9.,]+)\s*([\p{L}]+)?$') { throw "Unsupported time format: $value" } ... - $unit = $matches[2].Replace($microSign, 'u').Replace($muSign, 'u') + $unit = $matches[2] + if (-not $unit) { $unit = $unitHint } + if (-not $unit) { throw "Unsupported time format: $value" } + $unit = $unit.Replace($microSign, 'u').Replace($muSign, 'u')function Compare-Benchmarks { @@ - $rows = Import-Csv $CsvPath + $rows = Import-Csv $CsvPath + if (-not $rows) { throw "CSV has no rows: $CsvPath" } + $meanColumn = ($rows[0].PSObject.Properties.Name | Where-Object { $_ -like "Mean*" } | Select-Object -First 1) + if (-not $meanColumn) { throw "Mean column not found in $CsvPath" } + $meanUnit = $null + if ($meanColumn -match '\[(?<unit>[^\]]+)\]') { $meanUnit = $matches.unit } @@ - $aiNs = Convert-ToNanoseconds $row.Mean - $compNs = Convert-ToNanoseconds $compRow.Mean + $aiNs = Convert-ToNanoseconds -value $row.PSObject.Properties[$meanColumn].Value -unitHint $meanUnit + $compNs = Convert-ToNanoseconds -value $compRow.PSObject.Properties[$meanColumn].Value -unitHint $meanUnit
63-127: Verify keying if benchmarks use multiple params.
Keys are built fromMethod|sizeonly. If any CPU comparison benchmarks add params beyondSize(e.g.,M/N/K), rows can collide and comparisons may mismatch. Please confirm all relevant suites only varySize, or make the key configurable to include additional param columns.
- Fix 4 generic catch clauses in BlasProvider.cs by changing to catch (Exception) with explanatory comments about MKL.NET failures - Fix uninitialized allocation in VectorBase.cs by making zero-init the default behavior, with skipZeroInit parameter for performance-critical code that immediately overwrites all elements - Fix cross-platform dotnet CLI path resolution in FixedProjectFileToolchain by using RuntimeInformation.IsOSPlatform to determine correct executable name (dotnet.exe on Windows, dotnet on Linux/macOS) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Add [MethodImpl(MethodImplOptions.NoInlining)] attribute to MKL.NET helper methods to prevent JIT from trying to load MKL.NET assembly types when the containing method is compiled. This fixes FileNotFoundException on platforms where MKL.NET is not available. Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Summary\n- add BLAS-backed and blocked matmul paths for CPU tensors\n- reuse GPU buffers via pooling and support GPU-resident benchmarking\n- expand CPU/GPU benchmark coverage (TorchSharp, TensorFlow, ML.NET)\n\n## Testing\n- dotnet build -c Release