CUDA实现(A+B)/2

最新推荐文章于 2023-10-23 14:46:07 发布
最新推荐文章于 2023-10-23 14:46:07 发布
阅读量641 收藏 2
点赞数
分类专栏： CUDA 文章标签： cuda
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_29732003/article/details/104219818
版权

作者：老李
日期：2020-2-8

这是我的第一个CUDA程序，也是照着英伟达社区的讲座来写的，写这篇文章的目的是，我希望自己能够尽量把一个程序讲清楚。

Cuda的操作概括来说包含5个步骤：

  1. CPU在GPU上分配内存：cudaMalloc；
  2. CPU把数据发送到GPU：cudaMemcpy；
  3. CPU在GPU上启动内核（kernel），它是自己写的一段程序，在每个线程上运行；
  4. CPU把数据从GPU取回：cudaMemcpy；
  5. CPU释放GPU上的内存。

目标

//target: (A+B)/2 = C

我先把代码贴上：

代码如下

target: (A+B)/2 = C
//2-8
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdlib.h>
#include <stdio.h>
//定义常量
#define N (1024*1024)// size
#define FULL (N*20)
//定义核函数
__global__ void kernel(int* a, int* b, int* c)
{
 int index = threadIdx.x + blockDim.x * blockIdx.x;
 //dim是从1开始的；index是从0开始的(所以不用减1)；index对index
 if (index < N)
 {
  c[index] = (a[index] + b[index]) / 2;
 }
}
int main()
{
 //cuda中了解自己设备的方法
 cudaDeviceProp prop;//查看设备的属性支不支持多流并行的形式
 int whiceDevice;
 cudaGetDevice(&whiceDevice);//查找设备号
 cudaGetDeviceProperties(&prop, whiceDevice);
 if (!prop.deviceOverlap)
 {
  printf("Device will not support overlap\n");
 }
 //初始化计时器
 cudaEvent_t start, stop;
 float elapsedTime;
 //声明流和GPU Buffer的指针
 cudaStream_t stream0;
 cudaStream_t stream1;
 int *host_a, *host_b, *host_c;
 int *dev_a0, *dev_b0, *dev_c0;
 int *dev_a1, *dev_b1, *dev_c1;
 //创建计时器
 cudaEventCreate(&start);
 cudaEventCreate(&stop);
 //初始化流
 cudaStreamCreate(&stream0);
 cudaStreamCreate(&stream1);
 //在GPU端申请存储空间
 cudaMalloc((void**)&dev_a0, N * sizeof(int));
 cudaMalloc((void**)&dev_b0, N * sizeof(int));
 cudaMalloc((void**)&dev_c0, N * sizeof(int));
 cudaMalloc((void**)&dev_a1, N * sizeof(int));
 cudaMalloc((void**)&dev_b1, N * sizeof(int));
 cudaMalloc((void**)&dev_c1, N * sizeof(int));
 //在CPU端申请内存空间，要用到锁页内存(page-locked memory)
 cudaHostAlloc((void**)&host_a, FULL * sizeof(int), cudaHostAllocDefault);
 cudaHostAlloc((void**)&host_b, FULL * sizeof(int), cudaHostAllocDefault);
 cudaHostAlloc((void**)&host_c, FULL * sizeof(int), cudaHostAllocDefault);
 //初始化A,B向量
 for (int i = 0; i < FULL; i++)
 {
  host_a[i] = rand();
  host_b[i] = rand();
 }
 //开始计算
 cudaEventRecord(start, 0);
 for (int i = 0; i < FULL; i += 2 * N)
 {
  //将数据从CPU内存中传输给GPU显存
  cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
  cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
  cudaMemcpyAsync(dev_a1, host_a + i+N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
  cudaMemcpyAsync(dev_b1, host_b + i+N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
  kernel <<<N / 256, 256, 0, stream0 >>>(dev_a0, dev_b0, dev_c0);
  kernel <<<N / 256, 256, 0, stream1 >>>(dev_a1, dev_b1, dev_c1);
  cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
  cudaMemcpyAsync(host_c + i+N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
 }
 cudaStreamSynchronize(stream0);
 cudaStreamSynchronize(stream1);
 cudaEventRecord(stop, 0);
 cudaEventSynchronize(stop);
 cudaEventElapsedTime(&elapsedTime, start, stop);
 printf("time: %3.1f ms \n", elapsedTime);
 cudaFree(dev_a0);
 cudaFree(dev_b0);
 cudaFree(dev_c0);
 cudaFree(dev_a1);
 cudaFree(dev_b1);
 cudaFree(dev_c1);
 cudaFreeHost(host_a);
 cudaFreeHost(host_b);
 cudaFreeHost(host_c);
 cudaStreamDestroy(stream0);
 cudaStreamDestroy(stream1);
 return 0;
}

我一步一步进行描述
声明调用头文件后

1.kernel

  1. 在GPU上执行的函数通常称为核函数。

  2. 一般通过标识符__global__修饰，调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。

  3. 以线程格（Grid）的形式组织，每个线程格由若干个线程块（block）组成，而每个线程块又由若干个线程（thread）组成。

  4. 是以block为单位执行的。

  5. 只能在主机端代码中调用。

  6. 调用时必须声明内核函数的执行参数。

  7. 在编程时，必须先为kernel函数中用到的数组或变量分配好足够的空间，再调用kernel函数，否则在GPU计算时会发生错误，例如越界或报错，甚至导致蓝屏和死机。

我的核函数：

__global__ void kernel(int* a, int* b, int* c)
{
 int index = threadIdx.x + blockDim.x * blockIdx.x;
 //dim是从1开始的；index是从0开始的(所以不用减1)；index对index
 if (index < N)
 {
  c[index] = (a[index] + b[index]) / 2;//实现了目标的功能
 }
}

有两种函数修饰符
1.global，表明被修饰的函数在设备上执行，但在主机上调用。
2.device，表明被修饰的函数在设备上执行，但只能在其他__device__函数或者__global__函数中调用。

然后是对线程索引的计算：
三维结构类型

  1. dim3是基于uint3定义的矢量类型，相当亍由3个unsigned int型组成的结构体。
  2. uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z;
    可使用于一维、二维或三维的索引来标识线程，构成一维、二维或三维线程块。
  3. dim3结构类型变量用在核函数调用的<<<,>>>中。
  4. 相关的几个内置变量
    4.1. threadIdx，顾名思义获取线程thread的ID索引；如果线程是一维的那么就取threadIdx.x，二维的还可以多取到一个值threadIdx.y，以此类推到三维threadIdx.z。
    4.2. blockIdx，线程块的ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。
    4.3. blockDim，线程块的维度，同样有blockDim.x，blockDim.y，blockDim.z。
    4.4. gridDim，线程格的维度，同样有gridDim.x，gridDim.y，gridDim.z。
  5. 对于一维的block，线程的threadID=threadIdx.x。
  6. 对于大小为（blockDim.x, blockDim.y）的 二维block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。
  7. 对于大小为（blockDim.x, blockDim.y, blockDim.z）的 三维 block，线程的threadID=threadIdx.x+threadIdx.yblockDim.x+threadIdx.zblockDim.x*blockDim.y。
  8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride。

一张图可以很好的表示我的核函数的索引值的计算
在这里插入图片描述

2.主函数

我们首先看看自己的设备是否支持并行运算（是否具有重叠功能）

 //cuda中了解自己设备的方法
 cudaDeviceProp prop;//查看设备的属性支不支持多流并行的形式
 int whiceDevice;
 cudaGetDevice(&whiceDevice);//查找设备号
 cudaGetDeviceProperties(&prop, whiceDevice);
 if (!prop.deviceOverlap)
 {
  printf("Device will not support overlap\n");
 }

我需要计算该程序运行的时间（使用cudaEvent的方法）

// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// record start event on the default stream
cudaEventRecord(start);
// execute kernel
kernel<<<grid, block>>>(arguments);
// record stop event on the default stream
cudaEventRecord(stop);
// wait until the stop event completes
cudaEventSynchronize(stop);
// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);
// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);

然后是一些准备的工作，包括创建计时器，申明流，分配空间，初始化变量

 //初始化计时器
 cudaEvent_t start, stop;
 float elapsedTime;
 //声明流和GPU Buffer的指针
 cudaStream_t stream0;
 cudaStream_t stream1;
 int *host_a, *host_b, *host_c;
 int *dev_a0, *dev_b0, *dev_c0;
 int *dev_a1, *dev_b1, *dev_c1;
 //创建计时器
 cudaEventCreate(&start);
 cudaEventCreate(&stop);
 //初始化流
 cudaStreamCreate(&stream0);
 cudaStreamCreate(&stream1);
 //在GPU端申请存储空间
 cudaMalloc((void**)&dev_a0, N * sizeof(int));
 cudaMalloc((void**)&dev_b0, N * sizeof(int));
 cudaMalloc((void**)&dev_c0, N * sizeof(int));
 cudaMalloc((void**)&dev_a1, N * sizeof(int));
 cudaMalloc((void**)&dev_b1, N * sizeof(int));
 cudaMalloc((void**)&dev_c1, N * sizeof(int));
 //在CPU端申请内存空间，要用到锁页内存(page-locked memory)
 cudaHostAlloc((void**)&host_a, FULL * sizeof(int), cudaHostAllocDefault);
 cudaHostAlloc((void**)&host_b, FULL * sizeof(int), cudaHostAllocDefault);
 cudaHostAlloc((void**)&host_c, FULL * sizeof(int), cudaHostAllocDefault);
 //初始化A,B向量
 for (int i = 0; i < FULL; i++)
 {
  host_a[i] = rand();
  host_b[i] = rand();
 }

然后接下来的一步是很重要的，也是信息量最大的一步

//开始计算
 cudaEventRecord(start, 0);
 for (int i = 0; i < FULL; i += 2 * N)
 {
  //将数据从CPU内存中传输给GPU显存
  cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
  cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0);
  cudaMemcpyAsync(dev_a1, host_a + i+N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
  cudaMemcpyAsync(dev_b1, host_b + i+N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
  kernel <<<N / 256, 256, 0, stream0 >>>(dev_a0, dev_b0, dev_c0);
  kernel <<<N / 256, 256, 0, stream1 >>>(dev_a1, dev_b1, dev_c1);
  cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
  cudaMemcpyAsync(host_c + i+N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream1);
 }

在这个循环里我们做了3件事情：

  1. CPU把数据发送到GPU：cudaMemcpy；
  2. CPU在GPU上启动内核（kernel），它是自己写的一段程序，在每个线程上运行；
  3. CPU把数据从GPU取回：cudaMemcpy；

首先是kernel函数的调用方法

一个kernel结构如下：Kernel<<<Dg, Db, Ns, S>>>(param1, param2, …)
-Dg：grid的尺寸，说明一个grid含有多少个block，为dim3类型，一个grid最多含有655356553565535个block，Dg.x，Dg.y，Dg.z最大值为65535；
-Db：block的尺寸，说明一个block含有多上个thread，为dim3类型，一个block最多含有1024(cuda2.x版本)个threads，Db.x和Db.y最大值为1024，Db.z最大值64；
（举个例子，一个block的尺寸可以是：102411 | 25622 | 110241 | 2864 | 4464等）
-Ns：可选参数，如果kernel中由动态分配内存的shared memory，需要在此指定大小，以字节为单位；
-S：可选参数，表示该kernel处在哪个流当中。

然后我们说说数据的传输

这里用的是cudaMemcpyAsync而不是cudaMemcpy，意思是该步骤执行异步操作。

我们引出以下概念：

同步操作：主机向设备提交任务，主机将阻塞，直到设备将所提交任务完成，并将控制权交回主机。然后继续执行主机的程序。

异步操作：主机向设备提交任务，设备直接开始执行任务，但主机将不再阻塞，而是直接继续执行主机的程序。主机并不会等待设备执行任务完毕。

在CUDA当中，核函数kernel的执行总是异步的，而cudaMemcpy数据传输总是同步的。

特别需要注意的是主机在提交核函数之后，不会阻塞等待核函数执行完毕。在profiler CUDA程序时，一定要记得添加cudaDeviceSynchronize() 同步，或者添加一个数据传输(cudaMemcpy-隐含着同步操作) ，以保证核函数执行结束。

也就是说，如果不进行同步，cpu是不知道你的一个步骤是否执行完了，会接着向下执行，这时候如果不进行同步的话，cpu可能会因为缺失数据无法继续执行。

所以才有了后面的流的同步和事件的同步的操作，同时记录运行所耗费的时间

 cudaStreamSynchronize(stream0);
 cudaStreamSynchronize(stream1);
 cudaEventRecord(stop, 0);
 cudaEventSynchronize(stop);
 cudaEventElapsedTime(&elapsedTime, start, stop);
 printf("time: %3.1f ms \n", elapsedTime);

最后释放空间

 cudaFree(dev_a0);
 cudaFree(dev_b0);
 cudaFree(dev_c0);
 cudaFree(dev_a1);
 cudaFree(dev_b1);
 cudaFree(dev_c1);
 cudaFreeHost(host_a);
 cudaFreeHost(host_b);
 cudaFreeHost(host_c);
 cudaStreamDestroy(stream0);
 cudaStreamDestroy(stream1);
 return 0;
}

这是我的第一个cuda程序，水平不足，在理解上可能会有一些错误，如果有些不正确的地方，欢迎大家指正。

大家加油！

老李今天学习了吗
关注
专栏目录
CUDA：矢量加法C=A+B的实例
希望我的博客，能帮上你解决学习中工作中所遇到的问题
12-21 293
CUDA：矢量加法C=A+B的实例
【深度学习环境配置】Anaconda + Pycharm + CUDA + cuDNN + Pytorch + Opencv
shinuone的博客
02-23 5241
（1）【CPU版本】Anaconda（Python） + Pycharm + Pytorch（CPU） + Opencv （2）【GPU版本】Anaconda（Python） + Pycharm + CUDA + cuDNN + Pytorch（GPU） + Opencv
CUDA教程及其实例程序实现
07-29
CUDA教程及其实例程序实现，非常经典的资料，压缩包中有文档和代码实例，非常好的资料
几个cuda源代码示例
09-04
几个cuda源代码示例 1.cuda矩阵例子.cu 2.实现矩阵的乘法.cu 3.... /*****************************************************************************************/ /核心代码部分,实现了矩阵的乘法运算/核心代码部分,实现了矩阵的乘法运算 /*****************************************************************************************/ / // Matrix multiplication kernel –thread specification __global__ void MatrixMulKernel(Matrix M, Matrix N, Matrix P) { // 2D Thread ID int tx = threadIdx.x; int ty = threadIdx.y; // Pvalue is used to store the element of the matrix // that is computed by the thread float Pvalue = 0; // 并行计算体现在次,多个thread并行进行 // 每个thread计算一行与一列的乘积 for (int k = 0; k < M.width; ++k) { float Melement = M.elements[ty * M.pitch + k]; float Nelement = Nd.elements[k * N.pitch + tx]; Pvalue += Melement * Nelement; } // Write the matrix to device memory; // each thread writes one element P.elements[ty * P.pitch + tx] = Pvalue; }
一个基本的CUDA C程序
taifyang的博客
01-01 518
在详细介绍代码之前，有一件事你应该记得，主机代码是由标准C编译器编译的，设备代码是由NVIDIAGPU编译器来执行。基本上，它表示块的数量和将在设备上并行运行的线程数。虽然这不是对设备资源的最佳使用，但是理解在主机上执行的代码和在设备上执行的代码之间的区别是一个很好的起点。_global_是CUDA C在标准C中添加的一个限定符，它告诉编译器在这个限定符后面的函数定义应该在设备上而不是在主机上运行。在前面的代码中，myfirstkernel将运行在设备上而不是主机上，但是,在这段代码中,它是空的。
cuda相关函数
limengjuhanxin的专栏
03-11 1521
内存相关 函数 形参 返回值 意义 __host__ ​ __device__ ​cudaError_t cudaMalloc ( void** devPtr, size_t size ) devPtr:指向分配内存size:需要分配的大小 cudaSuccess = 0:成功cudaErrorInvalidValue = 11:形参不对cudaErrorMemoryAllocat...
基于OpenCV + CPU/GPU的导向滤波的实现
学以致用 知行合一
03-14 692
基于何恺明大神的论文中阐述，引导滤波器在各种计算机视觉和计算机图形应用中既有效又高效，包括降噪、细节平滑/增强（美颜里面的磨皮用导向滤波算法的比较多）、HDR压缩、图像抠图/羽化、雾霾去除和联合上采样。这里就不详述了，可以看论文，网上翻译的版本也非常多。我们着重看下面的代码实现。opencv中给出了导向滤波的cpu版本的函数，可以直接调用，但是对于效率要求非常高的场景，cpu版本不是很适合，会有一些慢。CV_64F格式（主要是里面的boxFilter没有实现，另外有些人的实现是使用blur函数）
CudaJacobi:Jacobi 方法的 CUDA 实现
06-21
Jacobi 方法的 CUDA 实现 该项目是在的范围内创建的，该课程由和。 问题 Jacobi 方法用于为Ax = b形式的线性方程组找到近似数值解。 该算法从x的初始估计开始，并迭代更新它直到收敛。 如果矩阵A是对角占优的，...
【WSL2笔记2】 搭建深度学习开发环境踩坑笔记 Ubuntu+CUDA+cuDNN+PyTorch+Tensorflow+ONNX
老富的博客
04-06 2495
错误的结果就是配置的所有虚拟环境都以base的python版本运行，无法配置每个虚拟环境使用不同python版本，失去了虚拟环境意义。至此，适用PyTorch、Tensorflow、cuDNN框架的深度学习开发环境搭建完成。WSL 上的 CUDA 用户指南。安装protobuf 3.2版本。Anaconda官网版本档案。下载Anaconda安装包。创建Python虚拟环境。设置Anaconda路径。需要注册账号登录下载。通过SSH传送WSL。
PyTorch 学习笔记 (3) PyTorch 1.0+ C++/CUDA extension
风海流的豪赌
11-20 3062
参考教程 PyTorch官方教程 本机系统 Ubuntu 18.04 LTS NVIDIA GeForce GTX 1080 with driver 430.50 CUDA V10.1.243 python 3.6.8 virtualenv torch.version == ‘1.2.0’ Visual Studio Code配置 由于我使用的是python 的virtual environme...
K-means CUDA实现
05-13
CUDA实现K-means算法，学习GPU编程入门级算法需要了解
​自定义卷积注意力算子的CUDA实现​
cyz0202的博客
06-22 1194
​自定义卷积注意力算子的CUDA实现​
cuda()的作用
热门推荐
Vodka_Lou的博客
06-01 2万+
文章目录前言一、cuda()函数的作用二、cuda()的用法1.引入库2.读入数据总结 前言 使用Pytorch框架时，会看到某个变量后面加上如a.cuda()，那么cuda()的作用是什么呢？ 一、cuda()函数的作用 在pytorch中，即使是有GPU的机器，它也不会自动使用GPU，而是需要在程序中显示指定。调用model.cuda()，可以将模型加载到GPU上去。这种方法不被提倡，而建议使用model.to(device)的方式，这样可以显示指定需要使用的计算资源，特别是有多个GPU的情况下。.
CUDA学习（一）
s1c1x1的专栏
04-27 422
//分配host内存 unsigned int numBytes = N * sizeof(float) float* h_A = (float*) malloc(numBytes);
【python】 6 cuda、cudaa、pytorch、tensorflow
孙谨谦的博客
07-27 428
1 机器、深度学习：共性基础笔记
tensorflow+cuda+cudaa对应表
DSCjxgzzd的博客
04-11 1886
https://tensorflow.google.cn/install/source
[CUDA 优化方法与技巧汇总] 源于 周斌 NVIDIA CUDA初级教程视频
xuke_2018的博客
07-25 681
CUDA 优化方法与技巧汇总
CUDA学习笔记（十二） CUDA库简介
最新发布
qq_45788429的博客
10-23 1173
上图是CUDA 库的位置，本文简要介绍cuSPARSE、cuBLAS、cuFFT和cuRAND，之后会介绍OpenACC。CUDA库和CPU编程所用到的库没有什么区别，都是一系列接口的集合，主要好处是，只需要编写host代码，调用相应API即可，可以节约很多开发时间。而且我们完全可以信任这些库能够达到很好的性能，写这些库的人都是在CUDA上的大能，一般人比不了。当然，完全依赖于这些库而对CUDA性能优化一无所知也是不行的，我们依然需要手动做一些改进来挖掘出更好的性能。
CUDA编程
Ling_i的博客
08-17 807
2006年，NVIDIA公司发布了CUDA，CUDA是一种通用的并行计算平台和编程模型，是在C语言基础上扩展的。借助于CUDA，可以像编写C语言程序一样实现并行算法。可以在NVIDIA的GPU平台上用CUDA为多种系统编写应用程序，范围从嵌入式设备、平板电脑、笔记本电脑、台式机、工作站到HPC集群。CUDA编程模型提供了一个计算机架构抽象作为应用程序和其可用硬件之间的桥梁。......
Win10+Qt5.8+CUDA10.0: Ubuntu下Qt Creator CUDA编程教程
可以参考CSDN博主的安装教程，如[博主A](https://blog.csdn.net/threefourly/article/details/84492563)和[博主B](https://blog.csdn.net/HaleyDong/article/details/86093520)的文章获取详细指导。 2. **项目初始...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值