码迷，mamicode.com
首页 > 其他好文 > 详细

CUDA编程（二） CUDA初始化与核函数

时间：2016-03-30 20:58:34      阅读：315      评论：0      收藏：0      [点我收藏+]

标签：

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》

CUDA编程（二） CUDA初始化与核函数

标签：

原文地址：http://blog.csdn.net/sunmc1204953974/article/details/51016816
踩
(0)
赞
(0)
   
举报
评论 一句话评论（0）
登录后才能评论！
分享档案
更多>
2021年07月29日 (22)
2021年07月28日 (40)
2021年07月27日 (32)
2021年07月26日 (79)
2021年07月23日 (29)
2021年07月22日 (30)
2021年07月21日 (42)
2021年07月20日 (16)
2021年07月19日 (90)
2021年07月16日 (35)
周排行
mamicode.com排行更多图片
更多
友情链接
兰亭集智   国之画   百度统计   站长统计   阿里云   chrome插件   新版天听网
关于我们 - 联系我们 - 留言反馈
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码！