《基于CUDA的并行程序设计》学习笔记(三)——下

这里写图片描述


第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.h
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;
}

__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) 弄好了后看到之前选择文件上有一个红色提示符
这里写图片描述

版权声明:本文为manong_wxd原创文章,遵循 CC 4.0 BY-SA 版权协议,转载请附上原文出处链接和本声明。
本文链接:https://blog.csdn.net/manong_wxd/article/details/78633840