========================================================================== 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 教程全部步骤通过。