上手CUDA编程

news2024/11/18 8:27:23

上手CUDA编程

文章目录

  • 上手CUDA编程
    • 矩阵加法例子
      • 编译
    • 查看本机GPU信息
    • 内存管理函数
      • 专门二维数组拷贝函数
    • Reference
      • >>>>> 欢迎关注公众号【三戒纪元】 <<<<<

矩阵加法例子

编写 CUDA C 程序时, 要将文件命名为 *.cu,cu文件支持 C/C++ 语法。

比对一下同样的矩阵相加,通过CPU和GPU计算,运算效率的差异。

#include <cuda_runtime.h>
#include <stdio.h>
#include "cudastart.h"

// CPU对照组,用于对比加速比
// nx ny 为矩阵长宽
void SumMatrix2DonCPU(float* Mat_A, float* Mat_B, float* Mat_C, int nx,
                      int ny) {
  float* a = Mat_A;
  float* b = Mat_B;
  float* c = Mat_C;
  for (int j = 0; j < ny; j++) {
    for (int i = 0; i < nx; i++) {
      c[i] = a[i] + b[i];
    }
    c += nx;
    b += nx;
    a += nx;
  }
}

// 核函数,每一个线程计算矩阵中的一个元素。
// nx ny 为矩阵长宽
__global__ void SumMatrix(float* Mat_A, float* Mat_B, float* Mat_C, int nx,
                          int ny) {
  int ix = threadIdx.x + blockDim.x * blockIdx.x;
  int iy = threadIdx.y + blockDim.y * blockIdx.y;
  int idx = ix + iy * ny;
  if (ix < nx && iy < ny) {
    Mat_C[idx] = Mat_A[idx] + Mat_B[idx];
  }
}

// 主函数
int main(int argc, char** argv) {
  // 设备初始化
  printf("Randy strating...\n");
  initDevice(0);

  // 输入二维矩阵,4096*4096,单精度浮点型。
  int nx = 1 << 12;
  int ny = 1 << 12;
  int nBytes = nx * ny * sizeof(float);

  // Malloc,开辟主机内存
  float* A_host = (float*)malloc(nBytes);
  float* B_host = (float*)malloc(nBytes);
  float* C_host = (float*)malloc(nBytes);
  float* C_cal_by_gpu = (float*)malloc(nBytes);
  InitialData(A_host, nx * ny);
  InitialData(B_host, nx * ny);

  // cudaMalloc,开辟设备内存
  float* A_dev = NULL;
  float* B_dev = NULL;
  float* C_dev = NULL;
  CHECK(cudaMalloc((void**)&A_dev, nBytes));
  CHECK(cudaMalloc((void**)&B_dev, nBytes));
  CHECK(cudaMalloc((void**)&C_dev, nBytes));

  // 输入数据从主机内存拷贝到设备内存
  CHECK(cudaMemcpy(A_dev, A_host, nBytes, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(B_dev, B_host, nBytes, cudaMemcpyHostToDevice));

  // 二维线程块,32×32
  dim3 block(32, 32);
  // 二维线程网格,128×128
  dim3 grid((nx - 1) / block.x + 1, (ny - 1) / block.y + 1);

  // 测试GPU执行时间
  double GpuStart = CpuSecond();
  // 将核函数放在线程网格中执行
  SumMatrix<<<grid, block>>>(A_dev, B_dev, C_dev, nx, ny);
  CHECK(cudaDeviceSynchronize());
  double gpuTime = CpuSecond() - GpuStart;
  printf("GPU Execution Time: %f sec\n", gpuTime);

  // 在CPU上完成相同的任务
  cudaMemcpy(C_cal_by_gpu, C_dev, nBytes, cudaMemcpyDeviceToHost);
  double cpuStart = CpuSecond();
  SumMatrix2DonCPU(A_host, B_host, C_host, nx, ny);
  double cpuTime = CpuSecond() - cpuStart;
  printf("CPU Execution Time: %f sec\n", cpuTime);

  // 检查GPU与CPU计算结果是否相同
  CHECK(cudaMemcpy(C_cal_by_gpu, C_dev, nBytes, cudaMemcpyDeviceToHost));
  checkResult(C_host, C_cal_by_gpu, nx * ny);

  cudaFree(A_dev);
  cudaFree(B_dev);
  cudaFree(C_dev);
  free(A_host);
  free(B_host);
  free(C_host);
  free(C_cal_by_gpu);
  cudaDeviceReset();
  return 0;
}

具体代码,原博主已经上传到github上了:https://github.com/ZihaoZhao/CUDA_study/tree/master/Sum_Matrix

主要思想就是先对设备进行初始化initDevice(0);,然后初始化主机内存和设备内存。在主机端初始化2个矩阵 A_host、B_host,然后把这2个数组内的值拷贝到设备内存 A_dev、B_dev中。

2维矩阵大小为 4096 × 4096 4096 \times 4096 4096×4096,grid网格大小为 128 × 128 128 \times 128 128×128,线程块大小为 32 × 32 32 \times 32 32×32,总线程为 4096 × 4096 4096 \times 4096 4096×4096,每个线程处理1个元素。

GPU上处理完后,将设备内存结果拷贝到主机,并与CPU结果进行比较。

编译

CUDA代码文件后缀名是.cu。

编译单文件CUDA程序命令

nvcc -o randy sum_martix.cu

之后会生成一个可执行文件test,使用下面的命令即可运行。

./randy

结果:

Randy strating...
Using device 0: NVIDIA GeForce RTX 3070 Laptop GPU
GPU Execution Time: 0.000816 sec
CPU Execution Time: 0.054139 sec
Check result success!

结果也能看出来,GPU的效率是CPU的 $0.054139 \div 0.000816= 66.35倍 $

查看本机GPU信息

#include <cuda.h>
#include <cuda_runtime.h>

#include <iostream>
int main() {
  int dev = 0;
  cudaDeviceProp devProp;
  cudaGetDeviceProperties(&devProp, dev);
  std::cout << "GPU Device Name" << dev << ": " << devProp.name << std::endl;
  std::cout << "SM Count: " << devProp.multiProcessorCount << std::endl;
  std::cout << "Shared Memory Size per Thread Block: "
            << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
  std::cout << "Threads per Thread Block: " << devProp.maxThreadsPerBlock
            << std::endl;
  std::cout << "Threads per SM: " << devProp.maxThreadsPerMultiProcessor
            << std::endl;
  std::cout << "Warps per SM: " << devProp.maxThreadsPerMultiProcessor / 32
            << std::endl;
  return 0;
}

结果:

GPU Device Name0: NVIDIA GeForce RTX 3070 Laptop GPU
SM Count: 40
Shared Memory Size per Thread Block: 48 KB
Threads per Thread Block: 1024
Threads per SM: 1536
Warps per SM: 48

本机电脑使用的是 英伟达的 GeForce RTX 3070显卡,有40个流式多处理器SM,每个线程块共享内存为48KB,每个线程块有1024个线程,每个SM有1536个线程,每个SM有48个线程束。

内存管理函数

主机和设备各自拥有独立的内存,C 语言通过标准库管理主机的内存,CUDA 提供 API 管理设备的内存:

C语言CUDA功能
malloccudaMalloc分配内存
memcpycudaMemcpy复制内存
memsetcudaMemset设置内存
freecudaFree释放内存

主机与设备的数据拷贝,可以使用 cudaMalloccudaMemcpycudaFree 函数,前篇已经介绍过了,看看具体怎么使用

#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#include <sys/time.h>
#include <iostream>
#include <vector>

double CpuSecond() {
  struct timeval tp;
  gettimeofday(&tp, NULL);
  return ((double)tp.tv_sec + (double)tp.tv_usec * 1e-6);
}

int main() {
  float mats[6][4] = {
      {23, 33, 43, 53}, {15, 25, 35, 45}, {12, 22, 32, 42},
      {39, 49, 59, 69}, {20, 30, 40, 50}, {8, 18, 28, 38}
  };
  // copy data to gpu
  std::cout << "mats size: " << sizeof(mats) << std::endl;

  double GpuStart = CpuSecond();
  float *dev_mats;
  cudaError_t err = cudaSuccess;
  err = cudaMalloc((void **)&dev_mats, sizeof(mats));
  if (err != cudaSuccess) {
    printf("Failed to cudaMalloc!");
    return 1;
  }
  cudaMemcpy(dev_mats, mats, sizeof(mats), cudaMemcpyHostToDevice);
  std::cout << "Copied data to GPU.\n";

  // get back copied cuda data
  float host_mats[sizeof(mats) / sizeof(float)];
  cudaMemcpy(&host_mats, dev_mats, sizeof(mats), cudaMemcpyDeviceToHost);
  double gpuTime = CpuSecond() - GpuStart;
  printf("GPU Execution Time: %f sec\n", gpuTime);

  std::cout << "Copied from cuda back to host.\n";
  std::cout << "host_mats size: " << sizeof(host_mats) << std::endl;
  for (int i = 0; i < sizeof(mats) / sizeof(float); i++) {
    std::cout << host_mats[i] << " ";
  }
  std::cout << std::endl;
  cudaFree(dev_mats);

  return 0;
}

结果:

mats size: 96
Copied data to GPU.
GPU Execution Time: 1.244266 sec
Copied from cuda back to host.
host_mats size: 96
23 33 43 53 15 25 35 45 12 22 32 42 39 49 59 69 20 30 40 50 8 18 28 38

专门二维数组拷贝函数

CUDA 避免2维数组在 kernel 运算时损失较高的性能,设计了专用的内存申请函数cudaMallocPitch,设计了cudaMemcpy2D函数在设备间内存拷贝,形参如下

__host__cudaError_t cudaMallocPitch ( void** devPtr, size_t* pitch, size_t width, size_t height )

__host__ cudaError_t cudaMemcpy2D ( void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, cudaMemcpyKind kind )

测试代码:

int main() {

  float mats[6][4] = {{23, 33, 43, 53}, {15, 25, 35, 45}, {12, 22, 32, 42},
                      {39, 49, 59, 69}, {20, 30, 40, 50}, {8, 18, 28, 38}};

  size_t width = 4;
  size_t height = 6;
  size_t pitch;

  std::cout << "mats size: " << sizeof(mats) << std::endl;

  double GpuStart = CpuSecond();
  float *dev_mats;
  cudaError_t err = cudaSuccess;
  err = cudaMallocPitch((void **)&dev_mats, &pitch, sizeof(float) * width,
                        height);
  if (err != cudaSuccess) {
    printf("cudaMalloc failed!");
    return 1;
  }
  // copy data to gpu
  cudaMemcpy2D(dev_mats, pitch, mats, sizeof(float) * width,
               sizeof(float) * width, height, cudaMemcpyHostToDevice);
  std::cout << "Copied data to GPU.\n";

  // get back copied cuda data
  float host_mats[sizeof(mats) / sizeof(float)];
  cudaMemcpy2D(&host_mats, sizeof(float) * width, dev_mats, pitch,
               sizeof(float) * width, height, cudaMemcpyDeviceToHost);
  double gpuTime = CpuSecond() - GpuStart;
  printf("GPU Execution Time: %f sec\n", gpuTime);

  std::cout << "Copied from cuda back to host.\n";
  std::cout << "host_mats size: " << sizeof(host_mats) << std::endl;
  for (int i = 0; i < width * height; i++) {
    std::cout << host_mats[i] << " ";
  }
  std::cout << std::endl;
  cudaFree(dev_mats);

  return 0;
}

结果:

mats size: 96
Copied data to GPU.
GPU Execution Time: 1.251606 sec
Copied from cuda back to host.
host_mats size: 96
23 33 43 53 15 25 35 45 12 22 32 42 39 49 59 69 20 30 40 50 8 18 28 38

这2个函数会使 kernel 的运行时间变短,因为 pitch 对齐后可实现 global 内存联合访问,但对齐也需要额外的时间,整体时间开销反而比不使用2函数时还多。

Reference

  • CUDA编程入门(三)从矩阵加法例程上手CUDA编程
  • CUDA 教程(三)CUDA C 编程简介

>>>>> 欢迎关注公众号【三戒纪元】 <<<<<

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.coloradmin.cn/o/737033.html

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈,一经查实,立即删除!

相关文章

新版危险废物标签二维码制作教程

生态环境部发布的《危险废物识别标志设置技术规范》已经在2023年7月1日正式实施&#xff0c;除了对危废标签格式、内容的规范&#xff0c;明确规定新版危废标签需要包含数字识别码和二维码&#xff0c;实现危险废物“一物一码”管理。 其中危险废物标签中的二维码部分&#xff…

【SuperPoint】语义SLAM深度学习用于特征提取

1. 概况 作者的写作思路很清晰&#xff0c;把各个技术点这么做的原因写的很清楚&#xff0c;一共三篇&#xff0c;另外两篇分别是2016年和2017年发表的&#xff0c;这三篇文章通读下来&#xff0c;可以看清作者在使用深度学习进行位姿估计这一方法上的思路演变过程&#xff0c…

IDEA中 jps+jmap+jconsole命令查看堆内存情况

结论 1.获取进程idjps2.jmap 某个时刻堆内存的情况jdk8之前jmap -heap pid 15876jdk8之后jhsdb jmap --heap --pid 158763.jconsole 动态查看堆内存情况&#xff0c;直接jconsole ,然后弹出可视化窗口jconsole其中12 要结合使用&#xff0c;且是静态的查看&#xff1b;3可以单…

【c++报错】无法打开自己的工程项目(C++ 无法打开文件“xxx.lib”)

问题&#xff1a; C 无法打开文件“xxx.lib” 问题分析&#xff1a; 在进行单个生成的时候&#xff0c;可以生成成功&#xff0c;也可以运行程序。但是点击全部重新生成时&#xff0c;就显示无法打开文件“xxx.lib”。 观察生成顺序&#xff0c;发现exe的程序&#xff08;调用…

基于”Python+”多技术融合在蒸散发与植被总初级生产力估算中的实践

熟悉蒸散发ET及其组分&#xff08;植被蒸腾Ec、土壤蒸发Es、冠层截留Ei&#xff09;、植被总初级生产力GPP的概念和碳水耦合的基本原理&#xff1b;掌握利用Python与ArcGIS工具进行课程相关的操作&#xff1b;熟练掌握国际上流行的Penman-Monteith模型&#xff0c;并能够应用该…

yarn 无法加载文件 CUsersAdministratorAppDataRoamingnpmyarn.ps1,因为在此系统上禁止运行脚本。的解决方案

yarn : 无法加载文件 C:\Users\Administrator\AppData\Roaming\npm\yarn.ps1&#xff0c;因为在此系统上禁止运行脚本。 1、问题描述 执行yarn相关命令时报错&#xff1a; yarn : 无法加载文件 C:\Users\Administrator\AppData\Roaming\npm\yarn.ps1&#xff0c;因为在此系统…

PHP 学生信息管理系统mysql数据库web结构apache计算机软件工程网页wamp

一、源码特点 PHP 学生信息管理系统 是一套完善的web设计系统&#xff0c;对理解php编程开发语言有帮助&#xff0c;系统具有完整的源代码和数据库&#xff0c;系统主要采用B/S模式开发。 代码下载 https://download.csdn.net/download/qq_41221322/88027229https://down…

7 个顶级免费网站在线图像压缩工具!

在将图像上传到网站之前对其进行压缩是缩短网站加载时间的最简单、最有效的方法之一&#xff0c;从而改善访问者的网站体验并提高搜索排名。 大图像会显着降低网站的性能&#xff0c;这总体来说是个坏消息。幸运的是&#xff0c;您可以使用一些很棒的工具来帮助您轻松优化图像…

实现windows系统文件传输到Linux系统中的工具

1、实现windows系统文件传输到Linux系统中的工具 yum -y install lrzsz然后就可以将windows中的文件&#xff0c;直接拖到Xshell窗口即可。

Springboot快速回顾(集成Dubbo)

Dubbo是实现远程调用的一个框架&#xff0c;阿里巴巴开源的。远程调用就是B服务器可以调用A服务器的方法。大型项目会被拆分成多个模块&#xff0c;部署在不同的服务器上。若将公共模块集中部署在一台服务器上&#xff0c;可以方便其他服务器调用。因此&#xff0c;需要Dubbo。…

深度学习实践大全

文章目录 1.可视化调试1.1 各通道相加可视化1.2 降维到3维或2维 1.可视化调试 可视化方法可分为&#xff1a;各通到相加可视化、 1.1 各通道相加可视化 def visualize_feature_map(img_batch,out_path,type,BI):feature_map torch.squeeze(img_batch)feature_map feature_…

C++教程(二)——第一个程序:编写hello world

1、点击左上角【文件】&#xff0c;再点击创建【项目】&#xff0c;设置项目名称&#xff0c;选择存储地址&#xff0c;再应用。 2、首先在解决方案资源管理器中点击【源文件】&#xff0c;右键【添加】--->【新建项】。 3、在弹出窗口中选择C文件(.cpp)&#xff0c;设置名称…

【考研思维题】【哈希表 || 什么时候用哈希表呢?快速查询的时候】【我们一起60天准备考研算法面试(大全)-第九天 9/60】

专注 效率 记忆 预习 笔记 复习 做题 欢迎观看我的博客&#xff0c;如有问题交流&#xff0c;欢迎评论区留言&#xff0c;一定尽快回复&#xff01;&#xff08;大家可以去看我的专栏&#xff0c;是所有文章的目录&#xff09;   文章字体风格&#xff1a; 红色文字表示&#…

【webrtc】ProcessThreadAttached

RegisterModule 调用所在的线程指针传递给ProcessThreadAttached ProcessThreadAttached 调用不是在worker thread 而是在 registers/deregister 这个module或者 start stop 这个module的时候 ** ** pacedsender是一个moudle -实现了

【Zabbix 监控设置】

目录 一、添加 zbx-agent01 客户端主机1、服务端和客户端都配置时间同步2、服务端和客户端都设置 hosts 解析3、设置 zabbix 的下载源&#xff0c;安装 zabbix-agent24、修改 agent2 配置文件5、启动 zabbix-agent26、在服务端验证 zabbix-agent2 的连通性1、常用的键值 7、在 …

Nginx 技术

Nginx (engine x) 是一个高性能的 HTTP 和反向代理 web 服务器&#xff0c;同时也提供了 IMAP/POP3/SMTP 服务。Nginx 是由伊戈尔赛索耶夫为俄罗斯访问量第二的 Rambler.ru 站点开发的&#xff0c;公开版本1.19.6发布于2020年12月15日。 其将源代码以类 BSD 许可证的形式发布&…

论文阅读 | UniFormer

UniFormer: Unified Multi-view Fusion Transformer for Spatial-Temporal Representation in Bird’s-Eye-View 文章目录 UniFormer: Unified Multi-view Fusion Transformer for Spatial-Temporal Representation in Bird’s-Eye-View摘要介绍Question: 说了半天这个时空融合…

开放式耳机漏音严重吗?开放式耳机哪个品牌好?

开放式耳机因其不入耳、不伤耳设计&#xff0c;深受大家喜欢&#xff0c;开放式耳机不会对耳朵产生任何物理伤害&#xff0c;也不会影响听觉神经&#xff0c;所以不会出现传统入耳式耳机音质损伤的问题。相信还有部分用户不知道什么是开放式耳机&#xff1f;开放式漏音严重吗&a…

剑指Offer-29-顺时针打印矩阵

剑指Offer-29题 题目描述&#xff1a;顺时针打印矩阵 输入一个矩阵&#xff0c;按照从外向里以顺时针的顺序依次打印出每一个数字。 **题解思路&#xff1a;**使用 模拟 的方法 定义四个边界变量表示当前要遍历的边界&#xff1a;上(top)、下(bottom)、左(left)、右(right)&am…

爬取新闻评论数据并进行情绪识别

一、为什么要爬取新闻评论数据并进行情绪识别&#xff1f; 爬取新闻评论数据并进行情绪识别的目的是为了从网页中抓取用户对新闻事件或话题的评价内容&#xff0c;并从中识别和提取用户的情绪或态度&#xff0c;如积极、消极、中立等。爬取新闻评论数据并进行情绪识别有以下几个…