sciBASIC#
↖
10 Tutorial — IL → CUDA
Write plain VB.NET math — the runtime decompiles it to CUDA C and launches it on the GPU.
No CUDA source, no nvcc project files. Six ordinary Shared functions — row sums, row square-sums, a Gram dot product, a Pearson-correlation cell, a Euclidean-distance cell and a scalar clamp — are reflected from the running script, lifted into an AST, re-emitted as .cu, compiled on the fly by NVRTC and launched against a 1024 × 1024 single-precision matrix on an NVIDIA RTX A4000. Every stage cross-checks itself: interpreted AST against the original method, GPU results against double-precision CPU references.
02 Pipeline
Four steps from IL bytecode to a GPU result
Step 1
Translate
IlCudaTranslator.Translate reflects each VB function, decompiles the IL into a
MethodSyntax AST and emits device + kernel .cu code; a CPU interpreter
re-evaluates the same AST as a self-check.
Step 2
Register
kernel.Register() hands every generated source to KernelSources —
before the engine is created.
Step 3
Launch
CudaEngine.TryCreate JIT-compiles via NVRTC (cubin, sm_86) and launches:
1-D kernels for row statistics, 16 × 16 2-D grids for the 1024 × 1024
Gram / correlation / distance matrices.
Step 4
Verify
GPU matrices are compared element-wise against double-precision CPU references — max
|error| 2.2e-7 (Pearson) and 2.1e-5 (distance). Every check prints
OK.
The pipeline maps VB functions onto kernels through three conventions, taken verbatim
from the script header: ① a parameter named
i → 1-D kernel,
i = blockIdx.x · blockDim.x + threadIdx.x; ② parameters i and
j → 2-D kernel, i on rows (blockIdx.y), j on
columns (blockIdx.x); ③ pure scalars with no array indexing → auto-wrapped as an
element-wise kernel over per-index arrays.01 The Script
Full demo source
The complete script exactly as executed by the sciBASIC# script engine (vbs.exe) — nothing elided.
cuda.vb · 581 linesDownload cuda.vb
#include "Microsoft.VisualBasic.Computing.ILCuda.dll"
#include "Microsoft.VisualBasic.ApplicationServices.Development.VisualStudio.dll"
imports System.Reflection
imports System.Text
imports System.Text.RegularExpressions
imports Microsoft.VisualBasic.Computing.ILCuda.Runtime
imports Microsoft.VisualBasic.Computing.ILCuda.IL2Cuda
imports Microsoft.VisualBasic.ApplicationServices.Development.VisualStudio.IL
' ============================================================================
' IL -> CUDA tutorial: define plain VB.NET math functions in this script and
' use them to compute a pairwise Pearson correlation matrix and a Euclidean
' distance matrix on the GPU.
'
' Run:
' vbs.exe tutorials\VBS\cuda.vb
'
' Pipeline:
' VB.NET function -> IL bytecode -> AST (MethodSyntax) -> .cu source
' -> NVRTC just-in-time compile -> cuLaunchKernel -> result matrix
'
' Three kernel-mapping conventions (see cuda\ILCuda\README.md for details):
' 1) a parameter named i -> 1D kernel, i = blockIdx.x * blockDim.x + threadIdx.x
' 2) parameters i and j -> 2D kernel, i indexes the row (blockIdx.y) and j the column (blockIdx.x)
' 3) pure scalars, no array element access -> automatically wrapped element-wise; every scalar
' parameter becomes an array that is read by the element index
'
' Note: because the VBS engine rewrites top-level Function blocks into anonymous
' functions, their Shared MethodInfo is not reachable, so every method to be
' decompiled is placed inside the Public Class below -- the type block is moved
' verbatim into the generated Module and can be reflected normally.
' ============================================================================
' ---------------------------------------------------------------------------
' Section 1: target functions to decompile
'
' Everything below is ordinary VB.NET and contains no CUDA concepts. IL2Cuda
' decompiles these methods into expression trees at runtime and then emits .cu
' source from them.
'
' For row i let sum_i = Σx, sq_i = Σx², dot_ij = Σx_ik·x_jk and n = column count:
' mean_i = sum_i / n
' var_i = sq_i - n * mean_i²
' corr_ij = (dot_ij - n * mean_i * mean_j) / sqrt(var_i * var_j)
' dist_ij = sqrt(sq_i + sq_j - 2 * dot_ij)
'
' On the diagonal corr is mathematically 1 and dist is 0; those cells are
' returned directly so that subtracting two close large numbers in single
' precision cannot amplify the rounding error (see the numeric notes in README).
' ---------------------------------------------------------------------------
Public Class PearsonMetrics
' Row sum; convention 1: the i parameter is the thread index -> 1D kernel
Public Shared Function RowSum(x As Single(), cols As Integer, i As Integer) As Single
Dim sum As Single = 0.0F
For k As Integer = 0 To cols - 1
sum += x(i * cols + k)
Next
Return sum
End Function
' Row sum of squares; also a 1D kernel
Public Shared Function RowSumSq(x As Single(), cols As Integer, i As Integer) As Single
Dim sumSq As Single = 0.0F
For k As Integer = 0 To cols - 1
Dim v As Single = x(i * cols + k)
sumSq += v * v
Next
Return sumSq
End Function
' Dot product of two rows; convention 2: both i and j are present -> 2D kernel
Public Shared Function GramDot(x As Single(), cols As Integer, i As Integer, j As Integer) As Single
Dim acc As Single = 0.0F
For k As Integer = 0 To cols - 1
acc += x(i * cols + k) * x(j * cols + k)
Next
Return acc
End Function
' Pearson correlation cell: guard clause + if/else diamond + MathF calls
Public Shared Function CorrelationCell(dot As Single(), rowSum As Single(), rowSumSq As Single(),
rows As Integer, cols As Integer,
i As Integer, j As Integer) As Single
If i = j Then
Return 1.0F
End If
Dim n As Single = CSng(cols)
Dim meanI As Single = rowSum(i) / n
Dim meanJ As Single = rowSum(j) / n
Dim varI As Single = rowSumSq(i) - n * meanI * meanI
Dim varJ As Single = rowSumSq(j) - n * meanJ * meanJ
Dim d As Single = dot(i * rows + j)
Dim cov As Single = d - n * meanI * meanJ
Dim denom As Single = MathF.Sqrt(MathF.Max(varI, 0.0F)) * MathF.Sqrt(MathF.Max(varJ, 0.0F))
Dim c As Single
If denom > 1.0E-12F Then
c = cov / denom
Else
c = 0.0F
End If
Return MathF.Min(1.0F, MathF.Max(-1.0F, c))
End Function
' Euclidean distance cell: guard clause + MathF.Sqrt
Public Shared Function DistanceCell(dot As Single(), rowSumSq As Single(),
rows As Integer, i As Integer, j As Integer) As Single
If i = j Then
Return 0.0F
End If
Dim d As Single = dot(i * rows + j)
Dim d2 As Single = MathF.Max(rowSumSq(i) + rowSumSq(j) - 2.0F * d, 0.0F)
Return MathF.Sqrt(d2)
End Function
' Pure scalar clamp; convention 3: no i/j and no array element access ->
' automatically wrapped element-wise
Public Shared Function PearsonClamp(cov As Single, denom As Single) As Single
Dim c As Single
If denom > 1.0E-12F Then
c = cov / denom
Else
c = 0.0F
End If
Return MathF.Min(1.0F, MathF.Max(-1.0F, c))
End Function
End Class
' ---------------------------------------------------------------------------
' Section 2: tutorial helpers (target method list / CPU reference / formatting / self-check)
' ---------------------------------------------------------------------------
Public Class TutorialKit
''' All target methods to be decompiled
Public Shared Function Targets() As MethodInfo()
Dim t As Type = GetType(PearsonMetrics)
Return New MethodInfo() {
t.GetMethod("RowSum"),
t.GetMethod("RowSumSq"),
t.GetMethod("GramDot"),
t.GetMethod("CorrelationCell"),
t.GetMethod("DistanceCell"),
t.GetMethod("PearsonClamp")
}
End Function
''' Build a reproducible random matrix (row-major)
Public Shared Function MakeMatrix(rows As Integer, cols As Integer, seed As Integer) As Single()
Dim rnd As New Random(seed)
Dim data(rows * cols - 1) As Single
For i As Integer = 0 To data.Length - 1
data(i) = CSng(rnd.NextDouble() * 2.0 - 1.0)
Next
Return data
End Function
''' Sample arguments for the interpreter self-check: valid indices, non-degenerate values
Public Shared Function SampleArgs(m As MethodInfo) As Object()
Dim ps As ParameterInfo() = m.GetParameters()
Dim args(ps.Length - 1) As Object
For i As Integer = 0 To ps.Length - 1
Dim p As ParameterInfo = ps(i)
Dim lower As String = p.Name.ToLowerInvariant()
If p.ParameterType.IsArray Then
Dim data(255) As Single
For k As Integer = 0 To 255
data(k) = CSng((k Mod 17) - 8) * 0.5F
Next
args(i) = data
ElseIf p.ParameterType = GetType(Single) Then
args(i) = 1.5F
ElseIf p.ParameterType = GetType(Integer) Then
Select Case lower
Case "cols" : args(i) = 8
Case "rows" : args(i) = 16
Case "i" : args(i) = 2
Case "j" : args(i) = 5
Case Else : args(i) = 3
End Select
End If
Next
Return args
End Function
''' CPU reference implementation: Pearson correlation matrix (double-precision two-pass)
Public Shared Function CpuCorrelation(data As Single(), rows As Integer, cols As Integer) As Double()
Dim n As Double = cols
Dim mean(rows - 1) As Double
Dim sd(rows - 1) As Double
For i As Integer = 0 To rows - 1
Dim s As Double = 0
Dim sq As Double = 0
For k As Integer = 0 To cols - 1
Dim v As Double = data(i * cols + k)
s += v
sq += v * v
Next
mean(i) = s / n
sd(i) = System.Math.Sqrt(System.Math.Max(sq - n * mean(i) * mean(i), 0.0))
Next
Dim out(rows * rows - 1) As Double
For i As Integer = 0 To rows - 1
For j As Integer = 0 To rows - 1
Dim d As Double = 0
For k As Integer = 0 To cols - 1
d += data(i * cols + k) * data(j * cols + k)
Next
Dim cov As Double = d - n * mean(i) * mean(j)
Dim den As Double = sd(i) * sd(j)
Dim c As Double = If(den > 0.0, cov / den, 0.0)
out(i * rows + j) = If(i = j, 1.0, System.Math.Min(1.0, System.Math.Max(-1.0, c)))
Next
Next
Return out
End Function
''' CPU reference implementation: Euclidean distance matrix
Public Shared Function CpuDistance(data As Single(), rows As Integer, cols As Integer) As Double()
Dim sq(rows - 1) As Double
For i As Integer = 0 To rows - 1
Dim s As Double = 0
For k As Integer = 0 To cols - 1
Dim v As Double = data(i * cols + k)
s += v * v
Next
sq(i) = s
Next
Dim out(rows * rows - 1) As Double
For i As Integer = 0 To rows - 1
For j As Integer = 0 To rows - 1
Dim d As Double = 0
For k As Integer = 0 To cols - 1
d += data(i * cols + k) * data(j * cols + k)
Next
out(i * rows + j) = If(i = j, 0.0,
System.Math.Sqrt(System.Math.Max(sq(i) + sq(j) - 2.0 * d, 0.0)))
Next
Next
Return out
End Function
''' Maximum absolute error between the GPU result and the CPU reference
Public Shared Function MaxError(got As Single(), expect As Double()) As Double
Dim maxDiff As Double = 0
For i As Integer = 0 To got.Length - 1
Dim diff As Double = System.Math.Abs(CDbl(got(i)) - expect(i))
If diff > maxDiff Then
maxDiff = diff
End If
Next
Return maxDiff
End Function
''' Maximum absolute error between the GPU result and a single-precision reference
Public Shared Function MaxErrorSingle(got As Single(), expect As Single()) As Double
Dim maxDiff As Double = 0
For i As Integer = 0 To got.Length - 1
Dim diff As Double = System.Math.Abs(CDbl(got(i)) - CDbl(expect(i)))
If diff > maxDiff Then
maxDiff = diff
End If
Next
Return maxDiff
End Function
''' Convert double-precision references to single precision so the same preview layout can be reused
Public Shared Function ToSingle(source As Double()) As Single()
Dim out(source.Length - 1) As Single
For i As Integer = 0 To source.Length - 1
out(i) = CSng(source(i))
Next
Return out
End Function
''' Section title
Public Shared Sub PrintTitle(text As String)
Call Console.WriteLine()
Call Console.WriteLine(New String("="c, 74))
Call Console.WriteLine(" " & text)
Call Console.WriteLine(New String("="c, 74))
End Sub
''' Print a multi-line text block with indentation (pseudo code / CUDA source)
Public Shared Sub PrintBlock(text As String, indent As String)
Dim lines As String() = Regex.Split(text, "\r\n|\r|\n")
For Each line As String In lines
Call Console.WriteLine(indent & line)
Next
End Sub
''' Print the top-left n x n block of a result matrix
Public Shared Sub PrintMatrix(title As String, m As Single(), rows As Integer, n As Integer)
Call Console.WriteLine(" " & title)
For i As Integer = 0 To n - 1
Dim line As New StringBuilder()
For j As Integer = 0 To n - 1
Call line.Append(m(i * rows + j).ToString("F4").PadLeft(10))
Next
Call Console.WriteLine(" " & line.ToString())
Next
End Sub
End Class
' ---------------------------------------------------------------------------
' Section 3: main flow
' ---------------------------------------------------------------------------
dim rows As Integer = 1024
dim cols As Integer = 1024
dim seed As Integer = 42
dim preview As Integer = 6
dim allOk As Boolean = True
call TutorialKit.PrintTitle("IL -> CUDA tutorial: Pearson correlation matrix and Euclidean distance matrix")
call console.WriteLine(" input size: " & rows & " rows x " & cols & " cols, random seed " & seed)
' ---------- Build the input matrix ----------
dim data As Single() = TutorialKit.MakeMatrix(rows, cols, seed)
' ---------- Step 1: IL -> AST -> .cu ----------
dim kernels As New Dictionary(Of String, IlCudaKernel)(StringComparer.Ordinal)
call TutorialKit.PrintTitle("Step 1: decompile the VB.NET functions into an AST and emit .cu source")
for each m As MethodInfo In TutorialKit.Targets()
dim errMsg As String = Nothing
dim kernel As IlCudaKernel = Nothing
try
kernel = IlCudaTranslator.Translate(m)
catch ex As Exception
errMsg = ex.GetBaseException().Message
end try
if errMsg IsNot Nothing Then
call console.WriteLine(" " & m.Name & " decompile failed: " & errMsg)
allOk = False
else
call kernels.Add(m.Name, kernel)
call console.WriteLine()
call console.WriteLine(" ---- " & m.Name & " ----")
call console.WriteLine(" index mode : " & kernel.IndexMode.ToString())
call console.WriteLine(" device func: " & kernel.DeviceFunctionName)
call console.WriteLine(" kernel : " & kernel.Name)
call console.WriteLine()
call console.WriteLine(" reconstructed pseudo code:")
call TutorialKit.PrintBlock(SyntaxWriter.WriteMethod(kernel.Syntax), " ")
call console.WriteLine()
call console.WriteLine(" generated CUDA source:")
call TutorialKit.PrintBlock(kernel.Source, " ")
' Interpreter self-check: run the same AST through the CPU interpreter and
' compare the result with a direct invocation of the original method
dim sampleArgs As Object() = TutorialKit.SampleArgs(m)
dim expected As Object = m.Invoke(Nothing, sampleArgs)
dim actual As Object = New AstInterpreter(kernel.Syntax).Invoke(sampleArgs)
dim diff As Double = System.Math.Abs(System.Convert.ToDouble(expected) - System.Convert.ToDouble(actual))
dim tol As Double = 1.0E-5 * System.Math.Max(1.0, System.Math.Abs(System.Convert.ToDouble(expected)))
dim pass As Boolean = diff <= tol
call console.WriteLine()
call console.WriteLine(" interpreter self-check: original=" & expected & ", AST=" & actual &
", diff=" & diff.ToString("E3") & " " & (if(pass, "OK", "FAIL")))
if Not pass Then
allOk = False
end if
end if
next
' ---------- Step 2: register the generated .cu into KernelSources ----------
call TutorialKit.PrintTitle("Step 2: register the kernel sources (must be done before creating the engine)")
for each k As IlCudaKernel In kernels.Values
call k.Register()
call console.WriteLine(" registered " & k.ToString())
next
' ---------- Step 3: NVRTC compile + launch the kernels ----------
call TutorialKit.PrintTitle("Step 3: compute the correlation and distance matrices on the GPU")
dim opts As New EngineOptions With {.DeviceOrdinal = 0}
dim engine As CudaEngine = CudaEngine.TryCreate(opts)
dim corrGPU As Single() = Nothing
dim distGPU As Single() = Nothing
dim gpuOk As Boolean = False
if Not allOk Then
call console.WriteLine(" Step 1 had failures, skipping the GPU computation.")
elseif engine Is Nothing Then
call console.WriteLine(" GPU unavailable: " & opts.ErrorMessage)
dim report = CudaEnvironment.Probe()
for each s As FixSuggestion In CudaEnvironment.Suggest(report)
call console.WriteLine(" suggestion: " & s.ToString())
call console.WriteLine(" " & s.Detail)
next
call console.WriteLine()
call console.WriteLine(" >> fell back to the CPU reference implementation; only CPU results are shown below.")
else
using engine
dim cells As Integer = rows * rows
call console.WriteLine(" device : " & engine.Device.Name)
call console.WriteLine(" kernel image: " & engine.Image.ToString())
using bufX As New DeviceBuffer(Of Single)(data.Length), _
bufSum As New DeviceBuffer(Of Single)(rows), _
bufSumSq As New DeviceBuffer(Of Single)(rows), _
bufDot As New DeviceBuffer(Of Single)(cells), _
bufCorr As New DeviceBuffer(Of Single)(cells), _
bufDist As New DeviceBuffer(Of Single)(cells)
call bufX.Write(data)
' 1) Row statistics: 1D kernel, one thread per row
call kernels("RowSum").Launch(engine, LaunchPlanner.For1D(rows, 256), bufX, cols, bufSum, rows)
call kernels("RowSumSq").Launch(engine, LaunchPlanner.For1D(rows, 256), bufX, cols, bufSumSq, rows)
' 2) Gram dot products: 2D kernel
call kernels("GramDot").Launch(engine, LaunchPlanner.For2D(rows, rows, 16, 16), bufX, cols, bufDot, rows, rows)
' 3) Reconstruct the correlation and distance matrices from the dot products
' and the row statistics
call kernels("CorrelationCell").Launch(engine, LaunchPlanner.For2D(rows, rows, 16, 16),
bufDot, bufSum, bufSumSq, rows, cols, bufCorr, rows, rows)
call kernels("DistanceCell").Launch(engine, LaunchPlanner.For2D(rows, rows, 16, 16),
bufDot, bufSumSq, rows, bufDist, rows, rows)
call engine.Synchronize()
corrGPU = bufCorr.Read()
distGPU = bufDist.Read()
gpuOk = True
' 4) Convention 3 demo: a pure scalar function is auto-wrapped into an
' element-wise kernel
dim hostDot As Single() = bufDot.Read()
dim hostSum As Single() = bufSum.Read()
dim hostSumSq As Single() = bufSumSq.Read()
dim n1 As Single = CSng(cols)
dim cov(cells - 1) As Single
dim denom(cells - 1) As Single
dim clampRef(cells - 1) As Single
for i As Integer = 0 To rows - 1
dim meanI As Single = hostSum(i) / n1
dim varI As Single = hostSumSq(i) - n1 * meanI * meanI
dim sdI As Single = MathF.Sqrt(MathF.Max(varI, 0.0F))
for j As Integer = 0 To rows - 1
dim meanJ As Single = hostSum(j) / n1
dim varJ As Single = hostSumSq(j) - n1 * meanJ * meanJ
dim idx As Integer = i * rows + j
cov(idx) = hostDot(idx) - n1 * meanI * meanJ
denom(idx) = sdI * MathF.Sqrt(MathF.Max(varJ, 0.0F))
clampRef(idx) = PearsonMetrics.PearsonClamp(cov(idx), denom(idx))
next
next
using bufCov As New DeviceBuffer(Of Single)(cells), _
bufDenom As New DeviceBuffer(Of Single)(cells), _
bufClamp As New DeviceBuffer(Of Single)(cells)
call bufCov.Write(cov)
call bufDenom.Write(denom)
call kernels("PearsonClamp").Launch(engine, LaunchPlanner.For1D(cells, 256),
bufCov, bufDenom, bufClamp, cells)
call engine.Synchronize()
dim clampErr As Double = TutorialKit.MaxErrorSingle(bufClamp.Read(), clampRef)
dim clampPass As Boolean = clampErr <= 1.0E-5
call console.WriteLine()
call console.WriteLine(" PearsonClamp (auto-wrapped element-wise) max abs error = " &
clampErr.ToString("E3") & " " & (if(clampPass, "OK", "FAIL")))
if Not clampPass Then
allOk = False
end if
end using
end using
end using
end if
' ---------- Step 4: compare against the CPU reference implementation ----------
call TutorialKit.PrintTitle("Step 4: compare with the CPU reference implementation + preview the results")
dim refCorr As Double() = TutorialKit.CpuCorrelation(data, rows, cols)
dim refDist As Double() = TutorialKit.CpuDistance(data, rows, cols)
if gpuOk Then
dim eCorr As Double = TutorialKit.MaxError(corrGPU, refCorr)
dim eDist As Double = TutorialKit.MaxError(distGPU, refDist)
dim corrPass As Boolean = eCorr <= 1.0E-3
dim distPass As Boolean = eDist <= 1.0E-2
call console.WriteLine(" Pearson correlation matrix max abs error = " & eCorr.ToString("E3") & " " & (if(corrPass, "OK", "FAIL")))
call console.WriteLine(" Euclidean distance matrix max abs error = " & eDist.ToString("E3") & " " & (if(distPass, "OK", "FAIL")))
if (Not corrPass) OrElse (Not distPass) Then
allOk = False
end if
call console.WriteLine()
call TutorialKit.PrintMatrix("Pearson correlation matrix (GPU, top-left " & preview & " x " & preview & "):", corrGPU, rows, preview)
call console.WriteLine()
call TutorialKit.PrintMatrix("Euclidean distance matrix (GPU, top-left " & preview & " x " & preview & "):", distGPU, rows, preview)
else
call console.WriteLine(" GPU results unavailable; the CPU reference results are shown below.")
call console.WriteLine()
call TutorialKit.PrintMatrix("Pearson correlation matrix (CPU, top-left " & preview & " x " & preview & "):", TutorialKit.ToSingle(refCorr), rows, preview)
call console.WriteLine()
call TutorialKit.PrintMatrix("Euclidean distance matrix (CPU, top-left " & preview & " x " & preview & "):", TutorialKit.ToSingle(refDist), rows, preview)
end if
call console.WriteLine()
if allOk Then
call console.WriteLine("All tutorial steps passed.")
else
call console.WriteLine("Some steps failed, please check the output above.")
end if
03 Results
Console output — stdout.txt
The full command-line log of the run on an NVIDIA RTX A4000 (NVRTC 13.3, compute capability
sm_86): decompiled pseudocode, the emitted CUDA source for each of the six kernels,
interpreter self-checks, kernel registration, device info and the final CPU-vs-GPU error report with a
6 × 6 preview of both result matrices. The complete log scrolls below.
vbs.exe cuda.vb · stdout
Download stdout.txt
==========================================================================
IL -> CUDA tutorial: Pearson correlation matrix and Euclidean distance matrix
==========================================================================
input size: 1024 rows x 1024 cols, random seed 42
==========================================================================
Step 1: decompile the VB.NET functions into an AST and emit .cu source
==========================================================================
---- RowSum ----
index mode : Grid1D
device func: il_RowSum_scalar
kernel : il_RowSum_kernel
reconstructed pseudo code:
Shared Function RowSum(p_x As Single(), p_cols As Integer, p_i As Integer) As Single
Dim V_0 As Single = 0F
Dim V_1 As Integer = 0
Dim V_2 As Integer = 0
Dim V_0_1 As Single
Dim V_2_1 As Integer
Dim V_0_2 As Single = 0F
Dim V_1_1 As Integer = (p_cols - 1)
Dim V_2_2 As Integer = 0
V_0_1 = V_0_2
For V_2_1 = V_2_2 While (V_2_1 <= V_1_1) Step ((V_2_1 + 1) - V_2_1)
Dim V_0_3 As Single = (V_0_1 + p_x[((p_i * p_cols) + V_2_1)])
V_0_1 = V_0_3
Next
Return V_0_1
End Function
generated CUDA source:
// ===== 本文件由 IL -> AST -> CUDA 流水线自动生成，请勿手工编辑 =====
// 来源方法 : RowSum
// 设备函数 : il_RowSum_scalar
// 内核 : il_RowSum_kernel（索引模式 Grid1D）
__device__ float il_RowSum_scalar(const float* __restrict__ p_x, int p_cols, int p_i) {
float V_0 = 0.0f;
int V_1 = 0;
int V_2 = 0;
float V_0_1;
int V_2_1;
float V_0_2 = 0.0f;
int V_1_1 = p_cols - 1;
int V_2_2 = 0;
V_0_1 = V_0_2;
for (V_2_1 = V_2_2; V_2_1 <= V_1_1; V_2_1 = V_2_1 + 1) {
float V_0_3 = V_0_1 + p_x[(p_i * p_cols) + V_2_1];
V_0_1 = V_0_3;
}
return V_0_1;
}
extern "C" __global__ void il_RowSum_kernel(const float* __restrict__ p_x, int p_cols, float* __restrict__ il_out, int il_n) {
int p_i = blockIdx.x * blockDim.x + threadIdx.x;
if (p_i >= il_n) return;
il_out[p_i] = il_RowSum_scalar(p_x, p_cols, p_i);
}
interpreter self-check: original=-13.5, AST=-13.5, diff=0.000E+000 OK
---- RowSumSq ----
index mode : Grid1D
device func: il_RowSumSq_scalar
kernel : il_RowSumSq_kernel
reconstructed pseudo code:
Shared Function RowSumSq(p_x As Single(), p_cols As Integer, p_i As Integer) As Single
Dim V_0 As Single = 0F
Dim V_1 As Integer = 0
Dim V_2 As Integer = 0
Dim V_3 As Single = 0F
Dim V_0_1 As Single
Dim V_2_1 As Integer
Dim V_3_1 As Single
Dim V_0_2 As Single = 0F
Dim V_1_1 As Integer = (p_cols - 1)
Dim V_2_2 As Integer = 0
V_0_1 = V_0_2
V_3_1 = V_3
For V_2_1 = V_2_2 While (V_2_1 <= V_1_1) Step ((V_2_1 + 1) - V_2_1)
Dim V_3_2 As Single = p_x[((p_i * p_cols) + V_2_1)]
Dim V_0_3 As Single = (V_0_1 + (V_3_2 * V_3_2))
V_0_1 = V_0_3
V_3_1 = V_3_2
Next
Return V_0_1
End Function
generated CUDA source:
// ===== 本文件由 IL -> AST -> CUDA 流水线自动生成，请勿手工编辑 =====
// 来源方法 : RowSumSq
// 设备函数 : il_RowSumSq_scalar
// 内核 : il_RowSumSq_kernel（索引模式 Grid1D）
__device__ float il_RowSumSq_scalar(const float* __restrict__ p_x, int p_cols, int p_i) {
float V_0 = 0.0f;
int V_1 = 0;
int V_2 = 0;
float V_3 = 0.0f;
float V_0_1;
int V_2_1;
float V_3_1;
float V_0_2 = 0.0f;
int V_1_1 = p_cols - 1;
int V_2_2 = 0;
V_0_1 = V_0_2;
V_3_1 = V_3;
for (V_2_1 = V_2_2; V_2_1 <= V_1_1; V_2_1 = V_2_1 + 1) {
float V_3_2 = p_x[(p_i * p_cols) + V_2_1];
float V_0_3 = V_0_1 + (V_3_2 * V_3_2);
V_0_1 = V_0_3;
V_3_1 = V_3_2;
}
return V_0_1;
}
extern "C" __global__ void il_RowSumSq_kernel(const float* __restrict__ p_x, int p_cols, float* __restrict__ il_out, int il_n) {
int p_i = blockIdx.x * blockDim.x + threadIdx.x;
if (p_i >= il_n) return;
il_out[p_i] = il_RowSumSq_scalar(p_x, p_cols, p_i);
}
interpreter self-check: original=66.75, AST=66.75, diff=0.000E+000 OK
---- GramDot ----
index mode : Grid2D
device func: il_GramDot_scalar
kernel : il_GramDot_kernel
reconstructed pseudo code:
Shared Function GramDot(p_x As Single(), p_cols As Integer, p_i As Integer, p_j As Integer) As Single
Dim V_0 As Single = 0F
Dim V_1 As Integer = 0
Dim V_2 As Integer = 0
Dim V_0_1 As Single
Dim V_2_1 As Integer
Dim V_0_2 As Single = 0F
Dim V_1_1 As Integer = (p_cols - 1)
Dim V_2_2 As Integer = 0
V_0_1 = V_0_2
For V_2_1 = V_2_2 While (V_2_1 <= V_1_1) Step ((V_2_1 + 1) - V_2_1)
Dim V_0_3 As Single = (V_0_1 + (p_x[((p_i * p_cols) + V_2_1)] * p_x[((p_j * p_cols) + V_2_1)]))
V_0_1 = V_0_3
Next
Return V_0_1
End Function
generated CUDA source:
// ===== 本文件由 IL -> AST -> CUDA 流水线自动生成，请勿手工编辑 =====
// 来源方法 : GramDot
// 设备函数 : il_GramDot_scalar
// 内核 : il_GramDot_kernel（索引模式 Grid2D）
__device__ float il_GramDot_scalar(const float* __restrict__ p_x, int p_cols, int p_i, int p_j) {
float V_0 = 0.0f;
int V_1 = 0;
int V_2 = 0;
float V_0_1;
int V_2_1;
float V_0_2 = 0.0f;
int V_1_1 = p_cols - 1;
int V_2_2 = 0;
V_0_1 = V_0_2;
for (V_2_1 = V_2_2; V_2_1 <= V_1_1; V_2_1 = V_2_1 + 1) {
float V_0_3 = V_0_1 + (p_x[(p_i * p_cols) + V_2_1] * p_x[(p_j * p_cols) + V_2_1]);
V_0_1 = V_0_3;
}
return V_0_1;
}
extern "C" __global__ void il_GramDot_kernel(const float* __restrict__ p_x, int p_cols, float* __restrict__ il_out, int il_nRows, int il_nCols) {
int p_i = blockIdx.y * blockDim.y + threadIdx.y;
int p_j = blockIdx.x * blockDim.x + threadIdx.x;
if (p_i >= il_nRows || p_j >= il_nCols) return;
il_out[p_i * il_nCols + p_j] = il_GramDot_scalar(p_x, p_cols, p_i, p_j);
}
interpreter self-check: original=-14.5, AST=-14.5, diff=0.000E+000 OK
---- CorrelationCell ----
index mode : Grid2D
device func: il_CorrelationCell_scalar
kernel : il_CorrelationCell_kernel
reconstructed pseudo code:
Shared Function CorrelationCell(p_dot As Single(), p_rowSum As Single(), p_rowSumSq As Single(), p_rows As Integer, p_cols As Integer, p_i As Integer, p_j As Integer) As Single
Dim V_0 As Single = 0F
Dim V_1 As Single = 0F
Dim V_2 As Single = 0F
Dim V_3 As Single = 0F
Dim V_4 As Single = 0F
Dim V_5 As Single = 0F
Dim V_6 As Single = 0F
Dim V_7 As Single = 0F
Dim V_0_1 As Single
Dim V_1_1 As Single
Dim V_2_1 As Single
Dim V_3_1 As Single
Dim V_4_1 As Single
Dim V_5_1 As Single
Dim V_6_1 As Single
Dim V_7_2 As Single
Dim V_7_1 As Single
If (p_i <> p_j) Then
Dim V_1_2 As Single = CType(p_cols, Single)
Dim V_2_2 As Single = (p_rowSum[p_i] / V_1_2)
Dim V_3_2 As Single = (p_rowSum[p_j] / V_1_2)
Dim V_4_2 As Single = (p_rowSumSq[p_j] - ((V_1_2 * V_3_2) * V_3_2))
Dim V_5_2 As Single = (p_dot[((p_i * p_rows) + p_j)] - ((V_1_2 * V_2_2) * V_3_2))
Dim V_6_2 As Single = (System.MathF.Sqrt(System.MathF.Max((p_rowSumSq[p_i] - ((V_1_2 * V_2_2) * V_2_2)), 0F)) * System.MathF.Sqrt(System.MathF.Max(V_4_2, 0F)))
If (V_6_2 <= 1E-12F) Then
Dim V_7_4 As Single = 0F
V_7_1 = V_7_4
Else
Dim V_7_3 As Single = (V_5_2 / V_6_2)
V_7_1 = V_7_3
End If
Dim V_0_3 As Single = System.MathF.Min(1F, System.MathF.Max(-1F, V_7_1))
V_0_1 = V_0_3
V_1_1 = V_1_2
V_2_1 = V_2_2
V_3_1 = V_3_2
V_4_1 = V_4_2
V_5_1 = V_5_2
V_6_1 = V_6_2
V_7_2 = V_7_1
Else
Dim V_0_2 As Single = 1F
V_0_1 = V_0_2
V_1_1 = V_1
V_2_1 = V_2
V_3_1 = V_3
V_4_1 = V_4
V_5_1 = V_5
V_6_1 = V_6
V_7_2 = V_7
End If
Return V_0_1
End Function
generated CUDA source:
// ===== 本文件由 IL -> AST -> CUDA 流水线自动生成，请勿手工编辑 =====
// 来源方法 : CorrelationCell
// 设备函数 : il_CorrelationCell_scalar
// 内核 : il_CorrelationCell_kernel（索引模式 Grid2D）
__device__ float il_CorrelationCell_scalar(const float* __restrict__ p_dot, const float* __restrict__ p_rowSum, const float* __restrict__ p_rowSumSq, int p_rows, int p_cols, int p_i, int p_j) {
float V_0 = 0.0f;
float V_1 = 0.0f;
float V_2 = 0.0f;
float V_3 = 0.0f;
float V_4 = 0.0f;
float V_5 = 0.0f;
float V_6 = 0.0f;
float V_7 = 0.0f;
float V_0_1;
float V_1_1;
float V_2_1;
float V_3_1;
float V_4_1;
float V_5_1;
float V_6_1;
float V_7_2;
float V_7_1;
if (p_i != p_j) {
float V_1_2 = ((float)p_cols);
float V_2_2 = p_rowSum[p_i] / V_1_2;
float V_3_2 = p_rowSum[p_j] / V_1_2;
float V_4_2 = p_rowSumSq[p_j] - ((V_1_2 * V_3_2) * V_3_2);
float V_5_2 = p_dot[(p_i * p_rows) + p_j] - ((V_1_2 * V_2_2) * V_3_2);
float V_6_2 = sqrtf(fmaxf(p_rowSumSq[p_i] - ((V_1_2 * V_2_2) * V_2_2), 0.0f)) * sqrtf(fmaxf(V_4_2, 0.0f));
if (V_6_2 <= 1E-12f) {
float V_7_4 = 0.0f;
V_7_1 = V_7_4;
} else {
float V_7_3 = V_5_2 / V_6_2;
V_7_1 = V_7_3;
}
float V_0_3 = fminf(1.0f, fmaxf(-1.0f, V_7_1));
V_0_1 = V_0_3;
V_1_1 = V_1_2;
V_2_1 = V_2_2;
V_3_1 = V_3_2;
V_4_1 = V_4_2;
V_5_1 = V_5_2;
V_6_1 = V_6_2;
V_7_2 = V_7_1;
} else {
float V_0_2 = 1.0f;
V_0_1 = V_0_2;
V_1_1 = V_1;
V_2_1 = V_2;
V_3_1 = V_3;
V_4_1 = V_4;
V_5_1 = V_5;
V_6_1 = V_6;
V_7_2 = V_7;
}
return V_0_1;
}
extern "C" __global__ void il_CorrelationCell_kernel(const float* __restrict__ p_dot, const float* __restrict__ p_rowSum, const float* __restrict__ p_rowSumSq, int p_rows, int p_cols, float* __restrict__ il_out, int il_nRows, int il_nCols) {
int p_i = blockIdx.y * blockDim.y + threadIdx.y;
int p_j = blockIdx.x * blockDim.x + threadIdx.x;
if (p_i >= il_nRows || p_j >= il_nCols) return;
il_out[p_i * il_nCols + p_j] = il_CorrelationCell_scalar(p_dot, p_rowSum, p_rowSumSq, p_rows, p_cols, p_i, p_j);
}
interpreter self-check: original=0, AST=0, diff=0.000E+000 OK
---- DistanceCell ----
index mode : Grid2D
device func: il_DistanceCell_scalar
kernel : il_DistanceCell_kernel
reconstructed pseudo code:
Shared Function DistanceCell(p_dot As Single(), p_rowSumSq As Single(), p_rows As Integer, p_i As Integer, p_j As Integer) As Single
Dim V_0 As Single = 0F
Dim V_1 As Single = 0F
Dim V_0_1 As Single
Dim V_1_1 As Single
If (p_i <> p_j) Then
Dim V_1_2 As Single = p_dot[((p_i * p_rows) + p_j)]
Dim V_0_3 As Single = System.MathF.Sqrt(System.MathF.Max(((p_rowSumSq[p_i] + p_rowSumSq[p_j]) - (2F * V_1_2)), 0F))
V_0_1 = V_0_3
V_1_1 = V_1_2
Else
Dim V_0_2 As Single = 0F
V_0_1 = V_0_2
V_1_1 = V_1
End If
Return V_0_1
End Function
generated CUDA source:
// ===== 本文件由 IL -> AST -> CUDA 流水线自动生成，请勿手工编辑 =====
// 来源方法 : DistanceCell
// 设备函数 : il_DistanceCell_scalar
// 内核 : il_DistanceCell_kernel（索引模式 Grid2D）
__device__ float il_DistanceCell_scalar(const float* __restrict__ p_dot, const float* __restrict__ p_rowSumSq, int p_rows, int p_i, int p_j) {
float V_0 = 0.0f;
float V_1 = 0.0f;
float V_0_1;
float V_1_1;
if (p_i != p_j) {
float V_1_2 = p_dot[(p_i * p_rows) + p_j];
float V_0_3 = sqrtf(fmaxf((p_rowSumSq[p_i] + p_rowSumSq[p_j]) - (2.0f * V_1_2), 0.0f));
V_0_1 = V_0_3;
V_1_1 = V_1_2;
} else {
float V_0_2 = 0.0f;
V_0_1 = V_0_2;
V_1_1 = V_1;
}
return V_0_1;
}
extern "C" __global__ void il_DistanceCell_kernel(const float* __restrict__ p_dot, const float* __restrict__ p_rowSumSq, int p_rows, float* __restrict__ il_out, int il_nRows, int il_nCols) {
int p_i = blockIdx.y * blockDim.y + threadIdx.y;
int p_j = blockIdx.x * blockDim.x + threadIdx.x;
if (p_i >= il_nRows || p_j >= il_nCols) return;
il_out[p_i * il_nCols + p_j] = il_DistanceCell_scalar(p_dot, p_rowSumSq, p_rows, p_i, p_j);
}
interpreter self-check: original=0.70710677, AST=0.70710677, diff=0.000E+000 OK
---- PearsonClamp ----
index mode : Grid1D
device func: il_PearsonClamp_scalar
kernel : il_PearsonClamp_kernel
reconstructed pseudo code:
Shared Function PearsonClamp(p_cov As Single, p_denom As Single) As Single
Dim V_0 As Single = 0F
Dim V_0_1 As Single
If (p_denom <= 1E-12F) Then
Dim V_0_3 As Single = 0F
V_0_1 = V_0_3
Else
Dim V_0_2 As Single = (p_cov / p_denom)
V_0_1 = V_0_2
End If
Return System.MathF.Min(1F, System.MathF.Max(-1F, V_0_1))
End Function
generated CUDA source:
// ===== 本文件由 IL -> AST -> CUDA 流水线自动生成，请勿手工编辑 =====
// 来源方法 : PearsonClamp
// 设备函数 : il_PearsonClamp_scalar
// 内核 : il_PearsonClamp_kernel（索引模式 Grid1D）
__device__ float il_PearsonClamp_scalar(float p_cov, float p_denom) {
float V_0 = 0.0f;
float V_0_1;
if (p_denom <= 1E-12f) {
float V_0_3 = 0.0f;
V_0_1 = V_0_3;
} else {
float V_0_2 = p_cov / p_denom;
V_0_1 = V_0_2;
}
return fminf(1.0f, fmaxf(-1.0f, V_0_1));
}
extern "C" __global__ void il_PearsonClamp_kernel(const float* __restrict__ p_cov, const float* __restrict__ p_denom, float* __restrict__ il_out, int il_n) {
int il_i = blockIdx.x * blockDim.x + threadIdx.x;
if (il_i >= il_n) return;
il_out[il_i] = il_PearsonClamp_scalar(p_cov[il_i], p_denom[il_i]);
}
interpreter self-check: original=1, AST=1, diff=0.000E+000 OK
==========================================================================
Step 2: register the kernel sources (must be done before creating the engine)
==========================================================================
registered il_RowSum_kernel [Grid1D] <- PearsonMetrics.RowSum
registered il_RowSumSq_kernel [Grid1D] <- PearsonMetrics.RowSumSq
registered il_GramDot_kernel [Grid2D] <- PearsonMetrics.GramDot
registered il_CorrelationCell_kernel [Grid2D] <- PearsonMetrics.CorrelationCell
registered il_DistanceCell_kernel [Grid2D] <- PearsonMetrics.DistanceCell
registered il_PearsonClamp_kernel [Grid1D] <- PearsonMetrics.PearsonClamp
==========================================================================
Step 3: compute the correlation and distance matrices on the GPU
==========================================================================
device : NVIDIA RTX A4000
kernel image: NVRTC 13.3 (cubin) arch=sm_86 [C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.3\bin\x64\nvrtc64_130_0.dll]
PearsonClamp (auto-wrapped element-wise) max abs error = 0.000E+000 OK
==========================================================================
Step 4: compare with the CPU reference implementation + preview the results
==========================================================================
Pearson correlation matrix max abs error = 2.221E-007 OK
Euclidean distance matrix max abs error = 2.069E-005 OK
Pearson correlation matrix (GPU, top-left 6 x 6):
1.0000 -0.0151 -0.0096 -0.0299 -0.0083 0.0006
-0.0151 1.0000 -0.0564 -0.0311 -0.0131 -0.0156
-0.0096 -0.0564 1.0000 0.0206 -0.0392 -0.0140
-0.0299 -0.0311 0.0206 1.0000 0.0250 -0.0570
-0.0083 -0.0131 -0.0392 0.0250 1.0000 -0.0185
0.0006 -0.0156 -0.0140 -0.0570 -0.0185 1.0000
Euclidean distance matrix (GPU, top-left 6 x 6):
0.0000 26.8603 26.5958 26.5702 26.4389 26.4184
26.8603 0.0000 27.1844 26.5662 26.4858 26.5902
26.5958 27.1844 0.0000 25.6775 26.5984 26.3983
26.5702 26.5662 25.6775 0.0000 25.4752 26.6538
26.4389 26.4858 26.5984 25.4752 0.0000 26.3193
26.4184 26.5902 26.3983 26.6538 26.3193 0.0000
All tutorial steps passed.Every verification line ends in
OK: the AST interpreter reproduces the
original method results bit-for-bit on all six kernels, the element-wise PearsonClamp wrapper
matches its host-side reference at 0.000E+000, and both GPU matrices agree with the
double-precision CPU baseline to single-precision rounding.