CUDA学习笔记(二)CUDA简介

本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/,仅用于学习。

CUDA是并行计算的平台和类C编程模型,我们能很容易的实现并行算法,就像写C代码一样。只要配备的NVIDIA GPU,就可以在许多设备上运行你的并行程序,无论是台式机、笔记本抑或平板电脑。熟悉C语言可以帮助你尽快掌握CUDA。

CUDA编程

CUDA编程允许你的程序执行在异构系统上,即CUP和GPU,二者有各自的存储空间,并由PCI-Express 总线区分开。因此,我们应该先注意二者术语上的区分:

  • Host:CPU and itsmemory (host memory)
  • Device: GPU and its memory (device memory)

代码中,一般用h_前缀表示host memory,d_表示device memory。

kernel是CUDA编程中的关键,他是跑在GPU的代码,用标示符__global__注明。

host可以独立于host进行大部分操作。当一个kernel启动后,控制权会立刻返还给CPU来执行其他额外的任务。所以,CUDA编程是异步的。一个典型的CUDA程序包含由并行代码补足的串行代码,串行代码由host执行,并行代码在device中执行。host端代码是标准C,device是CUDA C代码。我们可以把所有代码放到一个单独的源文件,也可以使用多个文件或库。NVIDIA C编译器(nvcc)可以编译host和device生成可执行程序。

这里再次说明下CUDA程序的处理流程:

  1. 从CPU拷贝数据到GPU。
  2. 调用kernel来操作存储在GPU的数据。
  3. 将操作结果从GPU拷贝至CPU。

Memory操作

cuda程序将系统区分成host和device,二者有各自的memory。kernel可以操作device memory,为了能很好的控制device端内存,CUDA提供了几个内存操作函数:

为了保证和易于学习,CUDA C 的风格跟C很接近,比如:

cudaError_t cudaMalloc ( void** devPtr, size_t size )

我们主要看看cudaMencpy,其函数原型为:

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )

其中cudaMemcpykind的可选类型有:

  1. cudaMemcpyHostToHost
  2. cudaMemcpyHossToDevice
  3. cudaMemcpyDeviceToHost
  4. cudaMemcpuDeviceToDevice

具体含义很好懂,就不多做解释了。

对于返回类型cudaError_t,如果正确调用,则返回cudaSuccess,否则返回cudaErrorMemoryAllocation。可以使用char* cudaGetErrorString(cudaError_t error)将其转化为易于理解的格式。

组织线程

掌握如何组织线程是CUDA编程的重要部分。CUDA线程分成Grid和Block两个层次。

由一个单独的kernel启动的所有线程组成一个grid,grid中所有线程共享global memory。一个grid由许多block组成,block由许多线程组成,grid和block都可以是一维二维或者三维,上图是一个二维grid和二维block。

这里介绍几个CUDA内置变量:

  • blockIdx:block的索引,blockIdx.x表示block的x坐标。
  • threadIdx:线程索引,同理blockIdx。
  • blockDim:block维度,上图中blockDim.x=5.
  • gridDim:grid维度,同理blockDim。

一般会把grid组织成2D,block为3D。grid和block都使用dim3作为声明,例如:

dim3 block(3);
// 后续博文会解释为何这样写grid
dim3 grid((nElem+block.x-1)/block.x);

需要注意的是,dim3仅为host端可见,其对应的device端类型为uint3。 

启动CUDA kernel

CUDA kernel的调用格式为:

kernel_name<<<grid, block>>>(argument list);

其中grid和block即为上文中介绍的类型为dim3的变量。通过这两个变量可以配置一个kernel的线程总和,以及线程的组织形式。例如:

kernel_name<<<4, 8>>>(argumentt list);

该行代码表明有grid为一维,有4个block,block为一维,每个block有8个线程,故此共有4*8=32个线程。

注意,不同于c函数的调用,所有CUDA kernel的启动都是异步的,当CUDA kernel被调用时,控制权会立即返回给CPU。

函数类型标示符

__device__ 和__host__可以组合使用。 

kernel的限制:

  • 仅能获取device memory 。
  • 必须返回void类型。
  • 不支持可变数目参数。
  • 不支持静态变量。
  • 不支持函数指针。
  • 异步。

代码分析

#include <cuda_runtime.h>
#include <stdio.h>
#define CHECK(call) \
{ \const cudaError_t error = call; \if (error != cudaSuccess) \{ \printf("Error: %s:%d, ", __FILE__, __LINE__); \printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \exit(1); \} \
}
void checkResult(float *hostRef, float *gpuRef, const int N) {double epsilon = 1.0E-8;bool match = 1;for (int i=0; i<N; i++) {if (abs(hostRef[i] - gpuRef[i]) > epsilon) {match = 0;printf("Arrays do not match!\n");printf("host %5.2f gpu %5.2f at current %d\n",hostRef[i],gpuRef[i],i);break;}}if (match) printf("Arrays match.\n\n");
}
void initialData(float *ip,int size) {// generate different seed for random numbertime_t t;srand((unsigned) time(&t));for (int i=0; i<size; i++) {ip[i] = (float)( rand() & 0xFF )/10.0f;}
}
void sumArraysOnHost(float *A, float *B, float *C, const int N) {for (int idx=0; idx<N; idx++)C[idx] = A[idx] + B[idx];
}
__global__ void sumArraysOnGPU(float *A, float *B, float *C) {int i = threadIdx.x;C[i] = A[i] + B[i];
}
int main(int argc, char **argv) {printf("%s Starting...\n", argv[0]);// set up deviceint dev = 0;cudaSetDevice(dev);// set up data size of vectorsint nElem = 32;printf("Vector size %d\n", nElem);// malloc host memorysize_t nBytes = nElem * sizeof(float);float *h_A, *h_B, *hostRef, *gpuRef;h_A = (float *)malloc(nBytes);h_B = (float *)malloc(nBytes);hostRef = (float *)malloc(nBytes);gpuRef = (float *)malloc(nBytes);// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef, 0, nBytes);// malloc device global memoryfloat *d_A, *d_B, *d_C;cudaMalloc((float**)&d_A, nBytes);cudaMalloc((float**)&d_B, nBytes);cudaMalloc((float**)&d_C, nBytes);// transfer data from host to devicecudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);// invoke kernel at host sidedim3 block (nElem);dim3 grid (nElem/block.x);sumArraysOnGPU<<< grid, block >>>(d_A, d_B, d_C);printf("Execution configuration <<<%d, %d>>>\n",grid.x,block.x);// copy kernel result back to host sidecudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);// add vector at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// check device resultscheckResult(hostRef, gpuRef, nElem);// free device global memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// free host memoryfree(h_A);free(h_B);free(hostRef);free(gpuRef);return(0);
}

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

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

相关文章

什么是马尔科夫随机场?

马尔科夫随机场&#xff0c;也称为马尔可夫网&#xff08;Markov Network&#xff09;&#xff0c;是一种概率图模型&#xff0c;用于表示随机变量之间的依赖关系。它是由若干个随机变量组成的无向图&#xff0c;其中节点代表随机变量&#xff0c;边代表它们之间的相互作用或依…

Node.js的crypto模块 加密

Node.js的crypto模块提供了许多加密和解密功能&#xff0c;包括对称加密、非对称加密、哈希函数等。在本篇文章中&#xff0c;我们将详细介绍Node.js的crypto模块的API、代码注释和举例。 加密和解密 对称加密 对称加密算法使用相同的密钥进行加密和解密&#xff0c;例如AES…

Android 13.0 进入recovery模式(等待用户选择recovery模式界面)进入自动恢复出厂设置模式

1.概述 在13.0的系统产品开发中,由于产品硬件有按钮,按钮执行恢复出厂设置功能,需要实现自动恢复出厂设置的功能,这就需要去掉等待输入recovery模式的相关代码,改成默认恢复出厂模式就实现这个功能了 2.进入recovery模式(等待用户选择recovery模式界面)进入自动恢复出厂设…

浅谈uniapp中开发安卓原生插件

其实官方文档介绍的比较清楚而且详细,但是有时候他太墨迹,你一下子找不到自己想要的,所以我总结了一下开发的提纲,也是为了自己方便下次使用。 1.第一步,下载官方提供的Android的示例工程,然后倒入UniPlugin-Hello-AS工程请在App离线SDK中查找,之后Android studio,编译运行项目…

自编efi文件测试vmware虚拟机如何进入UEFI环境

同事突然让帮忙编一下UEFI&#xff0c;之前完全没有接触过&#xff0c;在此粗鲁记录其过程。 UEFI的开源框架是edk2&#xff0c;开发环境配置起来还是有些麻烦&#xff0c;完全按照文档编译不过&#xff0c;经人帮助总算编译通过&#xff0c;但如何测试又是问题&#xff1b;网…

【T+】畅捷通T+增加会计科目提示执行超时已过期。

【问题描述】 在畅捷通T软件中&#xff0c; 增加会计科目的时候提示&#xff1a; 通过DataTable插入ext扩展表出错:执行超时已过期。 完成操作之前已超时或服务器未响应。 操作已被用户取消。 语句已终止。 【解决方法】 【方法一】 注销用户登录&#xff0c;回到软件登录界面…

FFmpeg和rtsp服务器搭建视频直播流服务

下面使用的是ubuntu的&#xff0c;window系统可以参考&#xff1a; 通过rtsp-simple-server和ffmpeg实现录屏并发布视频直播_rtsp simple server_病毒宇宇的博客-CSDN博客 一、安装rtsp-simple-server &#xff08;1&#xff09;下载rtsp-simple-server 下载地址&#xff1a;R…

1024啊啊啊啊啊啊

1024 程序员节快乐&#xff0c;没什么想发的&#xff0c;只是想要个1024胸章。

Kotlin中的Lambda表达式基本定义和使用

在Kotlin中&#xff0c;Lambda表达式是一种简洁的方式来定义匿名函数。Lambda表达式可以作为函数的实际参数或者返回值&#xff0c;使得函数成为高阶函数。本篇博客将介绍Lambda表达式的基本概念以及使用方法&#xff0c;并提供相关的示例代码。 Lambda表达式的基本概念 Lamb…

vsCode 格式化配置

学习目标&#xff1a; 基于 vsCode 配置格式化工具&#xff0c;提高&#xff08;React、Vue &#xff09;开发效率  1. vsCode 安装 prettier 插件并启用  2. 修改配置文件 setting.json setting.json 位置&#xff1a; 依次点击 替换内容&#xff1a;↓ {"git.enab…

智加科技与东风柳汽达成深度合作 自动驾驶重卡计划2024年初量产交付

&#xff08;2023年10月19日&#xff0c;苏州&#xff09;全球领先的重卡自动驾驶技术公司智加科技与东风柳汽宣布&#xff0c;双方共同开发的自动驾驶重卡H7计划2024年初实现量产交付。未来&#xff0c;双方将携手推出安全可靠、高性价比、性能卓越的自动驾驶重卡产品&#xf…

什么年代了,还在用FastQC?试试Falco吧

什么年代了&#xff0c;还在用FastQC&#xff1f;试试Falco吧 目前大部分的教程在质控上都是推荐的FastQC&#xff0c;然而它有一个不足&#xff0c;就是虽然名字上有一个Fast&#xff0c;但是它还不够Fast&#xff0c;真正的快&#xff0c;还得是Falco。 如何安装&#xff1…

STM32 HAL高级定时器正交编码模式案例

STM32 HAL高级定时器正交编码模式案例 &#x1f516;基于stm32F030RBT6单片机采用高级定时器1&#xff0c;编码器模式&#xff0c;测试EC11编码器。 &#x1f3ac;EC11测试效果&#xff1a; &#x1f33f;STM32定时器编码器有3种映射模式: ✨本次采用的是上面的模式3&#x…

postgresql14-模式的管理(三)

基本概念 postgresql成为数据库管理系统DBMS&#xff0c;在内存中以进程的形态运行起来成为一个实例&#xff0c;可管理多个database。 数据库databases&#xff1a;包含表、索引、视图、存储过程&#xff1b; 模式schema&#xff1a;多个对象组成一个模式&#xff0c;多个模…

2023年10月22日找工作面试交流遇到的基本问题

交叉编译解决的痛点问题 不同硬件体系结构之间的编译问题。嵌入式系统开发需要在主机上编写代码。提高效率和节省时间。软件移植和管理依赖关系。 不同硬件体系结构之间的编译问题&#xff1a;例如&#xff0c;你开发了一个针对Intel x86架构的应用程序&#xff0c;但想要在Ra…

学成在线第一天-课程内容管理服务搭建以及查询课程接口设计

目录 一、搭建课程内容管理服务 二、设计接口 三、面试题 四、总结 一、搭建课程内容管理服务 没什么好说的&#xff0c;直接就是创建内容模块 然后这个继承父模块&#xff0c;然后再课程内容模块下面创建三个子模块&#xff0c;model、sevice、controller model依赖base…

分享一下微信小程序的文章中怎么添加营销活动

在数字化时代&#xff0c;小程序已经成为企业营销的重要工具。通过小程序&#xff0c;企业可以提供更加便捷、高效的服务&#xff0c;吸引更多的用户和客户。本文将以小程序营销活动为主题&#xff0c;介绍如何在小程序文章中加入营销活动&#xff0c;提高品牌知名度和销售额。…

告别SQL优化和数据迁移备份烦恼,NineData--小白也能成为DBA的秘密武器!

1、概述 说起sql优化&#xff0c;大家可能首先想到的是创建索引、避免全表扫描、减少子查询及优化查询语句、避免隐式类型转换、慢查询日志记录慢SQL、explain分析SQL的执行计划及调优等等。当然&#xff0c;这也是我们常使用的一些调优手段&#xff0c;而且对开发人员也有一定…

【C语言】用函数实现模块化程序设计

前言&#xff1a;如果把所有的程序代码都写在一个主函数(main函数)中&#xff0c;就会使主函数变得庞杂、头绪不清&#xff0c;使阅读和维护程序变得困难。此外&#xff0c;有时程序中要多次实现某一功能&#xff0c;如果重新编写实现此功能就会使得程序冗长、不精炼。 &#x…