sciBASIC#
↖
02 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 · 572 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 教程: 用脚本里定义的 VB.NET 数学函数, 在 GPU 上算矩阵行间的
' 皮尔逊相关系数矩阵 与 欧氏距离矩阵
'
' 运行:
' vbs.exe tutorials\VBS\cuda.vb
'
' 流水线:
' VB.NET 函数 -> IL 字节码 -> AST(MethodSyntax) -> .cu 源码
' -> NVRTC 即时编译 -> cuLaunchKernel -> 结果矩阵
'
' 三条内核映射约定(详见 cuda\ILCuda\README.md):
' 1) 形参里有 i -> 一维内核, i = blockIdx.x * blockDim.x + threadIdx.x
' 2) 形参里有 i 和 j -> 二维内核, i 取行(blockIdx.y), j 取列(blockIdx.x)
' 3) 纯标量, 不做数组取元素 -> 逐元素自动包裹, 每个标量参数都变成按下标读取的数组
'
' 注意: 由于 VBS 引擎会把顶层的 Function 重写为匿名函数, 拿不到 Shared 的
' MethodInfo, 所以待反编译的目标函数一律写进下面的 Public Class 里 —— 类型块
' 会被原样搬到生成的 Module 中, 可以正常反射。
' ============================================================================
' ---------------------------------------------------------------------------
' 第 1 节: 待反编译的目标函数
'
' 下面全部是普通的 VB.NET 写法, 不含任何 CUDA 概念。IL2Cuda 会在运行时把它们
' 反编译成表达式树, 再发射成 .cu。
'
' 设第 i 行 sum_i = Σx, sq_i = Σx², dot_ij = Σx_ik·x_jk, n = 列数:
' 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)
'
' 对角线在数学上恒为 corr=1 / dist=0, 这里直接返回, 避免单精度下两个相近
' 大数相减把舍入误差放大(见 README 的数值注意点)。
' ---------------------------------------------------------------------------
Public Class PearsonMetrics
' 行求和, 对应约定 1: 形参 i 为线程索引 -> 一维内核
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
' 行平方和, 同样是一维内核
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
' 两行的点积, 对应约定 2: 同时有 i 和 j -> 二维内核
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
' 皮尔逊相关单元: guard clause + if/else 菱形 + MathF 调用
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
' 欧氏距离单元: 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
' 纯标量截断, 对应约定 3: 没有 i/j, 也没有数组取元素 -> 逐元素自动包裹
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
' ---------------------------------------------------------------------------
' 第 2 节: 教程用的辅助代码(目标方法清单 / CPU 参考实现 / 排版 / 自检)
' ---------------------------------------------------------------------------
Public Class TutorialKit
''' 全部待反编译的目标方法
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
''' 造一个可复现的随机矩阵(行主序)
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
''' 为解释求值自检造一组下标合法、数值不退化的样例参数
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 参考实现: 皮尔逊相关矩阵(双精度两遍算法)
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 参考实现: 欧氏距离矩阵
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
''' GPU 结果与 CPU 参考值之间的最大绝对误差
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
''' GPU 结果与同为单精度的参考值之间的最大绝对误差
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
''' 双精度参考值转单精度, 便于复用同一套预览排版
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
''' 分节标题
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
''' 缩进打印一段多行文本(伪代码 / CUDA 源码)
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
''' 打印结果矩阵的左上角 n x n
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
' ---------------------------------------------------------------------------
' 第 3 节: 主流程
' ---------------------------------------------------------------------------
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 教程: 皮尔逊相关矩阵 与 欧氏距离矩阵")
call console.WriteLine(" 输入规模: " & rows & " 行 x " & cols & " 列, 随机种子 " & seed)
' ---------- 造输入矩阵 ----------
dim data As Single() = TutorialKit.MakeMatrix(rows, cols, seed)
' ---------- 第 1 步: IL -> AST -> .cu ----------
dim kernels As New Dictionary(Of String, IlCudaKernel)(StringComparer.Ordinal)
call TutorialKit.PrintTitle("第 1 步: 把 VB.NET 函数反编译为 AST, 再发射成 .cu")
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 & " 反编译失败: " & errMsg)
allOk = False
else
call kernels.Add(m.Name, kernel)
call console.WriteLine()
call console.WriteLine(" ---- " & m.Name & " ----")
call console.WriteLine(" 索引模式 : " & kernel.IndexMode.ToString())
call console.WriteLine(" 设备函数 : " & kernel.DeviceFunctionName)
call console.WriteLine(" 内核 : " & kernel.Name)
call console.WriteLine()
call console.WriteLine(" 还原出的伪代码:")
call TutorialKit.PrintBlock(SyntaxWriter.WriteMethod(kernel.Syntax), " ")
call console.WriteLine()
call console.WriteLine(" 生成的 CUDA 源码:")
call TutorialKit.PrintBlock(kernel.Source, " ")
' 解释求值自检: 把同一棵 AST 交给 CPU 解释器执行, 与直接调用原方法比对
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(" 解释求值自检: 原方法=" & expected & ", AST=" & actual &
", 差=" & diff.ToString("E3") & " " & (if(pass, "OK", "FAIL")))
if Not pass Then
allOk = False
end if
end if
next
' ---------- 第 2 步: 把生成的 .cu 注册进 KernelSources ----------
call TutorialKit.PrintTitle("第 2 步: 注册内核源码(必须在创建引擎之前)")
for each k As IlCudaKernel In kernels.Values
call k.Register()
call console.WriteLine(" 已注册 " & k.ToString())
next
' ---------- 第 3 步: NVRTC 编译 + 启动内核 ----------
call TutorialKit.PrintTitle("第 3 步: 在 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(" 第 1 步存在失败项, 跳过 GPU 计算。")
elseif engine Is Nothing Then
call console.WriteLine(" GPU 不可用: " & opts.ErrorMessage)
dim report = CudaEnvironment.Probe()
for each s As FixSuggestion In CudaEnvironment.Suggest(report)
call console.WriteLine(" 建议: " & s.ToString())
call console.WriteLine(" " & s.Detail)
next
call console.WriteLine()
call console.WriteLine(" >> 已回退 CPU 参考实现, 下面只给出 CPU 结果。")
else
using engine
dim cells As Integer = rows * rows
call console.WriteLine(" 设备 : " & engine.Device.Name)
call console.WriteLine(" 内核镜像 : " & 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) 行统计量: 一维内核, 一个线程负责一行
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 点积: 二维内核
call kernels("GramDot").Launch(engine, LaunchPlanner.For2D(rows, rows, 16, 16), bufX, cols, bufDot, rows, rows)
' 3) 由点积与行统计量还原相关矩阵与距离矩阵
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) 约定 3 的演示: 纯标量函数被自动包裹成逐元素内核
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(逐元素自动包裹) 最大绝对误差 = " &
clampErr.ToString("E3") & " " & (if(clampPass, "OK", "FAIL")))
if Not clampPass Then
allOk = False
end if
end using
end using
end using
end if
' ---------- 第 4 步: 与 CPU 参考实现比对 ----------
call TutorialKit.PrintTitle("第 4 步: CPU 参考实现比对 + 结果预览")
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(" 皮尔逊相关矩阵 最大绝对误差 = " & eCorr.ToString("E3") & " " & (if(corrPass, "OK", "FAIL")))
call console.WriteLine(" 欧氏距离矩阵 最大绝对误差 = " & eDist.ToString("E3") & " " & (if(distPass, "OK", "FAIL")))
if (Not corrPass) OrElse (Not distPass) Then
allOk = False
end if
call console.WriteLine()
call TutorialKit.PrintMatrix("皮尔逊相关矩阵(GPU, 左上 " & preview & " x " & preview & "):", corrGPU, rows, preview)
call console.WriteLine()
call TutorialKit.PrintMatrix("欧氏距离矩阵(GPU, 左上 " & preview & " x " & preview & "):", distGPU, rows, preview)
else
call console.WriteLine(" GPU 结果不可用, 下面是 CPU 参考实现的结果。")
call console.WriteLine()
call TutorialKit.PrintMatrix("皮尔逊相关矩阵(CPU, 左上 " & preview & " x " & preview & "):", TutorialKit.ToSingle(refCorr), rows, preview)
call console.WriteLine()
call TutorialKit.PrintMatrix("欧氏距离矩阵(CPU, 左上 " & preview & " x " & preview & "):", TutorialKit.ToSingle(refDist), rows, preview)
end if
call console.WriteLine()
if allOk Then
call console.WriteLine("教程全部步骤通过。")
else
call console.WriteLine("存在失败项, 请查看上面的输出。")
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 教程: 皮尔逊相关矩阵 与 欧氏距离矩阵
==========================================================================
输入规模: 1024 行 x 1024 列, 随机种子 42
==========================================================================
第 1 步: 把 VB.NET 函数反编译为 AST, 再发射成 .cu
==========================================================================
---- RowSum ----
索引模式 : Grid1D
设备函数 : il_RowSum_scalar
内核 : il_RowSum_kernel
还原出的伪代码:
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
生成的 CUDA 源码:
// ===== 本文件由 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);
}
解释求值自检: 原方法=-13.5, AST=-13.5, 差=0.000E+000 OK
---- RowSumSq ----
索引模式 : Grid1D
设备函数 : il_RowSumSq_scalar
内核 : il_RowSumSq_kernel
还原出的伪代码:
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
生成的 CUDA 源码:
// ===== 本文件由 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);
}
解释求值自检: 原方法=66.75, AST=66.75, 差=0.000E+000 OK
---- GramDot ----
索引模式 : Grid2D
设备函数 : il_GramDot_scalar
内核 : il_GramDot_kernel
还原出的伪代码:
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
生成的 CUDA 源码:
// ===== 本文件由 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);
}
解释求值自检: 原方法=-14.5, AST=-14.5, 差=0.000E+000 OK
---- CorrelationCell ----
索引模式 : Grid2D
设备函数 : il_CorrelationCell_scalar
内核 : il_CorrelationCell_kernel
还原出的伪代码:
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
生成的 CUDA 源码:
// ===== 本文件由 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);
}
解释求值自检: 原方法=0, AST=0, 差=0.000E+000 OK
---- DistanceCell ----
索引模式 : Grid2D
设备函数 : il_DistanceCell_scalar
内核 : il_DistanceCell_kernel
还原出的伪代码:
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
生成的 CUDA 源码:
// ===== 本文件由 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);
}
解释求值自检: 原方法=0.70710677, AST=0.70710677, 差=0.000E+000 OK
---- PearsonClamp ----
索引模式 : Grid1D
设备函数 : il_PearsonClamp_scalar
内核 : il_PearsonClamp_kernel
还原出的伪代码:
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
生成的 CUDA 源码:
// ===== 本文件由 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]);
}
解释求值自检: 原方法=1, AST=1, 差=0.000E+000 OK
==========================================================================
第 2 步: 注册内核源码(必须在创建引擎之前)
==========================================================================
已注册 il_RowSum_kernel [Grid1D] <- PearsonMetrics.RowSum
已注册 il_RowSumSq_kernel [Grid1D] <- PearsonMetrics.RowSumSq
已注册 il_GramDot_kernel [Grid2D] <- PearsonMetrics.GramDot
已注册 il_CorrelationCell_kernel [Grid2D] <- PearsonMetrics.CorrelationCell
已注册 il_DistanceCell_kernel [Grid2D] <- PearsonMetrics.DistanceCell
已注册 il_PearsonClamp_kernel [Grid1D] <- PearsonMetrics.PearsonClamp
==========================================================================
第 3 步: 在 GPU 上计算相关矩阵与距离矩阵
==========================================================================
设备 : NVIDIA RTX A4000
内核镜像 : NVRTC 13.3 (cubin) arch=sm_86 [C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.3\bin\x64\nvrtc64_130_0.dll]
PearsonClamp(逐元素自动包裹) 最大绝对误差 = 0.000E+000 OK
==========================================================================
第 4 步: CPU 参考实现比对 + 结果预览
==========================================================================
皮尔逊相关矩阵 最大绝对误差 = 2.221E-007 OK
欧氏距离矩阵 最大绝对误差 = 2.069E-005 OK
皮尔逊相关矩阵(GPU, 左上 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
欧氏距离矩阵(GPU, 左上 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
教程全部步骤通过。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.