文章目录

前言
1. 向量点乘
2. C++ CUDA实现向量点乘
3. 执行结果
总结
学习资料

VS2017 CUDA编程学习1：CUDA编程两变量加法运算
VS2017 CUDA编程学习2：在GPU上执行线程
VS2017 CUDA编程学习3：CUDA获取设备上属性信息
VS2017 CUDA编程学习4：CUDA并行处理初探 - 向量加法实现
VS2017 CUDA编程学习5：CUDA并行执行-线程
VS2017 CUDA编程学习6: GPU存储器架构
VS2017 CUDA编程学习7：线程同步-共享内存
VS2017 CUDA编程学习8：线程同步-原子操作
VS2017 CUDA编程学习9：常量内存
VS2017 CUDA编程学习10：纹理内存

前言

这里继续跟大家分享CUDA编程的知识，这次就不介绍新知识了，主要是应用之前学的CUDA编程知识实现一个小例子：向量点乘运算

1. 向量点乘

这里有两个向量，v1=[a1,a2,...,an]v_{1}=[a_{1}, a_{2},..., a_{n}]v1=[a1,a2,...,an]和v2=[b1,b2,...,bn]v_{2}=[b_{1}, b_{2},..., b_{n}]v2=[b1,b2,...,bn],向量点乘的计算公式如下：
result=∑i=1n(ai∗bi)result=\sum_{i=1}^{n}(a_{i}*b_{i})result=∑i=1n(ai∗bi)

2. C++ CUDA实现向量点乘

这里使用共享内存实现向量点乘，主要是因为共享内存的访问速度比全局内存快一个数量级，使用多块多线程来加速计算过程，在源码中通过注释来解读实现过程。

详细代码如下所示：

#include <iostream>
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>#define N 1024
#define THREADS_PER_BLOCK 512__global__ void gpu_dot(float* d_a, float* d_b, float* d_c)
{//为每个块定义共享内存，不同块的共享内存没法通信__shared__ float partial_sum[THREADS_PER_BLOCK];//整个GPU设备的线程索引int tid = threadIdx.x + blockIdx.x * blockDim.x;//当前块的线程索引int index = threadIdx.x;//每个线程进行一个元素乘法操作，//比如v1=[a1,a2,a3], v2=[b1,b2,b3]//线程t1: a1*b1, 线程t2：a2*b2, ...float sum = 0;while (tid < N){//sum += d_a[tid] * d_b[tid];//多块多线程运行时，这里会造成计算结果出错sum = d_a[tid] * d_b[tid];tid += blockDim.x + gridDim.x;}//将上述部分和保存到共享内存中partial_sum[index] = sum;//线程同步，等待块中所有线程都执行完上述代码行__syncthreads();//使用当前块的共享内存中保存的数据循环计算部分和，//（1）将共享内存数组中数据二分，后面的数据累加到前面的数据中，//（2）然后舍弃后面的数据，在前面的数据中继续（1）的过程int i = blockDim.x / 2;while (i != 0){if (index < i)//每次仅在二分线程中的前半部分线程数进行累加{partial_sum[index] += partial_sum[index + i];}//等待所有线程执行完上述过程__syncthreads();i /= 2;}//将当前块的部分和计算结果保存到全局内存中if (index == 0){d_c[blockIdx.x] = partial_sum[0];}
}int main()
{float* h_a, *h_b, h_c, *partial_sum;float *d_a, *d_b, *d_partial_sum;//计算块的数量和线程数量int block_calc = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;int blocksPerGrid = block_calc > 32 ? 32 : block_calc;//在CPU上分配内存h_a = (float*)malloc(N * sizeof(float));h_b = (float*)malloc(N * sizeof(float));partial_sum = (float*)malloc(blocksPerGrid * sizeof(float));//在GPU上分配内存cudaMalloc(&d_a, N * sizeof(float));cudaMalloc(&d_b, N * sizeof(float));cudaMalloc(&d_partial_sum, blocksPerGrid * sizeof(float));//填充数据for (int i = 0; i < N; i++){h_a[i] = 1;h_b[i] = 2;}//将数据从CPU拷贝到GPU上cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);//调用内核函数，模块数为blocksPerGrid, 每个模块的线程数为THREADS_PER_BLOCKgpu_dot << <blocksPerGrid, THREADS_PER_BLOCK >> > (d_a, d_b, d_partial_sum);//将结果从GPU上拷贝到CPU上cudaMemcpy(partial_sum, d_partial_sum, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost);//计算最终的向量点乘结果：累加每个块的计算结果h_c = 0;for (int i = 0; i < blocksPerGrid; i++){h_c += partial_sum[i];}printf("GPU向量点乘计算结果为：%f\n", h_c);//释放GPU内存cudaFree(d_a);cudaFree(d_b);cudaFree(d_partial_sum);//释放GPU内存free(h_a);free(h_b);free(partial_sum);system("pause");return 0;
}

3. 执行结果

上述例子中两个输入向量一个为全1值，另一个为全2值，向量长度为1024，所以执行结果如下所示：

总结

这里并没有给出GPU上加速效果，大家如果感兴趣可以自行实现在CPU上执行的速度，因本人还是菜鸟一枚，上述过程如有错误，还请大神指正，谢谢！

学习资料

《基于GPU加速的计算机视觉编程》

VS2017 CUDA编程学习实例1：CUDA实现向量点乘相关推荐

VS2017 CUDA编程学习11：CUDA性能测量
文章目录前言 1. CUDA事件API 2. C++ 实现CUDA事件例子 3. CUDA事件例子的执行结果 4. NVIDIA Visual Profiler工具总结学习资料 VS2017 C ...
VS2017 CUDA编程学习4：CUDA并行处理初探 - 向量加法实现
文章目录前言 CUDA实现向量加法 2. 结果展示总结学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2:在GPU上执行线程 VS2017 ...
VS2017 CUDA编程学习5：CUDA并行执行-线程
文章目录前言 1. 线程 2. 代码实现总结学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2:在GPU上执行线程 VS2017 CUDA编 ...
VS2017 CUDA编程学习8：线程同步-原子操作
文章目录前言 1. 原子操作的理解 2. C++ CUDA实现原子操作 3. 执行结果总结学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2 ...
CUDA 编程上手指南：CUDA C 编程及 GPU 基本知识
作者丨科技猛兽编辑丨极市平台本文原创首发于极市平台,转载请获得授权并标明出处. 推荐大家关注极市平台公众号,每天都会更新最新的计算机视觉论文解读.综述盘点.调参攻略.面试经验等干货~ 目录 1 C ...
CUDA 编程简单入门 Advance CUDA 编程基础 (C++ programming)
Advance CUDA编程基础 (C++ programming) GPU 架构 CUDA 编程基础基本代码框架 CUDA Execution Model Case Study : Vector ...
CUDA 编程学习
0 简单的CUDA简介 1.简单教程 CUDA C ++只是使用CUDA创建大规模并行应用程序的方法之一.它允许您使用功能强大的C ++编程语言来开发由GPU上运行的数千个并行线程加速的高性能算法.许 ...
CUDA编程学习0——环境搭建环境详解
目录环境配置软件安装 1.支持最高的cuda版本查询,下载cuda开发软件: 3.配置环境(~/.bashrc添加环境变量) 4.后续维护查询补:关于windows下的cuda环境配置一.Vi ...
CUDA编程学习3——并行计算初窥CUDA的软硬件架构
目录并行计算概述查看GPU相关信息软硬件架构基础物理层(物理结构) 逻辑层(kernel组织) 物理层和逻辑层的总结参考并行计算概述所谓并行计算的概念定义同时多个计算资源一起工作(逻辑 ...
CUDA编程第三章: CUDA执行模型
前言: 本章内容: 通过配置文件驱动的方法优化内核理解线程束执行的本质增大GPU的并行性掌握网格和线程块的启发式配置学习多种CUDA的性能指标和事件了解动态并行与嵌套执行通过上一章的练习, ...

VS2017 CUDA编程学习实例1：CUDA实现向量点乘