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

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

∑
n−1
i
h = h i x , 则:
设二者在素阶数域上的乘积为
i=0

 s 0 + s n , i = 0



h i =  s i + s i+n + s i+n−1 , 1 ⩽ i ⩽ n−2 (2)


 s n−1 + s 2n−2 , i = n−1
展开后得到一般公式为:

 ∑ n−1
 f 0 g 0 +
 f n−k g k , i = 0


 k=1



 ∑ ∑ ∑
 i n−1 n−1

h i =  f i−k g k + f n+i−k g k + f n+i−1−k g k , 1 ⩽ i ⩽ n−2 (3)

 k=0 k=i+1 k=i



 ∑
 n−1


 f n−1−k g k + f n−1 g n−1 , i = n−1
k=0
进一步地, 用矩阵乘法来表示 h 的系数为:

      
f 0 f n−1 f n−2 ... f 1   0 0 0 ...
 0  g 0   h 0 
      
      
 ... f 2   0 ...    
 

 f 0 f n−1   f n−1 f n−2 f 1  g 1    h 1  
 f 1 
      
      
      
 ... f 3   0 0 ...    h 2  
 
 f 2 f 1 f 0 f n−1 f 2  g 2   (4)
  +    =  
      
 . . . .   . . . .  .   . 
 
 . .    
 
 . . . . .   . . . . .  .   . 
 . .    
 . . . .   . . . .  .   . 
      
      
      
f n−1 f n−2 f n−3 ... f 0 0 0 0 ... f n−1 g n−1 h n−1

f 2 f 1 f 0
g 2 g 1 g 0
f 2 g 0 f 1 g 0 f 0 g 0 f 2 g 0 f 1 g 0 f 0 g 0
= 0 0 0
f 2 g 1 f 1 g 1 f 0 g 1 f 1 g 1 f 0 g 1 +f 2 g 1 f 2 g 1 f 0 f 2 f 1 g 0
f 2 g 2 f 1 g 2 f 0 g 2
f 0 g 2 +f 2 g 2 f 2 g 2 +f 1 g 2 f 1 g 2 0 ×
f 1 f 0 f 2 f 2 f 1 g 1
0 0
f 2 f 1 f 0 f 2 g 2
h 2 h 1 h 0
s 4 s 3 s 2 s 1 s 0
v
M 1 M 2 M 3
图 2 素阶数域上多项式乘法的矩阵转化

3.2.2 基于 Tensor Core 的实现方法与细节
在 GPU 的 C++编程中, 和传统的 CUDA Core 相比, Tensor Core 能进一步加速 MMA 运算, 即 D = A×B+C,
其中 A 为 m×n 的矩阵, B 为 n×k 的矩阵, C 和 D 为 m×k 的矩阵. Tensor Core 使用 wmma::fragment 类来定义不同的
矩阵类型, 其中 A 为矩阵乘法的左矩阵, B 为矩阵乘法的右矩阵, C 和 D 为 MMA 操作中的累加器. Tensor Core 支
持多种混合精度的元素类型和矩阵大小, 即以低精度为输入, 高精度为输出, 表 3 给出了部分能够同时支持 3 种矩
阵大小 (16×16×16、32×8×16、8×32×16) 的左矩阵、右矩阵、累加器的组合精度示例.

表 3 Tensor Core 支持的部分组合精度 [29]

左矩阵 右矩阵 累加器
_half _half float
_half _half _half
unsigned char unsigned char int
signed char signed char int

Tensor Core 所支持的混合精度操作在 CUDA 中提供了相应的编程接口 wmma. 概括来说, 包括 1 个数据类
型 (wmma::fragment) 和 3 个操作 (wmma::load_matrix_sync、wmma::store_matrix_sync 和 wmma::mma_sync), 它
们都需要一个线程束协同完成. 例如, 当 m=16、n=16、k=16 时, 定义在核函数中的 wmma::fragment 对象, 运行时
会以 8 个 8 位整型的寄存器出现在每一个线程的运行空间中. wmma::load_matrix_sync 操作实现从连续内存区中
加载值到矩阵. wmma::store_matrix_sync 操作实现方向相反的操作, 即将矩阵中的值存储到连续的内存区域中. 上