《基于CUDA的并行程序设计》学习笔记(三)——下
来源:互联网 发布:卧龙大数据 上市 编辑:程序博客网 时间:2024/05/22 03:24
第3章 CUDA编程基础
3.5 “HelloWorld”CUDA 编程实例
安装完Visual Studio 2013软件并配置好CUDA开发环境。本节我们正式CUDA编程。
首先我们打开vs2013,点击新建工程,选中CUDA 7.5的模板。输入工程名等信息,完成工程的创建。
创建完工程后,会默认有一个kernel.cu的文件,其实现的是矩阵相加,这时候可以直接调试运行,如果没有报错则证明可以进行CUDA编程了。
可能会出现,下面这种无法识别<<<
符号的情况,但是目测不影响结果的输出。
运行结果如下:
如果想编写运行自己的程序,则需要先移除kernel.cu文件,不然会出现main函数重复声明。
我们现在开始编写自己的CUDA程序,首先移除原来的kernel.cu文件,然后右键工程名,选择添加新item。
选择CUDA C/C++ File
,然后输入文件名为main,就可以得到一个空的main.cu文件。在文件中输入如下代码:
#include <stdio.h>#include <cuda_runtime.h>#include <string>// 要使用 runtime API 的时候,需要 include cuda_runtime.hbool InitCUDA(){ int count; // 获取计算能力>=1.0的设备数量 cudaGetDeviceCount(&count); printf("%d\n", count); if (count == 0) { fprintf(stderr, "There is no device./n"); return false; } int i; for (i = 0; i < count; i++) { // 指定设备的属性 cudaDeviceProp prop; if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) { // 输出设备名称 char *c = prop.name; printf("%s\n",c); // 定义设备计算能力的主要修订编号 printf("%d",prop.major); // 定义设备计算能力的次要修订编号 printf("%d\n", prop.minor); if (prop.major >= 1) { break; } } } if (i == count) { fprintf(stderr, "There is no device supporting CUDA 1.x./n"); return false; } // 把设备设置为调用主线程的当前设备 cudaSetDevice(i); return true;}__global__ void kernel(char *str1, char *str2){ while ((*str1) != '\0') *str2++ = *str1++; *str2 = '\0';}int main(){ if (!InitCUDA()) { return 0; } char* host_str1; host_str1 = "HelloWorld"; char* host_str2 = new char[11]; int size = strlen(host_str1) + 1; char* dev_str1; char* dev_str2; cudaMalloc((void**)&dev_str1, size); cudaMalloc((void**)&dev_str2, size); cudaMemcpy(dev_str1, host_str1, size, cudaMemcpyHostToDevice); kernel <<<1, 1 >>>(dev_str1, dev_str2); cudaMemcpy(host_str2, dev_str2, size, cudaMemcpyDeviceToHost); printf("%s\n",host_str2); cudaFree(dev_str1); cudaFree(dev_str2); getchar(); return 0;}
InitCUDA()
中会先呼叫 cudaGetDeviceCount 函式,取得支持 CUDA 的装置的数目。如果系统上没有支持 CUDA 的装置,则它会传回 1,而 device 0 会是一个仿真的装置,但不支持 CUDA 1.0 以上的功能。所以,要确定系统上是否有支持 CUDA 的装置,需要对每个 device 呼叫 cudaGetDeviceProperties 函式,取得装置的各项数据,并判断装置支持的 CUDA 版本(prop.major 和 prop.minor 分别代表装置支持的版本号码,例如 1.0 则 prop.major 为 1 而 prop.minor 为 0)。
透过 cudaGetDeviceProperties 函式可以取得许多数据,除了装置支持的 CUDA 版本之外,还有装置的名称、内存的大小、最大的 thread 数目、执行单元的频率等等。详情可参考 NVIDIA 的 CUDA Programming Guide。
在找到支持 CUDA 1.0 以上的装置之后,就可以呼叫 cudaSetDevice 函式,把它设为目前要使用的装置。
编译后输出如下:
能正确输出上述结果,则代表第一个完整的CUDA程序到此结束了。但是这个项目的结构看上去一点也不规范化,所有的代码都写到一个源文件里,这样的代码不易读懂,也不好定位程序的错误。CUDA程序一般可以进行如下图所示的文件结构管理。
如上图所示,编写CUDA程序过程中,若某个功能适合串行,则编写串行代码在CPU端执行,若适合并行,则编写CUDA并行代码在GPU端执行。CUDA并行代码的机构应进行如下分离:
(1) CUDA程序调用接口封装了和CUDA相关的函数,一般位于.cpp文件中。
(2) CUDA主机端代码的主要功能包括选择计算设备,进行GPU端存储器的分配、主机端与设备端直接的数据复制及调用kernel函数等准备工作,一般位于一个独立的.cu文件中。
(3) CUDA设备端代码主要指GPU端执行的核心代码kernel函数。kernel函数可能有多个,各个kernel函数各自放在一个单独的以kernel函数功能命名的.cu文件下,因此kernel函数相关的文件可能有多个。
根据以上代码管理的方法,对HelloWordl程序改写如下:在源文件创建3个文件(main.cpp、kernel.cu和GPU_HelloWorld.cu)。main.cpp中调用核函数GPU_HelloWorld(host_str1, host_str2)
相当于上述结构中的“GPU函数”,kernel.cu文件相当于上述结构中的“kernel函数”。
main.cpp文件主要包含了一个完整程序执行的主框架,其代码改写如下:
#include<stdio.h>extern bool GPU_HelloWorld(char* host_str1, char* host_str2);int main(){ char* host_str1; host_str1 = "HelloWorld"; char* host_str2 = new char[11]; if(!GPU_HelloWorld(host_str1, host_str2)) return 0; printf("%s\n",host_str2); return 0;}
kernel.cu文件为GPU端执行函数的核心代码,其代码如下:
#include <cuda_runtime.h>__global__ void kernel(char *str1, char *str2){ while ((*str1) != '\0') *str2++ = *str1++; *str2 = '\0';}
GPU_HelloWorld.cu文件作用是为kernel函数选择可用的计算设备并进行调用环境准备,其代码如下:
#include "kernel.cu"#include <cuda_runtime.h>#include <string>bool InitCUDA(){ int count; // 获取计算能力>=1.0的设备数量 cudaGetDeviceCount(&count); printf("%d\n", count); if (count == 0) { fprintf(stderr, "There is no device./n"); return false; } int i; for (i = 0; i < count; i++) { // 指定设备的属性 cudaDeviceProp prop; if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) { // 输出设备名称 char *c = prop.name; printf("%s\n",c); // 定义设备计算能力的主要修订编号 printf("%d",prop.major); // 定义设备计算能力的次要修订编号 printf("%d\n", prop.minor); if (prop.major >= 1) { break; } } } if (i == count) { fprintf(stderr, "There is no device supporting CUDA 1.x./n"); return false; } // 把设备设置为调用主线程的当前设备 cudaSetDevice(i); return true;}bool GPU_HelloWorld(char* host_str1, char* host_str2){ if(!InitCUDA()) { return 0; } int size = strlen(host_str1)+1; char* dev_str1; char* dev_str2; cudaMalloc((void**)&dev_str1, size); cudaMalloc((void**)&dev_str2, size); cudaMemcpy(dev_str1, host_str1, size, cudaMemcpyHostToDevice); kernel <<<1, 1 >>>(dev_str1, dev_str2); cudaMemcpy(host_str2, dev_str2, size, cudaMemcpyDeviceToHost); cudaFree(dev_str1); cudaFree(dev_str2); return 1;}
编译运行结构修改后的项目,将得到和修改前一样的运行结果。编译时需要将kernel.cu文件从项目生成中排除,否则编译连接时会出错。
注:如何从生成项目中移除
(1) 右键想要移除的文件,选择属性
(2) 点击右侧的从从生成中移除的yes
,然后点击确定
(3) 弄好了后看到之前选择文件上有一个红色提示符
- 《基于CUDA的并行程序设计》学习笔记(三)——下
- 《基于CUDA的并行程序设计》学习笔记(三)——上
- 《基于CUDA的并行程序设计》学习笔记(三)——中
- 《基于CUDA的并行程序设计》学习笔记(二)
- 《基于CUDA的并行程序设计》学习笔记(一)
- 《基于CUDA的并行程序设计》阅读笔记(一)
- Cuda学习笔记(三)——Cuda编程Tips
- 并行程序设计---cuda memory
- CUDA学习笔记三
- CUDA学习笔记三
- cuda学习笔记之异步并行执行
- cuda学习笔记之异步并行执行
- cuda《学习笔记三》——共享内存和同步
- 《多核程序设计》学习笔记:冒泡排序的并行实现
- cuda并行结构下的冒泡排序
- CUDA学习--矩阵乘法的并行运算
- NVIDIA CUDA 6.5 基于CentOS 6.5 x86_64 开发平台的建立——GPU并行计算前提
- 基于CUDA的并行lammps编译及测试
- 节日_100
- 深度学习概述
- 基础作业篇
- 91. Decode Ways
- 持续集成-git使用
- 《基于CUDA的并行程序设计》学习笔记(三)——下
- OJ(Online Judge)系统及ACM测试题库大全
- 第一篇博客 仅做测试
- win10 VS2017 安装OpenSSL
- 我想进阿里
- 编写程序求500 以内的勾股弦数,即满足 c2=b2+a2的3个数,要求b>a。将所有符合要求的组合存入文本文件中,每个组合占一行。
- 设计模式—原型模式
- Tensorflow项目构建流程
- Android aapt 工具环境配置