深入浅出谈CUDA(二)
矩阵乘法
|
for(j = 0; j < n; j++) {
C[i][j] = 0;
for(k = 0; k < n; k++) {
C[i][j] += A[i][k] * B[k][j];
}
}
}
{
float *a, *b, *c, *d;
int n = 1000;
if(!InitCUDA()) return 0;
a = (float*) malloc(sizeof(float) * n * n);
b = (float*) malloc(sizeof(float) * n * n);
c = (float*) malloc(sizeof(float) * n * n);
d = (float*) malloc(sizeof(float) * n * n);
srand(0);
matgen(a, n, n);
matgen(b, n, n);
clock_t time = matmultCUDA(a, n, b, n, c, n, n);
matmult(a, n, b, n, d, n, n);
compare_mat(c, n, d, n, n);
double sec = (double) time / CLOCKS_PER_SEC;
printf("Time used: %.2f (%.2lf GFLOPS)\n", sec,
2.0 * n * n * n / (sec * 1E9));
return 0;
}
{
int i, j;
for(i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
a[i * lda + j] = (float) rand() / RAND_MAX +
(float) rand() / (RAND_MAX * RAND_MAX);
}
}
}
float* c, int ldc, int n)
{
int i, j, k;
for(i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
double t = 0;
for(k = 0; k < n; k++) {
t += a[i * lda + k] * b[k * ldb + j];
}
c[i * ldc + j] = t;
}
}
}
const float* b, int ldb, int n)
{
float max_err = 0;
float average_err = 0;
int i, j;
for(i = 0; i < n; i++) {
for(j = 0; j < n; j++) {
if(b[i * ldb + j] != 0) {
float err = fabs((a[i * lda + j] -
b[i * ldb + j]) / b[i * ldb + j]);
if(max_err < err) max_err = err;
average_err += err;
}
}
}
printf("Max error: %g Average error: %g\n",
max_err, average_err / (n * n));
}
clock_t matmultCUDA(const float* a, int lda,
const float* b, int ldb, float* c, int ldc, int n)
{
float *ac, *bc, *cc;
clock_t start, end;
start = clock();
cudaMalloc((void**) &ac, sizeof(float) * n * n);
cudaMalloc((void**) &bc, sizeof(float) * n * n);
cudaMalloc((void**) &cc, sizeof(float) * n * n);
cudaMemcpy2D(ac, sizeof(float) * n, a, sizeof(float) * lda,
sizeof(float) * n, n, cudaMemcpyHostToDevice);
cudaMemcpy2D(bc, sizeof(float) * n, b, sizeof(float) * ldb,
sizeof(float) * n, n, cudaMemcpyHostToDevice);
int blocks = (n + NUM_THREADS - 1) / NUM_THREADS;
matMultCUDA<<<blocks * n, NUM_THREADS>>>
(ac, n, bc, n, cc, n, n);
cudaMemcpy2D(c, sizeof(float) * ldc, cc, sizeof(float) * n,
sizeof(float) * n, n, cudaMemcpyDeviceToHost);
cudaFree(ac);
cudaFree(bc);
cudaFree(cc);
end = clock();
return end - start;
}
const float* b, size_t ldb, float* c, size_t ldc, int n)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int idx = bid * blockDim.x + tid;
const int row = idx / n;
const int column = idx % n;
int i;
if(row < n && column < n) {
float t = 0;
for(i = 0; i < n; i++) {
t += a[row * lda + i] * b[i * ldb + column];
}
c[row * ldc + column] = t;
}
}
Time used: 1.1560 (1.73 GFLOPS)
float t = 0;
float y = 0;
for(i = 0; i < n; i++) {
float r;
y -= a[row * lda + i] * b[i * ldb + column];
r = t - y;
y = (r - t) + y;
t = r;
}
}
Time used: 1.1560 (1.73 GFLOPS)
第一个改良
|
(ac, n, bc, n, cc, n, n);
const float* b, size_t ldb, float* c, size_t ldc, int n)
{
extern __shared__ float data[];
const int tid = threadIdx.x;
const int row = blockIdx.x;
int i, j;
for(i = tid; i < n; i += blockDim.x) {
data[i] = a[row * lda + i];
}
__syncthreads();
for(j = tid; j < n; j += blockDim.x) {
float t = 0;
float y = 0;
for(i = 0; i < n; i++) {
float r;
y -= data[i] * b[i * ldb + j];
r = t - y;
y = (r - t) + y;
t = r;
}
c[row * ldc + j] = t;
}
}
第一个部份先把整个 row 读到 shared memory 中,而第二个部份则进行计算,并没有太大的变化。主要的差别是现在一个 row 只由一个 block 进行计算。
Time used: 0.4220 (4.74 GFLOPS)
cudaMallocPitch((void**) &ac, &pitch_a, sizeof(float) * n, n);
cudaMallocPitch((void**) &bc, &pitch_b, sizeof(float) * n, n);
cudaMallocPitch((void**) &cc, &pitch_c, sizeof(float) * n, n);
sizeof(float) * n, n, cudaMemcpyHostToDevice);
cudaMemcpy2D(bc, pitch_b, b, sizeof(float) * ldb,
sizeof(float) * n, n, cudaMemcpyHostToDevice);
(ac, pitch_a / sizeof(float), bc, pitch_b / sizeof(float),
cc, pitch_c / sizeof(float), n);
sizeof(float) * n, n, cudaMemcpyDeviceToHost);
Time used: 0.1250 (16.00 GFLOPS)
进一步的改良
|
+ A(0~15, 32~47) * B(32~47, 0~15) + ...
dim3 blocks(bx, bx);
dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
matMultCUDA<<<blocks, threads>>>(ac, pitch_a / sizeof(float),
bc, pitch_b / sizeof(float), cc, pitch_c / sizeof(float), n);
const float* b, size_t ldb, float* c, size_t ldc, int n)
{
__shared__ float matA[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float matB[BLOCK_SIZE][BLOCK_SIZE];
const int tidc = threadIdx.x;
const int tidr = threadIdx.y;
const int bidc = blockIdx.x * BLOCK_SIZE;
const int bidr = blockIdx.y * BLOCK_SIZE;
int i, j;
float results = 0;
float comp = 0;
for(j = 0; j < n; j += BLOCK_SIZE) {
if(tidr + bidr < n && tidc + j < n) {
matA[tidr][tidc] = a[(tidr + bidr) * lda + tidc + j];
}
else {
matA[tidr][tidc] = 0;
}
if(tidr + j < n && tidc + bidc < n) {
matB[tidr][tidc] = b[(tidr + j) * ldb + tidc + bidc];
}
else {
matB[tidr][tidc] = 0;
}
__syncthreads();
for(i = 0; i < BLOCK_SIZE; i++) {
float t;
comp -= matA[tidr][i] * matB[i][tidc];
t = results - comp;
comp = (t - results) + comp;
results = t;
}
__syncthreads();
}
if(tidr + bidr < n && tidc + bidc < n) {
c[(tidr + bidr) * ldc + tidc + bidc] = results;
}
}
Time used: 0.0780 (25.64 GFLOPS)
cudaMallocPitch((void**) &ac, &pitch_a,
sizeof(float) * newn, newn);
cudaMallocPitch((void**) &bc, &pitch_b,
sizeof(float) * newn, newn);
cudaMallocPitch((void**) &cc, &pitch_c,
sizeof(float) * newn, newn);
cudaMemset(ac, 0, pitch_a * newn);
cudaMemset(bc, 0, pitch_b * newn);
const float* b, size_t ldb, float* c, size_t ldc, int n)
{
__shared__ float matA[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float matB[BLOCK_SIZE][BLOCK_SIZE];
const int tidc = threadIdx.x;
const int tidr = threadIdx.y;
const int bidc = blockIdx.x * BLOCK_SIZE;
const int bidr = blockIdx.y * BLOCK_SIZE;
int i, j;
float results = 0;
float comp = 0;
for(j = 0; j < n; j += BLOCK_SIZE) {
matA[tidr][tidc] = a[(tidr + bidr) * lda + tidc + j];
matB[tidr][tidc] = b[(tidr + j) * ldb + tidc + bidc];
__syncthreads();
for(i = 0; i < BLOCK_SIZE; i++) {
float t;
comp -= matA[tidr][i] * matB[i][tidc];
t = results - comp;
comp = (t - results) + comp;
results = t;
}
__syncthreads();
}
c[(tidr + bidr) * ldc + tidc + bidc] = results;
}
Time used: 0.0780 (25.64 GFLOPS)
结论
|
GPU 的基本介绍
|
![](../p_w_upload/200907/200907151247646676906.jpg)
执行过程
|
Shared memory
|
...
int number = data[16 * tid];
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];
Global memory
|
int number = data[tid];
}
...
__global__ void func(struct vec3d* data, float* output)
{
output[tid] = data[tid].x * data[tid].x +
data[tid].y * data[tid].y +
data[tid].z * data[tid].z;
}
{
output[tid] = x[tid] * x[tid] + y[tid] * y[tid] +
z[tid] * z[tid];
}
{
__shared__ float temp[THREAD_NUM * 3];
const float* fdata = (float*) data;
temp[tid] = fdata[tid];
temp[tid + THREAD_NUM] = fdata[tid + THREAD_NUM];
temp[tid + THREAD_NUM*2] = fdata[tid + THREAD_NUM*2];
__syncthreads();
output[tid] = temp[tid*3] * temp[tid*3] +
temp[tid*3+1] * temp[tid*3+1] +
temp[tid*3+2] * temp[tid*3+2];
}
Texture
|
运算单元
|
和主内存间的数据传输
|
转载于:https://blog.51cto.com/390820/178589
深入浅出谈CUDA(二)相关推荐
- 深入浅出谈cuda 书_年终大促 | 让书做你最好的陪伴
年终大促(11.21--11.30) 全场图书8折优惠 叠加使用文惠券 满400减100满200减50满80减20 ﹀ ﹀ ﹀ 最高可以享受300元的满减优惠! 新书推荐 1.<国王的餐桌&g ...
- 深入浅出谈cuda 书_入门和基础——9本关于美学的书
<朱光潜谈美> 作者:朱光潜 <谈美>是朱光潜先生于1932年写的一本美学入门书,谈了美从哪里来.美是什么以及美的本质等问题.本书从净化人心.美化生活出发,提出人生的艺术化,倡 ...
- 深入浅出谈CUDA(一)
[转载]http://www.pcinlife.com/article/graphics/2008-06-04/1212575164d532.html 第一个CUDA程序 CUDA 目前有两种不同的 ...
- CUDA 深入浅出谈[转]
CUDA 深入浅出谈 "CUDA 是 NVIDIA 的 GPGPU 模型,它使用 C 语言为基础,可以直接以大多数人熟悉的 C 语言,写出在显示芯片上执行的程序,而不需要 ...
- CUDA 深入浅出谈
CUDA 深入浅出谈 "CUDA 是 NVIDIA 的 GPGPU 模型,它使用 C 语言为基础,可以直接以大多数人熟悉的 C 语言,写出在显示芯片上执行的程序,而不需要 ...
- 深入浅出Websocket(二)分布式Websocket集群
前言 最近在构建两个系统的实时通信部分,总结一下所学. 这是一个系列文章,暂时主要构思四个部分 深入浅出Websocket(一)Websocket协议 深入浅出Websocket(二)分布式Webso ...
- 隐私计算--37--演讲实录:深入浅出谈联邦学习
一.前言 前段时间受CSDN邀请,为CSDN和易观分析主办的<隐私计算-Meet-up>做隐私计算相关的演讲,最终选题<深入浅出谈联邦学习>,本次分享的内容主要分为三部分,第一 ...
- C++ 浅谈之二叉搜索树
C++ 浅谈之二叉搜索树 HELLO,各位博友好,我是阿呆
- IT领域标准化浅谈(二):中国IT领域标准制定工作程序
IT领域标准化浅谈(二):中国IT领域标准制定工作程序 我国国家标准制定阶段的划分和流程如下表所示: 阶段代码 阶段名称 阶段任务 阶段成果 完成周期 主要工作 00 预阶段 提出新工作项目建议 P ...
最新文章
- phpmyadmin另类拿shell
- Python 3.x 引入了函数注释
- docker容器的重启策略:通过--restart来指定
- 滤波器开发之四:基于算术平均的中值滤波器
- Linux实战教学笔记22:企业级NFS网络文件共享服务
- 字符串过滤非数字c语言,【新手】【求思路】如何判断用户输入的字符串中是否含有非数字?...
- 华为根本没有鸿蒙系统,【图片】你看不明白的鸿蒙系统,才是华为缔造未来的“伟大”!华为并没有把系统划分为手机操作系统,我们就能知道华为想的并不是那么简单【手机吧】_百度贴吧...
- Redmine(Ruby)配置经验
- native2ascii.exe unicode
- Java函数式编程详解
- 【机器人学】机器人运动学基础
- android 双拼输入法,高效输入解决方案——双拼输入法
- px4 Linux下环境的配置 gcc-arm装不上 OpenGL=3.1 问题
- wps打印错误未定义书签怎么解决_word文档打印时候目录出现错:未定义书签,在wps里怎么修改?...
- 电脑自动出现html文件,电脑自动弹出很多网页怎么办?
- EXCEL 删除重复项并保留最大最小值
- 个人面经整理深信服测开
- 【水汐のc++】建立一个会员管理程序, 每个会员的登记内容包括会员编号、 会员卡号、累计消费金额,可以分别按会员编号、 会员卡号进行查询,也可以增加或删除会员信息。
- 阿里云盘终于可以分享文文件了!!!
- wordpress最佳架构_如何在2019年选择WordPress主题:最佳选择