Page 338 - 《软件学报》2026年第7期

胡晓雯 等: 基于 CUDA Core 和 Tensor Core 的 CTRU-Prime 高吞吐量实现 3023

warpgroup 计算待访问 G h 、 G l 的位置偏移, 以 warpgroup/2 计算待访问 F 的位置偏移. 其中, G h 、 G l 以行主序存
储于全局内存, F 则以列主序存储于全局内存.

算法 7. Twmma.

输入: 矩阵 A (大小为 A x ×A y ) 对应的连续内存 array_a; 矩阵 B (大小为 B x ×B y ) 对应的连续内存 array_b;
输出: 存储结果 C = A×B 的连续内存 array_c.

1. wmma::fragment<wmma::matrix_a, m, n, k, int8_t, wmma::row_major> a_frag;
2. wmma::fragment<wmma::matrix_b, m, n, k, int8_t, wmma::col_major> b_frag;
3. wmma::fragment<wmma::accumulator, m, n, k, int32_t> c_frag;
4. tot_warpID = (blockIdx.x × blockDim.x + threadIdx.x) / 32;
5. warpgroup = tot_warpID / WRAP_NUMS;
6. warpID = tot_warpID – warpgroup × WRAP_NUMS;
7. row_idx = warpID % (( B y ) / k) × k;
8. col_idx = warpID / (( B y ) / k) × m;
9. st_offset = col_idx + row_idx × A x ;
10. wmma::fill_fragment(c_frag, 0);
11. for i from 0 to ( A y ) / n do
12. ldA_offset = col_idx × ( A y ) + i × n;
13. ldB_offset = row_idx × ( B x ) + i × n;
×
14. wmma::load_matrix_sync(a_frag, &array_a[ A x A y × warpgroup] + ldA_offset, A y );
B x B y × (warpgroup / 2)] + ldB_offset, B x );
×
15. wmma::load_matrix_sync(b_frag, &array_b[
16. wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
17. end for
×
18. wmma::store_matrix_sync(&array_c[ A x B y × warpgroup] + st_offset, c_frag, A x , wmma::mem_col_major);
在 IoT 场景下, 由于资源受限, 所有节点在特定会话中会使用同一公私钥对进行密钥封装与解封装, 这意味着
在一次通信会话中使用相同的公私钥对执行数百到数千次 KEM 是很常见的 [30] . 考虑到上述方法需填充 0, 为增大
计算吞吐量, 提出面向相同公私钥的多项式乘法实现. 在该过程中, 相同的公私钥将展开为公式 (4) 中的矩阵,
CTRU-Prime.PKE.Dec 中的密文 c 和 CTRU-Prime.PKE.Enc 中的随机多项式 r 将展开为公式 (4) 中的向量, 多个向
量可以无缝拼接为矩阵而无需填充 0, 从而避免计算资源的浪费.
4 面向吞吐量的 CTRU-Prime 总体架构与优化实现

在本节中, 主要阐述面向吞吐量的 CTRU-Prime 实现方案的总体架构、优化策略以及多项式乘法以外的主要
底层模块的 GPU 实现方案. 本文面向吞吐量设计, 基于 GPU 的不同特性, 提出不同的优化策略, 以最大化利用
GPU 的高并行性特点, 充分释放 GPU 的潜能.
4.1 设计概述
针对 CTRU-Prime 中的不同操作, 基于简单模式, 本文提取其并行度, 在最优数据划分的基础上, 使用多线程
协同完成以减少操作延迟. 在核函数设计过程中, 以优化占用率为目标, 合理分配核函数使用的寄存器数、共享内
存数和线程数, 从而最大化 SM 上活跃线程束数. 此外, 本文使用合并内存访问、转移数据交换到寄存器和共享内
存、使用页锁定内存等内存管理技术, 进一步减少核函数的数据访问延迟, 提高核函数整体执行效率.
在此基础上, 本文充分考虑各操作之间的控制流和数据流关系. 鉴于 CTRU-Prime 无分支语句, 控制流较为简