并行计算之路<3>——CUDA与CPP文件联姻

最新推荐文章于 2023-11-14 19:53:21 发布

xx116213

最新推荐文章于 2023-11-14 19:53:21 发布

阅读量1.6k

点赞数

分类专栏：并行计算文章标签： cuda gpgpu cpp

本文链接：https://blog.csdn.net/xx116213/article/details/50674855

版权

并行计算专栏收录该内容

5 篇文章 2 订阅

订阅专栏

原内容来源于《GPGPU编程技术——从GLSL、CUDA到OpenCL》的7.1.3节。 有修改。

目的

之前都是单一的*.cu文件。一个较为复杂的工程，往往需要多个文件来实现，合理安排这些文件的结构会使得工程文件尽然有序。《GPGPU编程技术——从GLSL、CUDA到OpenCL》提供了三种方法来实现将CUDA C代码集成到C++中。，

内核代理：将CUDA C程序集成在面向对象的CPP项目中。在CPP语言的帮助下，CUDA C的程序模块也可以呈现面向对象的特性。使用内核代理的基本原理如下：

把CUDA C代码从CPP程序中提取出来，并用文件分割两者，使CPP中类的成员函数“看不到”CUDA C代码的存在。被提取出来的CUDA C代码被一些代理函数封装。于是，代理函数调用CUDA内核，而同时CPP函数调用代理函数。这些代理函数并不包含任何对功能的实现，只是起到了将调用重定向的作用，这样一来，CUDA内核和CPP代码就可以被隔离开。这样做的目的是，尽量封装CUDA C代码，并用nvcc来编译它们，而用CPP编译剩余的CPP代码。

可能画个图会比较好理解。

架构

1）C++应用程序：application.cpp

#include <iostream>

#include "class.cuh"

using namespace std;

int main(int argc, char **argv)
{
    CHelloWorld* hello = new CHelloWorld(argc, argv);
    hello->sayHello();

    delete hello;

    return 0;
}

2）类的定义：class.cuh

#ifndef CLASS_CUH
#define CLASS_CUH

#include <cstdlib>
#include <iostream>

#include "kernel.cuh"

#define BLOCKSIZE 8u

using namespace std;

class CHelloWorld {
public:
    CHelloWorld(int argc, char **argv) {
        _argc = argc;
        _argv = argv;
        init();
    }

    ~CHelloWorld();

    void sayHello(void);

private:
    void init(void);
    int _argc;
    char** _argv;
    unsigned _nBlockSize;
};


#endif

3）类的实现：class.cu

#include "class.cuh"
#include "kernel.cuh"

void CHelloWorld::init()
{
    _nBlockSize = (unsigned)BLOCKSIZE;

    //TODO:...
}

CHelloWorld::~CHelloWorld()
{
    //TODO:...
}

void CHelloWorld::sayHello(void)
{
    sayHello_agent(_nBlockSize);
}

4）内核头文件：kernel.cuh

#ifndef KERNEL_CUH
#define KERNEL_CUH


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

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

__global__ void sayHello_kernel(void);

void sayHello_agent(unsigned unBlockSize);

#endif

5）内核实现：kernel.cu

#include "kernel.cuh"

__global__ void sayHello_kernel(void)
{
    printf("Hello from thread %d\n", threadIdx.x);
}


void sayHello_agent(unsigned unBlockSize)
{
    dim3 dimBlock;
    dim3 dimGrid;

    dimBlock.x = unBlockSize;
    dimBlock.y = 1;
    dimBlock.z = 1;
    dimGrid.x = 1;
    dimGrid.y = 1;
    dimGrid.z = 1;

    // 代理函数调用内核
    sayHello_kernel <<< dimGrid, dimBlock >>>();
    cudaThreadExit();
}

输出结果如下：

Hello from thread 0
Hello from thread 1
Hello from thread 2
Hello from thread 3
Hello from thread 4
Hello from thread 5
Hello from thread 6
Hello from thread 7