工程文件和nvvp：
https://pan.baidu.com/s/1G6n4SWzcRIM2z0NUYsR85A
昨晚有空做了测试，汇总一下结果。
首先是两者超参数下的性能。当然这两者规格区别大，横向对比就免了，看看自己与自己比。要横比等我搞到rtx titan先（体彩已买。）
超参数如下：
int batch_size_hparam[1] = { 128 };
int input_channel_hparam[1] = { 32 };
int input_height_hparam[3] = { 64, 128, 256 };
复制代码时间宝贵，只安排了三组。
结果：
titanv.JPG (116.72 KB, 下载次数: 0)
rtx2080.JPG (122.61 KB, 下载次数: 0)
可以看得出nn.forward性能titan v的tensor core有sp的一倍速度，而rtx2080基本上与sp的性能相当。是否能验证turing卡的tensor core fp16的peak mac（乘积与累加）性能被阉割一半不可知（cublas的代码已经写好了，还有些错误调下，以后在测试，下面提供源码）。
其实这两代卡的sp性能也非常强悍，在alu群的负载不饱满的情况下，即便是计算单精度浮点，sp性能也是优于tensor core的。
turing的多倍int推理性能能否可用？
目前测试的结果是，不可用，而且按照说明书中的说法，这个多倍int性能目前只为sm72准备，注意，既不是volta的sm70，也不是turing的sm75，大概是xavier还是什么推理专用核心使用的。
这应该只是软件上的问题，期待后续cuda版本更新开放给sm70/75。
相同的问题：
https://www.reddit.com/r/nvidia/ ... n_data_int8x32_for/
titanv与rtx2080的nvprof分析。
主要解析两者计算过程。
1，单精度浮点计算过程对比：
ffspvs.png (347.13 KB, 下载次数: 0)
其中的scudnn代表使用的是sp进行计算，首字母s代表单精度，半精度计算时为hcudnn，没截图，就不一一展示了，后面会提供nvprof的nvvp文件。
计算时，同规格超参数下原子级计算的延迟对比：
rtx2080：
ffsprtx2080 lag.png (125.98 KB, 下载次数: 0)
titanv:
titanvffsp lag.png (292.78 KB, 下载次数: 0)
不用吐槽名字的开头，命名都是以volta开头说明这个函数从volta时代没有更新而已。
总之延迟对得起规格的差距。没有负优化。
2，tensor core的半精度运算（半精度输入，半精度输出）。
h h tc.png (418.22 KB, 下载次数: 0)
这里显示的h885cudnn/h1688cudnn即代表目前在计算的是tensor core。
参考：
http://tadaoyamaoka.hatenablog.com/entry/2018/10/08/185803
https://mxnet.incubator.apache.org/faq/float16.html
延迟：
titan v：
titanv hhtc lag.png (128.15 KB, 下载次数: 0)
rtx2080：
rtx2080 hh tc lag.png (123.43 KB, 下载次数: 0)
感觉截错了。
2，tensor core的半精度运算（半精度输入，单精度输出）。
hftc vs.png (279.53 KB, 下载次数: 0)
比较有意思的是volta的算法先进一点（v1），计算的同时进行了矩阵的resharp（把channel放到了最后，据说可以加快运行速度），而turing目前还需要一个单独的操作。
不过turing也会更新的。
总结：
cuda目前还是很多原型方法，每次更新都会改不少东西，不太稳定。在跑运算的时候，选择的算法经常乱跳（后来统一固定为了CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM），内部根据不同的数据片大小，还会内部不停优化。
其实用cublas的gemmex测试更加原子。代码如下：
cublas源码（尚未测试）：
typedef float the_data_type;
#include "cuda_runtime.h"
#include "cublas_v2.h"
#include <iostream>
// Linux
#include <sys/time.h>
#include <boost/multi_array.hpp>
using namespace std;
int const width = 8;
int const height = 16;
int result_edge = 8;
// cublas use column for leading dim
int lda = 16, ldb = 8, ldc = 8;
double cublas_get_wall_time() {
struct timeval time;
if (gettimeofday(&time, NULL)) {
return 0;
}
return (double) time.tv_sec + (double) time.tv_usec * .000001;
}
int cublasmain() {
cublasStatus_t status;
float *h_matrix_A = (float*) malloc(width * height * sizeof(float));
float *h_matrix_B = (float*) malloc(height * width * sizeof(float));
float *h_result = (float*) malloc(
result_edge * result_edge * sizeof(float));
for (int i = 0; i < width * height; i++) {
h_matrix_A[i] = (float) (rand() % 10 + 1) * 0.1;
h_matrix_B[i] = (float) (rand() % 10 + 1) * 1;
}
cout << "A:" << endl;
for (int i = 0; i < width * height; i++) {
cout << h_matrix_A[i] << " ";
if ((i + 1) % width == 0)
cout << endl;
}
cout << endl;
cout << "B:" << endl;
for (int i = 0; i < width * height; i++) {
cout << h_matrix_B[i] << " ";
if ((i + 1) % height == 0)
cout << endl;
}
cout << endl;
cublasHandle_t handle;
status = cublasCreate(&handle);
cublasSetMathMode(handle, CUBLAS_TENSOR_OP_MATH);
if (status != CUBLAS_STATUS_SUCCESS) {
if (status == CUBLAS_STATUS_NOT_INITIALIZED) {
cout << "CUBLAS_STATUS_NOT_INITIALIZED" << endl;
}
getchar();
return EXIT_FAILURE;
}
the_data_type* d_matrix_A { nullptr };
the_data_type* d_matrix_B { nullptr };
the_data_type* d_result { nullptr };
int matrix_bytes = width * height * sizeof(the_data_type);
cudaMalloc(&d_matrix_A, matrix_bytes);
// cudaMemcpy(d_matrix_A, h_matrix_A, matrix_bytes, cudaMemcpyHostToDevice);
cublasSetMatrix(width, height, sizeof(float), h_matrix_A, width,
d_matrix_A, width);
/* check A */
float *h_matrix_A_check = (float*) malloc(width * height * sizeof(float));
cublasGetMatrix(width, height, sizeof(float), d_matrix_A,
width, h_matrix_A_check, width);
cout << "check A锛? << endl;
for (int i = 0; i < width * height; i++) {
cout << h_matrix_A_check[i] << " ";
if ((i + 1) % width == 0)
cout << endl;
}
cout << endl;
cudaMalloc(&d_matrix_B, matrix_bytes);
// cudaMemcpy(d_matrix_B, h_matrix_B, matrix_bytes, cudaMemcpyHostToDevice);
cublasSetMatrix(height, width, sizeof(float), h_matrix_B, height,
d_matrix_B, height);
/* check B */
float *h_matrix_B_check = (float*) malloc(width * height * sizeof(float));
cublasGetMatrix(height, width, sizeof(float), d_matrix_B,
height, h_matrix_B_check, height);
cout << "check B锛? << endl;
for (int i = 0; i < width * height; i++) {
cout << h_matrix_B_check[i] << " ";
if ((i + 1) % height == 0)
cout << endl;
}
cout << endl;
int result_bytes = result_edge * result_edge * sizeof(the_data_type);
cudaMalloc(&d_result, result_bytes);
// cudaMemset(d_result, 0, result_bytes);
cublasSetMatrix(8, 8, sizeof(float), h_result, 8, d_result, 8);
int a = 1, b = 0;
float alpha = 1, beta = 0;
// double start_time_1 = cublas_get_wall_time();
//
// for (int i = 0; i < 1000000; i++) {
//
// cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_T, width,/*row number of matrix_A*/
// width,/*column number of matrix B, should be equal to row number of matrix_A*/
// height,/*column number of matrix_B, should be equal to column number of matrix_A*/
// &alpha, d_matrix_A, lda,/*number of long edge in matrix_A*/
// d_matrix_B, ldb,/*number of long edge in matrix_A*/
// &beta, d_result, ldc/*number of long edge in result*/
// );
//
// }
//
// double end_time_1 = cublas_get_wall_time();
// cout << end_time_1 - start_time_1 << endl;
//
// // can use cublasGetMatrix
// cudaMemcpy(h_result, d_result, result_bytes, cudaMemcpyDeviceToHost);
//
// cout << "result with cuda core锛? << endl;
// for (int i = 0; i < result_edge * result_edge; i++) {
// cout << h_result[i] << " ";
// if ((i + 1) % result_edge == 0)
// cout << endl;
// }
// cout << endl;
double start_time_2 = cublas_get_wall_time();
for (int i = 0; i < 1; i++) {
cublasGemmEx(handle, CUBLAS_OP_T, CUBLAS_OP_T, 8,/*row number of matrix_A*/
8,/*column number of matrix B, should be equal to row number of matrix_A*/
16,/*column number of matrix_B, should be equal to column number of matrix_A*/
&a, d_matrix_A, CUDA_R_16F, 16, d_matrix_B, CUDA_R_16F, 8, &b,
d_result, CUDA_R_16F, 8, CUDA_R_32F,
CUBLAS_GEMM_DEFAULT_TENSOR_OP);
}
double end_time_2 = cublas_get_wall_time();
cout << end_time_2 - start_time_2 << endl;
cublasGetMatrix(8, 8, sizeof(float), d_result,
8, h_result, 8);
cout << "result with tensor core锛? << endl;
for (int i = 0; i < result_edge * result_edge; i++) {
cout << h_result[i] << " ";
if ((i + 1) % result_edge == 0)
cout << endl;
}
cout << endl;
free(h_matrix_A);
free(h_matrix_B);
free(h_result);
cudaFree(d_matrix_A);
cudaFree(d_matrix_B);
cudaFree(d_result);
cublasDestroy(handle);
return 0;
}
复制代码
评论
大佬，先膜.
评论
太专业的看不懂，能通俗点翻译一下吗
评论
树导，我真看不懂啊。。。。。。。。。。。。。
评论
这有啥看不懂的，前两张图是重点，titan v比rtx2080好一些，建议买。
评论
TITAN RTX霸王龙要出了。
评论
图灵cuda部分是改进版伏打，并不是马甲，ttv只有少数游戏能超过2080ti
https://www.tomshardware.com/rev ... x-2080-gaming-oc-8g,5879-3.html
评论
这个要综合对比 ttv 2080ti rtx tt 性能价格综合比对比较好，2080感觉不是一个档次的。
希望rtx tt能成为众望所归的爆款。
话说树导应该不会错过这个新旗舰产品的
评论
不是干这一行的，我就来发个邪恶，凑凑热闹
评论
反过来说的话。
考虑到价格。2080的表现还不错
评论
显存太小了，而且gddr6的延迟貌似比hbm2高。
评论
rtx titan性能可以预测，其实没啥惊喜的。
评论
别拆台嘛。一下就暴露出我不懂了xddd。
话说树导我上次加你origin咋没反应233.
评论
貌似目前老黄对titan系列的定位就是用于DL的专业卡了
评论
貌似个啥？反正你和树导都是要真香的人，到时测测不就知道了？
评论
不是很想买了。。毕竟我又不玩DL
等宽带树收藏啊！她一定会买的，满血核心！
讲道理我想把2080这个残废出掉换Titan的
评论
加了你了，经常看你玩战地5多人游戏，我就玩玩单机。
评论
战地单机没意思啊。
又短又逗。还是多人好玩
评论
Titan V，好贵鸭
评论
单位的两张2080ti在路上了，其他的都期货买不到啊
评论
2张2080的价格，真香。。。。我再想等RTX TT出来还是先入2080ti。。。
评论
其实这还是游戏卡 不是专业卡 计算卡 老黄不会傻到泰坦就是廉价专业卡 它家专业卡 计算卡别卖了 其实专业卡 计算卡老黄也不是真零售一般都是卖配套服务 软件优化
评论
的确是，专业应用没nv的服务根本玩不转，那个nvporf中tensor core调用的显示说明书里面根本没写。google一搜0条记录。。要不是mxnet写文档的时候随便带了一句谁知道这个。
评论
而且不止ecc ai计算也可能出错 但是泰坦出错率更高 还没有ecc 最后发现泰坦只是一张不划算的游戏旗舰公版卡但是带了专业卡 计算卡东西 但是扯谈他是专业卡 计算卡 就搞笑了
评论
老黄算盘打的精 泰坦就是给个人工作站 小企业的廉价计算卡 但是一旦碰上真正问题造成亏本 那可怪不了老黄
评论
跑深度学习不需要高精度，偶尔出错也没什么，因为这类统计学算法的本质就是修正误差的过程，自己等于是个大号的纠错器。顶多影响一点收敛性能。
这是游戏卡能做深度学习的原因，要求高精度和ecc的算法都不能用游戏卡算。
评论
就喜欢看你们聊专业人士怎么来花钱的
评论
如果是真生产力专业卡 计算卡只能买高价昂贵东西 泰坦只能客串来点真格的直接亏本
评论
https://pan.baidu.com/s/1G6n4SWzcRIM2z0NUYsR85A
工程文件，包含所有代码和nvvp，想看的想自己跑跑的可以参考下。
评论
所以我最近特纠结到底是买RX VEGA FE还是WX7100 囧
赚钱高价专业卡 只是自己乐呵 vega 16g足够
评论
奇怪 看lambda的测试结果, 2080Ti和Titan V的FP16性能差不多, 都能提升60%左右的效率, 但是2080真的没有多少fp16的提升, 老黄这是什么操作? 感觉2080单独被阉割了.
https://lambdalabs.com/blog/best ... -1080-ti-benchmark/
顺便吐槽一下RTX Titan这散热咋放服务器/工作站? 都不能临着放....还不如买2080Ti turbo....
评论
楼主专业的!
评论
每一个数字和字母都认识，连起来看不懂
评论
先前Titan V開箱文的樓主某壕就"欽定"Titan V是專業卡
某壕只跑benchmark玩耍遊戲，不跑繪圖軟件不做運算不做DL或AI相關的東西還買一張披著"有點專業"性能的遊戲卡
我只能說消費從沒降級，呵呵
评论
啥时候测一下RTX8000和6000就好了
评论
因为老黄没有更大规格的核心更大显存游戏卡给你非要搞出个游戏卡带专业卡 计算卡一点东西 但是专业方面又各种限制的多功能专业卡 有点像 以前iu 至强 i7多核 指令集有点小区别 针对特殊专业工作软件
评论
其實老黃不是硬件思維而是驅動軟件思維，用軟體強迫搭配硬件
現在沒有人做驅動破解了，剛需只能乖乖繳保護費給老黃了
评论
术业有专攻
评论
哎呦喂，这不是那个国文不好但又特喜欢天天看SPECviewperf的那位吗？肯定你的键盘特好用，对吧？ 我稍微看了一下你的记录，除了说你键盘好用就是说别人的不好用，就没干别的了，还好意思到处说？tw省真的是太牛了
评论
鬼島小民的中文程度怎樣都不差的，倒是有人整天拿benchmark說事這點...能不能請阿鬼好好用中文？
Specviewperf的公信力早就被質疑多年了，但是沒在繪圖的人會懂？
TTV就一張公公卡，跑跑現實世界真正拿來生產用的軟件再來討論好嗎？口袋裡有錢跟認知程度的高低並沒有相關性與可比性的
评论
原来以为你就中文不好，就喜欢说自己键盘好而已，原来英文也看不懂了，我那图是你说的那样吗？
唉，跟你说多了拉低我的键盘了，随你便吧，你爱怎么说你键盘有多好是你的权利，我尊重你有这权利！
反正你天天干的就是说自己键盘好，别人的键盘不好，没别的干了，妨碍你这爱好是我的不是
评论
这个测试太不原子了。里面包含大量数据交换和队列的串行计算，这样显存只有8G的2080会很吃亏的。可能还有别的因素，这样的测试很难分析原因。
评论
我试着编译了cudnn的代码在2080Ti跑了一下, 结果好像特别惨烈 (我直接没用过nvcc, 可能是我用错了), 看起来2080Ti的TC性能也确实被阉割了, lambda的实验确实不够原子.
我把结果贴在下面了
谢谢!
DEFAULT_MATH (HALF, HALF):
CUDA runtime: 10000 | CuDNN: 7301
h h sp
selected algorithm: 7
13.3564
selected algorithm: 7
53.6991
selected algorithm: 7
220.735
DEFAULT_MATH (FLOAT, FLOAT):
selected algorithm: 1
7.82729
selected algorithm: 1
31.0167
selected algorithm: 6
127.38
CUDNN_TENSOR_OP_MATH (HALF, HALF):
selected algorithm: 7
7.45516
selected algorithm: 7
29.4903
selected algorithm: 7
122.945
我用的编译命令
nvcc -o cudnn cudnnmain.cu -lcublas -lcudadevrt -lcudnn复制代码
评论
你的第一个测试是用sp算半精度，理论上turing的sp都有双倍半精度能力，实际上测试结果这个是最差的，甚至比单精度还慢的多。这个可能是编码上的问题，可能需要在进行计算前，显式的把数据都转换为半精度。我的结果也是这样的，所以我就没用采用这个数据。
在用tensor core计算时性能最好的应该时half float，因为这是最标准的tensor core的mac，两个半精度相乘累加到一个单精度位上作为输出。
如果时half half，最后还要多一道工序，就是将单精度再量化为半精度输出，速度会比half float慢一点。
评论
感谢指点, 用了V100试了下确实half float能提升10%左右,TC也是接近2倍性能. 图灵游戏卡在TensorCore被阉割的情况下还是可以双倍半精度吗? 因为看起来就算有TC也就是双倍FP32性能的样子.
V100: CUDA 9 + CUDNN 7.4.1
ff sp
selected algorithm: 1
7.07143
selected algorithm: 1
26.6963
selected algorithm: 1
105.592
hf tc
selected algorithm: 1
3.80806
selected algorithm: 1
13.9783
selected algorithm: 1
55.6275
评论
用PyTorch写了个简单脚本, 2080Ti 开启Apex (Nvidia出的一个混合精度训练插件)的情况下, DenseNet可以获得80%的速度提升(FP16的batch size x 2的情况下), 估计如果没有NormLayer可以更快一点. 虽然这个测试不太原子, 但可以证明确实图灵游戏卡可以用FP16在深度学习上获得很大的性能提升. 1080系没有这个功能, 相比之下2080的性价比还有救.
https://gist.github.com/matthew- ... 35780af361fab6ac82f
评论
TFBENCHMARK.png (45.31 KB, 下载次数: 0)
放一个测了一半的TFBenchmark结果
环境：CUDA9.0 + CUDNN7.4 Tensorflow-GPU 1.12（不带AVX2）
CPU:AMD ThreadRipper 1950X
RAM:4*DDR4 2400 8GB
评论
横坐标是examples/sec吗
为啥你这1080Ti在FP16上有提升, 很多评测都显示10系FP16性能提升基本就是1-3%这样. 代码能贴一下吗? 我感觉你是不是input pipeline没prefetch或者模型参数还是fp32.
我那个脚本(synthetic input) 刚在1080Ti上测试了下, 在Res50 v2的FP16和FP32的速度是190和186, 几乎没提升. 在2080Ti上FP16和FP32的速度分别是489和248, 感觉跟你的fp16结果好像出入有点大.
PS: 因为fp16只占一半的空间, 所以我FP16下batch size增大了一倍
我环境是1900X+2133 DDR4..... CUDA 10 + CUDNN 7.4.1
评论
横坐标是examples/sec
只是跑了一遍https://github.com/tensorflow/benchmarks，粗略地看了一下似乎有prefetch
评论
估计这个repo的fp16没写对? 或者是你驱动/CUDA不够新? 看issue里作者说fp16在V100 Fp16 Res50能有3x的提升....https://github.com/tensorflow/benchmarks/issues/273
评论
驱动应该足够新(416.94，测试时的最高版本)...cuDNN也是最新的，难道是CUDA9.0的问题？
评论
应该没什么问题，阉割的只是tensor core，sp的能力不会阉割。半精度图像渲染也用的上。
评论
试试编译一个CUDA 10? BTW, 我试了下V100 ResNet50的速度
FP32: 322
FP16: 835
确实恐怖, 勉强到3了
评论
试试编译一个CUDA 10的版本? BTW, 我试了下V100 ResNet50的速度
FP32: 322
FP16: 835
确实恐怖, 勉强到3了, 目测RTX Titan可以上900....
评论
白天有空试试，似乎两张2080TI闷在一个罐子里把网卡搞过热了，这会儿怎么都连不上实验室的机器
评论
https://docs.nvidia.com/cuda/cud ... hmetic-instructions
评论
折腾了一番Windows下编译tf1.12+cuda10似乎有问题
".\tensorflow/compiler/tf2xla/cpu_function_runtime.h(71): error C2338:"
评论
cublas线代库可以测试，有人写了代码。2080的半精度是正常的，为单精度的两倍，没阉割。当然，alu群的利用率要起来，否则效率可能还不如单精度。观察发现半精度只省一点显存，2G空间只省200mb左右。
https://github.com/hma02/cublasHgemm-P100
cublas:
running cublasSgemm test
running with min_m_k_n: 2 max_m_k_n: 8192 repeats: 1000
allocating device variables
float32; size 2 average: 1.95369e-05 s
float32; size 4 average: 1.97402e-05 s
float32; size 8 average: 2.24297e-05 s
float32; size 16 average: 2.58005e-05 s
float32; size 32 average: 2.86604e-05 s
float32; size 64 average: 2.62252e-05 s
float32; size 128 average: 2.37135e-05 s
float32; size 256 average: 2.70684e-05 s
float32; size 512 average: 6.13298e-05 s
float32; size 1024 average: 0.000246053 s
float32; size 2048 average: 0.00178317 s
float32; size 4096 average: 0.013566 s
float32; size 8192 average: 0.110101 s
running cublasHgemm test
running with min_m_k_n: 2 max_m_k_n: 8192 repeats: 1000
allocating device variables
float16; size 2 average: 2.79607e-05 s
float16; size 4 average: 2.59722e-05 s
float16; size 8 average: 2.65561e-05 s
float16; size 16 average: 2.79748e-05 s
float16; size 32 average: 4.39696e-05 s
float16; size 64 average: 3.24604e-05 s
float16; size 128 average: 4.2616e-05 s
float16; size 256 average: 4.70025e-05 s
float16; size 512 average: 5.7532e-05 s
float16; size 1024 average: 0.000158499 s
float16; size 2048 average: 0.00089542 s
float16; size 4096 average: 0.00679688 s
float16; size 8192 average: 0.0533021 s
评论
按照之前NV Turing架构白皮书上的spec，Geforce RTX2080/2080TI被阉割的部分似乎是FP16 with Tensor Core
是的。
评论
分享一个cublas tc代码
// build with command nvcc -lcublas -lcudart -lcurand -arch=sm_70 half.cu
// use max clock
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <curand.h>
#include <cublas_v2.h>
// Must be multiples of 16 to fit TensorCore
#define SIZE 8192 // 4096 8192 10240 16384 24576
#define MATRIX_M SIZE
#define MATRIX_N SIZE
#define MATRIX_K SIZE
#define num_clock 1530
#define num_SM 80
#define num_TC 8
#define num_FMA 2
#define num_mma 64
#define FP16_OP num_clock*num_SM*num_TC*num_FMA*num_mma
#define TOTAL_OP MATRIX_M * MATRIX_N * MATRIX_K * 2
#define TOTAL_OP2 (MATRIX_M*MATRIX_N) * (2*MATRIX_K+2)
__global__ void convertFp32ToFp16 (half *out, float *in, int n);
__global__ void convertFp16ToFp32 (float *out, half *in, int n);
int main(int argc, char* argv[]) {
printf("FP32 Matrix Memory Size : %f \n", (float) (sizeof(float) * (float) (MATRIX_M*MATRIX_M) / ( 1024 * 1024 ) ) );
printf("FP16 Matrix Memory Size : %f \n", (float) (sizeof(half) * (float) (MATRIX_M*MATRIX_M) / ( 1024 * 1024 ) ) );
float *a_fp32;
float *b_fp32;
float *c_fp32;
half *a_fp16;
half *b_fp16;
half *c_fp16;
float *c_cublas_fp32;
float *c_host_cublas_fp32;
half *c_cublas_fp16;
half *c_host_cublas_fp16;
printf(" Step1. Initialize GPU API handles...\n");
curandGenerator_t gen;
cublasHandle_t cublasHandle;
cublasCreate(&cublasHandle);
cudaEvent_t startcublas;
cudaEvent_t stopcublas;
cudaEventCreate(&startcublas);
cudaEventCreate(&stopcublas);
// Use tensor cores
cublasSetMathMode(cublasHandle, CUBLAS_TENSOR_OP_MATH);
printf(" Step2. Memory Mallocation ...\n");
cudaMalloc((void**)&a_fp32, MATRIX_M * MATRIX_K * sizeof(float));
cudaMalloc((void**)&b_fp32, MATRIX_K * MATRIX_N * sizeof(float));
cudaMalloc((void**)&c_fp32, MATRIX_M * MATRIX_N * sizeof(float));
cudaMalloc((void**)&a_fp16, MATRIX_M * MATRIX_K * sizeof(half));
cudaMalloc((void**)&b_fp16, MATRIX_K * MATRIX_N * sizeof(half));
cudaMalloc((void**)&c_fp16, MATRIX_M * MATRIX_N * sizeof(half));
cudaMalloc((void**)&c_cublas_fp32, MATRIX_M * MATRIX_N * sizeof(float));
cudaMalloc((void**)&c_cublas_fp16, MATRIX_M * MATRIX_N * sizeof(half));
c_host_cublas_fp32 = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float));
printf(" Step3. Data init with cuRAND ...\n");
curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(gen, 1337ULL);
curandGenerateUniform(gen, a_fp32, MATRIX_M * MATRIX_K);
curandGenerateUniform(gen, b_fp32, MATRIX_K * MATRIX_N);
curandGenerateUniform(gen, c_fp32, MATRIX_M * MATRIX_N);
cudaMemcpy(c_cublas_fp32, c_fp32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToDevice);
printf(" Stsep4. convert FP32 to FP16 for FP16 benchmark...\n");
// curand doesn't currently support fp16 so we generate in fp32 and convert to fp16.
convertFp32ToFp16 <<< (MATRIX_M * MATRIX_K + 255) / 256, 256 >>> (a_fp16, a_fp32, MATRIX_M * MATRIX_K);
convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N);
convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_fp32, MATRIX_M * MATRIX_N);
curandDestroyGenerator(gen);
half alpha = 2.0f;
half beta = 2.0f;
printf(" Step5. Ready to Run...\n");
printf("\nM = %d, N = %d, K = %d. alpha = %f, beta = %f\n\n", MATRIX_M, MATRIX_N, MATRIX_K, alpha, beta);
// Now using cuBLAS
printf("warm up...");
cublasHgemm(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N,
MATRIX_M, MATRIX_N, MATRIX_K,
&alpha,
a_fp16, MATRIX_M,
b_fp16, MATRIX_K,
&beta,
c_fp16, MATRIX_M);
printf(" Step6. Running with cuBLAS...\n");
cudaEventRecord(startcublas);
cublasHgemm(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N,
MATRIX_M, MATRIX_N, MATRIX_K,
&alpha,
a_fp16, MATRIX_M,
b_fp16, MATRIX_K,
&beta,
c_cublas_fp16, MATRIX_M);
cudaEventRecord(stopcublas);
// Error checking
printf(" Step7. Download results...\n");
convertFp16ToFp32 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_cublas_fp32, c_cublas_fp16, MATRIX_M * MATRIX_N);
cudaMemcpy( c_host_cublas_fp32, c_cublas_fp32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost);
float cublasTime;
cudaEventSynchronize(stopcublas);
cudaEventElapsedTime(&cublasTime, startcublas, stopcublas);
printf("cublas took %fms", cublasTime);
printf(" with Operation %.2f\n", (double) TOTAL_OP );
printf(" RPeak FP16 TFLOPS: %.2f with max clock %d Mhz \n", (double) FP16_OP /(1000000) , num_clock );
printf(" RMax FP16 TFLOPS %.2f\n", (double) ( ((double)TOTAL_OP / (double) (1000000) ) / ((double) cublasTime)/1000 ) );
cudaEventDestroy(startcublas);
cudaEventDestroy(stopcublas);
cudaFree(a_fp32);
cudaFree(b_fp32);
cudaFree(c_fp32);
cudaFree(a_fp16);
cudaFree(b_fp16);
cudaFree(c_fp16);
cudaFree(c_cublas_fp32);
cudaFree(c_cublas_fp16);
free(c_host_cublas_fp32);
cudaDeviceReset();
return 0;
}
__global__ void convertFp32ToFp16 (half *out, float *in, int n) {
int idx = blockDim.x * blockIdx.x + threadIdx.x;
if (idx < n) {
out[idx] = in[idx];
}
}
__global__ void convertFp16ToFp32 (float *out, half *in, int n) {
int idx = blockDim.x * blockIdx.x + threadIdx.x;
if (idx < n) {
out[idx] = in[idx];
}
}复制代码
评论
这代码比较老了，现在官方提供了fp16和fp32转换的函数。在cuda_fp16.h里面。还提供了一个新的half类型叫half2。 电路 电子 维修 我现在把定影部分拆出来了。想换下滚，因为卡纸。但是我发现灯管挡住了。拆不了。不会拆。论坛里的高手拆解过吗？ 评论 认真看，认真瞧。果然有收 电路 电子 维修 求创维42c08RD电路图 评论 电视的图纸很少见 评论 电视的图纸很少见 评论 创维的图纸你要说 版号，不然无能为力 评论 板号5800-p42ALM-0050 168P-P42CLM-01
·日本中文新闻 唐田绘里香为新剧《极恶女王》剃光头 展现演员决心
·日本中文新闻 真子小室夫妇新居引发隐私担忧
·日本中文新闻 前AKB48成员柏木由纪与搞笑艺人交往曝光
·日本学校 ｛日本国际学校｝梅田インターナショナルスクール
·日本学校 LINE：sm287 陳雨菲、20歳、台湾からの留学生、東京に来たばかり
·日本留学生活 出售平成22年走行48000km 代步小车
·日本华人网络交流 円相場 一時1ドル＝140円台まで上昇？
·日本华人网络交流 问日本华人一个问题
·日本旅游代购 富山接机
·生活百科 英国转澳大利亚转换插头
·汽车 【求助】修车遇到困难怎么办？