Cuda1.0

新建cuda第一个程序后,对其内置代码进行解析。

1、初始定义。

程序首先定义了一个函数addWithCuda,它是调用GPU运算的入口函数,返回类型是cudaError_t。

cudaError_t是一个枚举类型,可以作为几乎所有CUDA函数的返回类型,用来检测函数执行期间发生的不同类型的错误,一共有80多个错误类型,可以在driver_types.h头文件中查看每一个整型对应的错误类型,如果返回0,代表执行成功。


定义函数

2、核函数解析。

函数addKernel在最前有一个修饰符“__global__”,这个修饰符告诉编译器,被修饰的函数应该编译为在GPU而不是在CPU上运行,所以这个函数将被交给编译设备代码的编译器——NVCC编译器来处理,其他普通的函数或语句将交给主机编译器处理。

这里“设备”的概念可以理解为GPU和其显存组成的运算单元,“主机”可以理解为CPU和系统内存组成的运算单元。在GPU上执行的函数称为核函数


核函数

注:关于threadIdx.x的说明。

CUDA中的线程(thread)是设备中并行运算结构中的最小单位,类似于主机中的线程的概念,thread可以以一维、二维、三维的形式组织在一起,threadIdx.x表示的是thread在x方向的索引号,还可能存在thread在y和z方向的索引号threadIdx.y和threadIdx.z。

一维、二维或三维的thread组成一个线程块(Block),一维、二维或三维的线程块(Block)组合成一个线程块网格(Grid),线程块网格(Grid)可以是一维或二维的。通过网格块(Grid)->线程块(Block)->线程(thread)的 顺序可以定位到每一个并且唯一的线程。

addKernel函数会被GPU上的多个线程同时执行一次,线程间彼此没有通信,相互独立。到底会有多少个线程来分别执行核函数,是在“<<< >>>”符号里定义的。“<<< >>>”表示运行时配置符号,在本程序中的定义是<<<1,size>>>,表示分配了一个线程块(Block),每个线程块有分配了size个线程“<<<>>>”中的 参数并不是传递给设备代码的参数,而是定义主机代码运行时如何启动设备代码。以上定义的这些线程都是一个维度上的,可以通过thredaIdx.x来获取执行当前计算任务的线程的ID号。

3、其他函数解析。

cudaSetDevice函数用来设置要在哪个GPU上执行,如果只有一个GPU,设置为cudaSetDevice(0);

cudaMalloc函数用来为参与运算的数据分配显存空间,函数原型:cudaError_t cudaMalloc(void **p, size_t s);

cudaMemcpy函数用于主机内存和设备显存以及主机与主机之间,设备与设备之间相互拷贝数据,函数原型:

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

第一个参数dst是目标数据地址,第二个参数src是源数据地址,第三个参数count是数据大小,第四个参数kind定义数据拷贝的类型,有如下几类枚举类型:

enum __device_builtin__ cudaMemcpyKind  

{  

cudaMemcpyHostToHost          =   0,/**< Host   -> Host */  

cudaMemcpyHostToDevice        =   1,/**< Host   -> Device */  

cudaMemcpyDeviceToHost        =   2,/**< Device -> Host */  

cudaMemcpyDeviceToDevice      =   3,/**< Device -> Device */  

cudaMemcpyDefault             =   4/**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */  

};

cudaGetLastError函数用于返回最新的一个运行时调用错误,对于任何CUDA错误,都可以通过函数cudaGetErrorString函数来获取错误的详细信息。

cudaDeviceSynchronize函数提供了一个阻塞,用于等待所有的线程都执行完各自的计算任务,然后继续往下执行。

cudaFree函数用于释放申请的显存空间。

cudaDeviceReset函数用于释放所有申请的显存空间和重置设备状态;

©著作权归作者所有,转载或内容合作请联系作者
  • 序言:七十年代末,一起剥皮案震惊了整个滨河市,随后出现的几起案子,更是在滨河造成了极大的恐慌,老刑警刘岩,带你破解...
    沈念sama阅读 212,686评论 6 492
  • 序言:滨河连续发生了三起死亡事件,死亡现场离奇诡异,居然都是意外死亡,警方通过查阅死者的电脑和手机,发现死者居然都...
    沈念sama阅读 90,668评论 3 385
  • 文/潘晓璐 我一进店门,熙熙楼的掌柜王于贵愁眉苦脸地迎上来,“玉大人,你说我怎么就摊上这事。” “怎么了?”我有些...
    开封第一讲书人阅读 158,160评论 0 348
  • 文/不坏的土叔 我叫张陵,是天一观的道长。 经常有香客问我,道长,这世上最难降的妖魔是什么? 我笑而不...
    开封第一讲书人阅读 56,736评论 1 284
  • 正文 为了忘掉前任,我火速办了婚礼,结果婚礼上,老公的妹妹穿的比我还像新娘。我一直安慰自己,他们只是感情好,可当我...
    茶点故事阅读 65,847评论 6 386
  • 文/花漫 我一把揭开白布。 她就那样静静地躺着,像睡着了一般。 火红的嫁衣衬着肌肤如雪。 梳的纹丝不乱的头发上,一...
    开封第一讲书人阅读 50,043评论 1 291
  • 那天,我揣着相机与录音,去河边找鬼。 笑死,一个胖子当着我的面吹牛,可吹牛的内容都是我干的。 我是一名探鬼主播,决...
    沈念sama阅读 39,129评论 3 410
  • 文/苍兰香墨 我猛地睁开眼,长吁一口气:“原来是场噩梦啊……” “哼!你这毒妇竟也来了?” 一声冷哼从身侧响起,我...
    开封第一讲书人阅读 37,872评论 0 268
  • 序言:老挝万荣一对情侣失踪,失踪者是张志新(化名)和其女友刘颖,没想到半个月后,有当地人在树林里发现了一具尸体,经...
    沈念sama阅读 44,318评论 1 303
  • 正文 独居荒郊野岭守林人离奇死亡,尸身上长有42处带血的脓包…… 初始之章·张勋 以下内容为张勋视角 年9月15日...
    茶点故事阅读 36,645评论 2 327
  • 正文 我和宋清朗相恋三年,在试婚纱的时候发现自己被绿了。 大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
    茶点故事阅读 38,777评论 1 341
  • 序言:一个原本活蹦乱跳的男人离奇死亡,死状恐怖,灵堂内的尸体忽然破棺而出,到底是诈尸还是另有隐情,我是刑警宁泽,带...
    沈念sama阅读 34,470评论 4 333
  • 正文 年R本政府宣布,位于F岛的核电站,受9级特大地震影响,放射性物质发生泄漏。R本人自食恶果不足惜,却给世界环境...
    茶点故事阅读 40,126评论 3 317
  • 文/蒙蒙 一、第九天 我趴在偏房一处隐蔽的房顶上张望。 院中可真热闹,春花似锦、人声如沸。这庄子的主人今日做“春日...
    开封第一讲书人阅读 30,861评论 0 21
  • 文/苍兰香墨 我抬头看了看天上的太阳。三九已至,却和暖如春,着一层夹袄步出监牢的瞬间,已是汗流浃背。 一阵脚步声响...
    开封第一讲书人阅读 32,095评论 1 267
  • 我被黑心中介骗来泰国打工, 没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留,地道东北人。 一个月前我还...
    沈念sama阅读 46,589评论 2 362
  • 正文 我出身青楼,却偏偏与公主长得像,于是被迫代替她去往敌国和亲。 传闻我的和亲对象是个残疾皇子,可洞房花烛夜当晚...
    茶点故事阅读 43,687评论 2 351

推荐阅读更多精彩内容

  • CUDA从入门到精通(零):写在前面 本文原版链接: 在老板的要求下,本博主从2012年上高性能计算课程开始接触C...
    Pitfalls阅读 3,604评论 1 3
  • 1. CPU vs. GPU 1.1 四种计算机模型 GPU设计的初衷就是为了减轻CPU计算的负载,将一部分图形计...
    王侦阅读 20,859评论 3 20
  • CUDA是一种新的操作GPU计算的硬件和软件架构,它将GPU视作一个数据并行计算设备,而且无需把这些计算映射到图形...
    ai领域阅读 9,059评论 0 8
  • CUDA是什么 CUDA,ComputeUnifiedDeviceArchitecture的简称,是由NVIDIA...
    Pitfalls阅读 9,474评论 0 1
  • 这是可羧为你献上的第一部影评 你的喜欢是我最高的目标 看不见的缘分 ...
    可羧阅读 3,193评论 13 32