开发者社区> mingchaosun> 正文
阿里云
为了无法计算的价值
打开APP
阿里云APP内打开

CUDA编程(二) CUDA初始化与核函数

简介:
+关注继续查看

CUDA编程(二)

CUDA初始化与核函数

CUDA初始化

在上一次中已经说过了,CUDA安装成功之后,新建一个工程还是十分简单的,直接在新建项目的时候选择NVIDIA CUDA项目就可以了,我们先新建一个MyCudaTest 工程,删掉自带的示例kernel.cu,然后新建项,新建一个CUDA C/C++ File ,我们首先看一下如何初始化CUDA,因此我命名为InitCuda.cu

这里写图片描述

这里写图片描述

首先我们要使用CUDA的RunTime API 所以 我们需要include cuda_runtime.h

#include <stdio.h> 

//CUDA RunTime API
#include <cuda_runtime.h>

接下来这个函数会调用 runtime API 中 有关初始化CUDA的内容

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的装置的数目
    cudaGetDeviceCount(&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) {
            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;
}

这段程序首先会调用cudaGetDeviceCount 函数,获得支持 CUDA 的GPU的数量,如果计算机上没有支持 CUDA 的装置,则会传回 1,而这个1是device 0 ,device0 只是一个仿真装置,但是CUDA的很多功能都不支持(不支持CUDA1.0以上版本),因此我们要真正确定系统上是否有支持CUDA的装置,需要对每个device调用cudaGetDeviceProperties,来获得它们的具体参数,以及所支持的CUDA版本(prop.major 和 prop.minor 分别代表装置支持的版本号码,例如 6.5 则 prop.major 为 6 而prop.minor 为 5)

cudaGetDeviceProperties除了可以获得装置支持的 CUDA 版本之外,还有装置的名称、内存的大小、最大的 thread 数目、执行单元的频率等等。详情可参考NVIDIA 的 CUDA Programming Guide。

在找到支持 CUDA 1.0 以上的装置之后,就可以呼叫 cudaSetDevice 函式,把它设为目前要使用的显卡。

下面我们在Main函数中调用InitCUDA函数,由于我们使用VS,所以直接ctrl+F5编译执行就可以了,执行时如果系统上有支持 CUDA 的装置,应该会显示 CUDA initialized。

int main() 
{

    if (!InitCUDA()) 
    { 
        return 0; 
    }

    printf("CUDA initialized.\n"); return 0;

}

这里写图片描述

完整程序:

#include <stdio.h> 

//CUDA RunTime API
#include <cuda_runtime.h>

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的装置的数目
    cudaGetDeviceCount(&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) {
            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;
}

int main() 
{

    if (!InitCUDA()) 
    { 
        return 0; 
    }

    printf("CUDA initialized.\n"); return 0;

}

CUDA核函数

完成了CUDA的初始化检查操作,下面我们就可以使用CUDA完成一些简单计算了,这里我们打算计算一系列数字的立方和。

所以我们先写了一个随机函数:

#define DATA_SIZE 1048576

int data[DATA_SIZE];

//产生大量0-9之间的随机数
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//生成随机数(main中调用)
//GenerateNumbers(data, DATA_SIZE);

该函数会产生一大堆 0 ~ 9 之间的随机数,然后我们要对他们进行立方和操作。

那么我们如何让这个工作在显卡上完成呢?首先第一件事很显而易见,这些数字不能放在内存里了,而是要复制到GPU的显存上。下面我们就来看一下数据复制的部分。

Host&Device架构:

这里写图片描述

上一次已经讲过关于CUDA架构的一些基础了,这里再稍微复习一下,在 CUDA 的架构下,一个程序分为两个部份:host 端和 device 端。Host 端是指在 CPU 上执行的部份,而 device 端则是在显示芯片上执行的部份。Device 端的程序又称为 “kernel”。通常 host 端程序会将数据准备好后,复制到显卡的内存中,再由显示芯片执行 device 端程序,完成后再由 host 端程序将结果从显卡的内存中取回。

我们需要把产生的数据复制到Device端的RAM,才能在显卡上完成计算,因此我们首先开辟一块合适的显存,然后把随机数从内存复制进去。

    //生成随机数
    GenerateNumbers(data, DATA_SIZE);

    /*把数据复制到显卡内存中*/

    int* gpudata, *result;

    //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int));

    //cudaMemcpy 将产生的随机数复制到显卡内存中 
    //cudaMemcpyHostToDevice - 从内存复制到显卡内存
    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE,cudaMemcpyHostToDevice);

注释已经写得比较明白了。cudaMalloc 和 cudaMemcpy 的用法和一般的 malloc 及 memcpy 类似,不过 cudaMemcpy 则多出一个参数,指示复制内存的方向。在这里因为是从主内存复制到显卡内存,所以使用 cudaMemcpyHostToDevice。如果是从显卡内存到主内存,则使用cudaMemcpyDeviceToHost。

完成了从内存到显存的数据拷贝之后,我们接下来就要在显卡上完成计算了,如何让程序跑在显卡上?答案是核函数。

CUDA核函数:

要写在显示芯片上执行的程序。在 CUDA 中,在函数前面加上__global__ 表示这个函式是要在显示芯片上执行的,所以我们只要在正常函数之前加上一个__global__就行了:

// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
    int sum = 0;

    int i;

    for (i = 0; i < DATA_SIZE; i++) {

        sum += num[i] * num[i] * num[i];

    }

    *result = sum;

}

在显示芯片上执行的程序有一些限制,首先最明显的一个限制——不能有传回值,还有一些其他的限制,后面会慢慢提到。

执行核函数:

写好核函数之后需要让CUDA执行这个函数。

在 CUDA 中,要执行一个核函数,使用以下的语法:

    函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);

这里我们先不去并行,只是单纯地完成GPU计算,所以我们让block = 1,thread = 1,share memory = 0

    sumOfSquares<<<1, 1, 0>>>(gpudata, result);

计算完了,千万别忘了还要把结果从显示芯片复制回主内存上,然后释放掉内存~

    int sum;

    //cudaMemcpy 将结果从显存中复制回内存
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);

最后我们把结果打印出来就大功告成了:

    printf("GPUsum: %d \n", sum);

之后我们再用CPU计算一下来验证一下上面的过程是否有错,这一步还是十分必要的:

    sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i] * data[i];
    }

    printf("CPUsum: %d \n", sum);

完整程序:

程序代码:

#include <stdio.h>
#include <stdlib.h>

//CUDA RunTime API
#include <cuda_runtime.h>

#define DATA_SIZE 1048576

int data[DATA_SIZE];

//产生大量0-9之间的随机数
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的装置的数目
    cudaGetDeviceCount(&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) {
            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__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
    int sum = 0;

    int i;

    for (i = 0; i < DATA_SIZE; i++) {

        sum += num[i] * num[i] * num[i];

    }

    *result = sum;

}





int main()
{

    //CUDA 初始化
    if (!InitCUDA()) {
        return 0;
    }

    //生成随机数
    GenerateNumbers(data, DATA_SIZE);

    /*把数据复制到显卡内存中*/

    int* gpudata, *result;

    //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int));

    //cudaMemcpy 将产生的随机数复制到显卡内存中 
    //cudaMemcpyHostToDevice - 从内存复制到显卡内存
    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);

    // 在CUDA 中执行函数 语法:函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    sumOfSquares << <1, 1, 0 >> >(gpudata, result);


    /*把结果从显示芯片复制回主内存*/

    int sum;

    //cudaMemcpy 将结果从显存中复制回内存
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);

    printf("GPUsum: %d \n", sum);

    sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i] * data[i];
    }

    printf("CPUsum: %d \n", sum);

    return 0;
}

运行结果:

这里写图片描述

总结:

这次给大家介绍了CUDA的初始化和如何在显卡上运行程序,即先将数据从内存复制到显存,再写好运算的核函数,之后用CUDA调用核函数,完成GPU上的计算,之后当然不要忘记将结果复制回内存,释放掉显存。

总的来说一个CUDA程序的骨架已经搭建起来了,而GPU计算的重中之重即并行加速还没有进行介绍,不过在加速之前我们还有一件非常重要的事情需要考虑,那就是我们的程序到底有没有加速,也就是我们要输出程序的运行时间,这个时间我们需要使用CUDA提供的一个Clock函数,可以取得GPU执行单元的频率,所以下一篇博客我将主要讲解这个函数~希望能给大家的学习带来帮助~

参考资料:《深入浅出谈CUDA》

版权声明:本文内容由阿里云实名注册用户自发贡献,版权归原作者所有,阿里云开发者社区不拥有其著作权,亦不承担相应法律责任。具体规则请查看《阿里云开发者社区用户服务协议》和《阿里云开发者社区知识产权保护指引》。如果您发现本社区中有涉嫌抄袭的内容,填写侵权投诉表单进行举报,一经查实,本社区将立刻删除涉嫌侵权内容。

相关文章
iOS的App模块化编程的框架实现方案——BeeHive
BeeHive是用于iOS的App模块化编程的框架实现方案,吸收了Spring框架Service的理念来实现模块间的API耦合。
319 0
Java初学者作业——编写 Java 程序,让用户输入指定数字实现产生随机数。
Java初学者作业——编写 Java 程序,让用户输入指定数字实现产生随机数。
28 0
Java初学者作业——编写程序计算实发工资(实践1)
Java初学者作业——编写程序计算实发工资(实践1)
53 0
自己动手编写IOC框架(二)
  万事开头难,上篇已经起了一个头,之后的事情相对就简单了。上次定义了框架所需的dtd也就是规定了xml中该怎么写,有哪些元素。并且我们也让dtd和xml绑定在了一起,使dtd对xml的格式进行校验,并且在DocumentHolder中提供了根据xml文件路径获取xml文件的Document对象。
623 0
+关注
77
文章
0
问答
文章排行榜
最热
最新
相关电子书
更多
低代码开发师(初级)实战教程
立即下载
阿里巴巴DevOps 最佳实践手册
立即下载
冬季实战营第三期:MySQL数据库进阶实战
立即下载