美国橡树岭国家实验室对nvidia v100的性能做了实测，发现只能达到标称理论性能（112T tensor 性能）的74%，即83T tensor性能。我简单解读了一些要点，大家愿意看可以进入下面的链接：
https://arxiv.org/abs/1803.04014
1,性能
首先我们先了解下到底什么是tensor计算，nvidia用新名词偷换概念也是一把好手。
tensor性能，应该是一种混合精度运算性能，oak ridge实验室认为其就是matrix-multiply-and-accumulate（tensor core的情况是在一个时钟周期处理4×4矩阵）性能，即矩阵相乘累加，类似高维版的fma指令。计算形式为C = αAB + βC with α = 1.0 and β = 1.0。
Screenshot from 2018-05-09 18-05-30.png (38.01 KB, 下载次数: 0)
所谓tensor计算如图。在乘法阶段为半精度，累加阶段为单精度。
使用cuda实现该算法有多种方式。
1,汇编以上最底层的cuda操作，cuda kernel function及wmma api，在这个层面上，你可以手动控制GPU的thread/block/grid/warp的调度，极限压榨gpu性能。
// Calculate AB with NVIDIA Tensor Cores
// Kernel executed by 1 Warp (32 Threads)
__global__ void tensorOp(float ∗ D, half ∗ A, half ∗ B) {
// 1. Declare the fragments
wmma::fragment<wmma::matrix_a, M, N, K, half, wmma::colmajor> Amat;
wmma::fragment<wmma::matrix_b, M, N, K, half, wmma::colmajor> Bmat;
wmma::fragment<wmma::accumulator, M, N, K, float, void> Cmat;
// 2. Initialize the output to zero
wmma::fill_fragment(Cmat, 0.0f);
// 3. Load the inputs into the fragments
wmma::load_matrix_sync(Amat, A, M);
wmma::load_matrix_sync(Bmat, B, K);
// 4. Perform the matrix multiplication
wmma::mma_sync(Cmat, Amat, Bmat, Cmat);
// 5. Store the result from fragment to global
wmma::store_matrix_sync(D, Cmat, M, wmma::mem_col_major);
}复制代码
2,底层次封装，cuda 模板类CUTLASS，cuda线性代数模板类。
/// CUTLASS SGEMM example
__global__ void gemm_kernel(
float *C,
float const *A,
float const *B,
int M,
int N,
int K) {
// Define the GEMM tile sizes - discussed in next section
typedef block_task_policy <
128, // BlockItemsY: Height in rows of a tile
32, // BlockItemsX - Width in columns of a tile
8, // ThreadItemsY - Height in rows of a thread-tile
4, // ThreadItemsX - Width in columns of a thread-tile
8, // BlockItemsK - Depth of a tile
true, // UseDoubleScratchTiles - whether to double-buffer SMEM
block_raster_enum::Default // Block rasterization strategy
> block_task_policy_t;
// Define the epilogue functor
typedef gemm::blas_scaled_epilogue<float, float, float> epilogue_op_t ;
// Define the block_task type.
typedef block_task <
block_task_policy_t,
float,
float,
matrix_transform_t::NonTranspose,
4,
matrix_transform_t::NonTranspose,
4,
epilogue_op_t,
4,
true
> block_task_t;
// Declare statically-allocated shared storage
__shared__ block_task_t::scratch_storage_t smem;
// Construct and run the task
block_task_t(
reinterpret_cast(&smem),
&smem,
A,
B,
C,
epilogue_op_t(1, 0),
M,
N,
K).run();
}复制代码
3,高层次封装，cublas/cudnn中的gemm乘法函数。
// First, create a cuBLAS handle:
cublasStatus_t cublasStat = cublasCreate(&handle);
// Set the math mode to allow cuBLAS to use Tensor Cores:
cublasStat = cublasSetMathMode(handle, CUBLAS_TENSOR_OP_MATH);
// Allocate and initialize your matrices (only the A matrix is shown):
size_t matrixSizeA = (size_t)rowsA * colsA;
T_ELEM_IN **devPtrA = 0;
cudaMalloc((void**)&devPtrA[0], matrixSizeA * sizeof(devPtrA[0][0]));
T_ELEM_IN A = (T_ELEM_IN *)malloc(matrixSizeA * sizeof(A[0]));
memset( A, 0xFF, matrixSizeA* sizeof(A[0]));
status1 = cublasSetMatrix(rowsA, colsA, sizeof(A[0]), A, rowsA, devPtrA[i], rowsA);
// ... allocate and initialize B and C matrices (not shown) ...
// Invoke the GEMM, ensuring k, lda, ldb, and ldc are all multiples of 8,
// and m is a multiple of 4:
cublasStat = cublasGemmEx(handle, transa, transb, m, n, k, alpha,
A, CUDA_R_16F, lda,
B, CUDA_R_16F, ldb,
beta, C, CUDA_R_16F, ldc, CUDA_R_32F, algo);复制代码
cublas的batched gemm无法调用tensor core，任何batched方法都未进行优化。
当N=16384时，模板类性能开始比高封装的cuBlas更加有优势。那是因为模板类可以选择不同的tile大小和配置，等于是开放了部分底层参数的封装函数。
N=8192时性能达到高峰，为标称理论性能（112T tensor 性能）的74%，即83T tensor性能。
其他性能对比如下：
Screenshot from 2018-05-09 17-31-43.png (76.67 KB, 下载次数: 0)
tensor core在gemm中提供了比native sp更高的性能，为6倍单精度，3倍半精度
使用native sp，哪怕是使用最底层的wmma实现，在性能上也没有太多惊喜。
Screenshot from 2018-05-09 17-39-23.png (48.74 KB, 下载次数: 0)
使用batched（批量提交）函数，他们无法适用于tensor core，但可以把native sp的性能提升到就4T附近。
2,精度损失
gemm不仅用于深度学习，且适用于几乎任何科学计算，虽然深度学习对精度不敏感，但不代表其他应用也不敏感（比如Montecarlo code，一种基于梦特拉罗法的粒子仿真模拟），为了因对其他类型的计算，精度损失的测量也在此报告中。
在文中，为了压制ground truth之间的误差，该实验室使用了一种Precision refinement机制，简单的说，既然运算器（tensor core的alu）只支持16bit的乘法，那么就将32bit的数值拆分为2个16bit number进行计算。这样，乘法运算变成了两个二项式相乘，但是运算的元素都是半精度的，运算过程虽然复杂化了，但是可以利用tensor core加速，并保持精度。
A single B single = (R A + A half )(R B + B half ) = R A R B + A half R B + R A B half + A half B half .
但是，这种拆分 - 乘法 -合并的过程，并不完全等价直接乘法，所以较ground truth仍然有少量误差。
Screenshot from 2018-05-09 17-55-48.png (58.05 KB, 下载次数: 0)
文章中测试了只precision refine一侧，和precision refine两侧的比较，结果显示只做一侧没有太大意义。两边同时precision refine后，效果拔群。
最后文章还比较了各层次编程的区别，不过没什么太大意义，要用的终究会用。
评论
意思老黄Tsensor的运算能力就像当年VLIW一样，是极限理论性能嘛？
评论
不知道他是怎么算出来了，理论上估计就是用时钟周期乘出来的，但他公布的v100两个版本（pcie和smx2）理论值又不一样，smx2为125T。好像又是考虑了io的。
评论
两个版本（PCI-E和SMX2，不考虑3月份那个）的Max Boost Clock不一样，SMX2的是1455，PCI-E的是1370，PCI-E的TDP低了50W，差异应该源自这里
评论
可能吧，不过差了10%的性能，这频率才差这么点。这数字八成是随便说说的。
评论
4x4矩阵D=AxB+C有64乘法64加法，共128
有640个Tensor Core，128*640*1370 = 112T
评论
这种算法太随意了。
评论
smx2是120T，记错了。
评论
理论浮点性能都是这么算的啊。不管哪个厂商在标理论性能时都是按照每周期能做的最大操作数量乘以时钟频率去算。
评论
GPU的那些浮点性能，像素，纹理填充率 也是都这么按时钟周期算的吧
评论
我可以超到2G。。。我要夸老黄标的实在是太保守了
评论
还好，和我没关系！！
评论
浮点性能上依标的1.5G频率计算，老黄确实很保守了
评论
NVIDIA TITAN V计算出错，后续如何解决的？
评论
目前来看没有大规模爆发，只是次品吧，个例。老实说如果都买去跑深度学习，即使出错也很难看出来。
评论
严格来讲专业卡系列里面没有 泰坦v 泰坦v也没有ecc显存解决我觉得等于没有
评论
那个不是ecc的问题，首先他的错误率达到了10%，除非把你的机器放到福岛核电站里面才有可能出这么多错。ecc修正的是外部物理环境引起的位翻转。
这就是一个设计缺陷/制造缺陷的问题，现在还不确定。就算titan v不用ecc显存，他所预计出的错也是设计范围内的。
评论
位翻转是不是那种比如超新星爆发产生的大量射线引起的错误？
评论
也包括这个。小到比如你摸了一下机器，导致电平信号紊乱也可能产生。
评论
咦……这玩意不是一直标的是理论极限嘛？一开始就和实测值没啥关系啊
评论
对的
厂商想隐藏在阴影里的真相就这么被正义君发现了
要是厂商标的全都能信
那就不需要测试样品了
评论
因为不是专业卡干专业卡的活出错就是正常
评论
那全世界90%做ai的都不专业。。
评论
评论
工商也没辙
评论
我觉得这个问题，可能就跟A卡理论浮点性能高，游戏帧率低，是一个意思
评论
A卡那个是总的流水架构问题。后端容易爆，前端就发挥不出来。
评论
专门跑去问了一下专家，按照橡树岭实验室的测试方法，能达到74%的效率不算差的
评论
当然不算差，这种只激发一个函数的测试方法，本来是作弊的，要跑实际应用，我做了算法复杂度解析后的overall利用率计算，预测支持向量机类应用效率在67%，传统神经网络类应用在49-55%，lstm类神经网络在43%，胶囊神经网络在44%。
可以说在绝大部分应用里面能用到50%就合格了。
评论
这么说的话，其实全世界计算设备的理论指标/性能都在虚标...
《超算性能，到底能够虚标多少》
《CPU性能，到底能够虚标多少》
好像也能按这个套路写出来？
话不能这么说，跑科学计算cpu的利用率是很高的，即使是用显卡的流处理器跑gemm，利用率肯定也是大于90%的。唯独这个tensor core利用率虚标严重。
评论
我第一反映是，居然有74%的效率，那么高。。。sp效率>80%，都上了啥优化？？？linpack上，Intel的xeon CPU跑gemm也不过80-85%的效率。
评论
你这是用库跑，橡树岭测试是用核函数，下面就是汇编了，核函数里面也可以嵌汇编优化代码。
评论
不对啊，他们不是跑的大矩阵乘法么？我记得sandra 2017跑GPU的sgemm，结果也只有60%不到的效率啊，这才是比较符合我的理解啊。
评论
核函数跑矩阵乘法，参考第一段代码。
评论
这里的核函数不是支持向量机的核函数，是指cuda中直接调度核心的kernel function。。。 电路 电子 维修 求创维42c08RD电路图 评论 电视的图纸很少见 评论 电视的图纸很少见 评论 创维的图纸你要说 版号，不然无能为力 评论 板号5800-p42ALM-0050 168P-P42CLM-01 电路 电子 维修 我现在把定影部分拆出来了。想换下滚，因为卡纸。但是我发现灯管挡住了。拆不了。不会拆。论坛里的高手拆解过吗？ 评论 认真看，认真瞧。果然有收
·日本中文新闻 唐田绘里香为新剧《极恶女王》剃光头 展现演员决心
·日本中文新闻 真子小室夫妇新居引发隐私担忧
·日本中文新闻 前AKB48成员柏木由纪与搞笑艺人交往曝光
·日本学校 ｛日本国际学校｝梅田インターナショナルスクール
·日本学校 LINE：sm287 陳雨菲、20歳、台湾からの留学生、東京に来たばかり
·日本留学生活 出售平成22年走行48000km 代步小车
·日本华人网络交流 円相場 一時1ドル＝140円台まで上昇？
·日本华人网络交流 问日本华人一个问题
·日本旅游代购 富山接机
·生活百科 英国转澳大利亚转换插头
·汽车 【求助】修车遇到困难怎么办？