pytorch的c++/cuda扩展，CUDA编程

xx_xjm

已于 2023-06-12 17:31:21 修改

阅读量1.8k

点赞数 5

分类专栏： pytorch扩展文章标签： pytorch 人工智能 python

于 2023-04-25 18:48:40 首次发布

本文链接：https://blog.csdn.net/xx_xjm/article/details/130367955

版权

pytorch扩展专栏收录该内容

1 篇文章 0 订阅

订阅专栏

1：首先，我们写一个单纯的cpp代码，用于计算两个相同长度的一维数组对应元素之和，代码如下：

2：接下来，我们将上面的cpp代码转化为.cu代码：

3：cuda相关api函数讲解：

a: cudaMalloc

b: cudaFree

c：cudaMemcpy:主机和设备之间的数据传输

一：三大命名空间Aten，c10，torch

二：扩展流程

主要分两部分来讲，一是cuda编程，二是pytorch的c++/cuda扩展；

先说cuda编程

一：什么是cuda编程

我们知道C++，C这类的编程语言是为了让计算机执行我们的指令，确切一点是让计算机的cpu执行我们的执行，现在cuda编程则是要让显卡中的计算核心执行我们的指令；

所以，cuda编程其实就是编写显卡中计算核心执行指令。为了区别于.cpp，.c这样的文件，我们取.cu后缀来指明当前的代码文件是给显卡用的；

二：.cu文件的构成

（这部分基于个人理解，没有严格证明，仅方便理解，如有不正之处，欢迎指正）

因为显卡上面大量的都是计算单元，只有少量的控制单元，我们是无法像C++，C直接指挥CPU那样直接把代码编译成显卡能识别的东西，简单点说，就是，我们不能像C++，C这样写完代码后，编译链接完CPU就可以直接运行了。为此，我们是需要通过CPU转达我们的指令给到显卡；所以说到底，.cu文件也是写给CPU的，只不过你告诉了CPU要怎么去指挥显卡工作，从这个角度来看，其实.cu文件和.cpp等文件一样，都是C++的代码；

因此，.cu文件中包含两部分的内容，一是告诉CPU怎么传递信息和指挥显卡工作，二是指明具体显卡上面应该怎么操作；

三：.cu文件的编译

上面说了，.cu文件包含两部分内容，一部分是写给CPU的，一部分是写给GPU的，所以编译的时候自然需要先区分开这两部分代码，分别编译，然后再合起来一起链接，形成相应的可调用的库（动态库或者静态库），.cu文件的编译需要用到nvcc，这部分详细的可以参考：

(3条消息) CUDA学习（一）-NVCC的编译过程_nvcc编译_Scott f的博客-CSDN博客

nvcc编译网上有很多教程，这一片看不懂的话，大家也可以自查

四：.cu文件的编写

这部分内容主要参考《CUDA编程基础与实践》---清华大学出版社，讲的非常好，强推！

我们前面说了，.cu文件中需要实现cpu对显卡（也就是设备端，后面若无特指，均以设备代替gpu）进行调用，这一步是通过核函数来实现的；

一个典型的，简单的cuda程序结构如下：

注：本文所用代码实例无特殊说明均来自《CUDA编程基础与实践》

int main(void)
{
    主机代码
    核函数调用
    主机代码
    return 0;
}

1：什么是核函数？

我们上面说了，.cu本质上还是写给CPU的，所以核函数其实也是C++函数的一种，只不过有一个特殊的限定词"__global__"，用以指明“这个C++函数，是用来调用显卡的！”。

一个简单的核函数如下，其中__global__和void的顺序可以互换，另外，核函数的返回值必须是空类型：

__global__ void hello_world()
{
    printf("hello world")
}

2：核函数的调用

核函数虽然是C++函数的一种，但是它的调用有一点区别，举例如下：

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

__global__ void hello_world()
{
    printf("hello world");
}

int main(void)
{
    hello_world<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

可以看到，我们这里调用核函数的时候有一个<<<1,1>>>,这其实是指明核函数中的线程数目和排列情况，具体的可以自行百度或者参考书籍，简单理解就是分配多少计算资源来进行运算；

五：主机和设备之间的数据传输

在第四节中，我们简单用主机通过核函数来指挥设备进行工作，但没有进行数据的传输，在这一节，我们更细致的讲解主机和设备间的工作流程；

1：首先，我们写一个单纯的cpp代码，用于计算两个相同长度的一维数组对应元素之和，代码如下：

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

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;

void add(const double *x, const double *y, double *z, const int N);
void check(const double *z, const int N);

int main(void){
    const int N = 100000000;
    const int M = sizeof(double) * N;

    double *x = (double *) malloc(M);
    double *y = (double *) malloc(M);
    double *z = (double *) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        x[n] = a;
        y[n] = b;
    }

    add(x,y,z,N);
    check(z, N);

    free(x);
    free(y);
    free(z);

    return 0;
}

void add(const double *x, const double *y, double *z, const int N)
{
    for (int n=0; n <N; ++n)
    {
        z[n] = x[n] +y[n];
    }
}

void check(const double *z, const int N)
{
    bool has_error = false;
    for (int n = 0; n<N; ++n)
    {
        if (fabs(z[n] - c)> EPSILON)
        {
            has_error = true;
        }
    }
    printf("%s\n", has_error ? "has errors" : "no errors");
}

double *x = (double *) malloc(M)这里;

malloc返回的是一个void类型的指针，也即任意类型的指针，这个指针指向了分配的内存，所以*)malloc(sizeof(double) * 5)就是取出了malloc返回的指针中存储的内存的地址，把这个地址强制转化为double类型，并赋值给p指针，因此，此时的p指针指向的就是所分配的地址了。

接下来，我们将add函数在设备端执行，在这种情况下，一个典型的CUDA程序基本框架为：

头文件包含
常量定义/宏定义
C++自定义函数和CUDA核函数声明（原型）

int main(void)
{
    分配主机和设备内存；
    初始化主机中的数据；
    将某些数据从主机复制到设备；
    调用核函数在设备中进行计算；
    将某些数据从设备复制到主机；
    释放主机与设备内存
}

C++自定义函数和CUDA核函数的定义（实现）

2：接下来，我们将上面的cpp代码转化为.cu代码：

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

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;

void __global__ add(const double *x, const double *y, double *z);
void check(const double *z, const int N);

int main(void){
    const int N = 100000000;
    const int M = sizeof(double) * N;

    double *h_x = (double *) malloc(M);
    double *h_y = (double *) malloc(M);
    double *h_z = (double *) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        h_x[n] = a;
        h_y[n] = b;
    }

    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);

    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

    const int block_size = 128;
    const int grid_size = N / block_size;
    add<<<block_size, grid_size>>>(d_x, d_y, d_z);

    cudaMemcpy(d_z, h_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);

    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);

    return 0;
}

void __global__ add(const double *x, const double *y, double *z, const int N)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    z[n] = x[n] +y[n];

}

void check(const double *z, const int N)
{
    bool has_error = false;
    for (int n = 0; n<N; ++n)
    {
        if (fabs(z[n] - c)> EPSILON)
        {
            has_error = true;
        }
    }
    printf("%s\n", has_error ? "has errors" : "no errors");
}

3：cuda相关api函数讲解：

cuda编程主要用到cuda运行时api函数，所有cuda运行时api函数都是以cuda开头；

a: cudaMalloc

对应于c++中的malloc，用于手动分配设备端的内存，该函数原型如下：

cudaError_t cudaMalloc(void **address, size_t size);

第一个参数address是待分配设备内存的指针，注意，因为内存（地址）本身就是一个指针，所以待分配设备内存的指针是一个指针的指针，即双重指针。

用上面的例子来说明，首先double *d_x, *d_y, *d_z;就是三个指针了，只不过此时d_x，d_y，d_z都还没有指向内容，此时&d_x其实就是取得d_x的地址，(void **)&d_x 相当于(void **) a;此时&d_x就是一个变量，只不过这个变量的名字是一个内存的地址。

那么上面说的，内存地址本身是一个指针怎么理解呢？比如内存地址110，则110此时相当于d_x这样的变量名，它有一个自己的内存地址b，这个内存b的内容指向内存110，也就是内存b存放了内存110的地址，所以说内存地址本身就是一个指针！

所以 *&d_x即取出了d_x这个内存地址所占用的内存地址，也就是上面的内存b，**&d_x则是取出了这个内存b的内容，也就是d_x的地址，所以(void **)&d_x可以直接用&d_x代替！！！

第二个参数size是待分配内存的字节数；

那么这个函数的作用即：让主机中的指针，指向设备端的内存，也即分配设备内存！

b: cudaFree

这个比较简单，对应free，释放设备内存

c：cudaMemcpy:主机和设备之间的数据传输

分配了设备内存以后，就需要把主机上面的数据传送到设备端进行计算；这个函数的原型如下：

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);

第一个参数dst:目标地址

第二个参数src:源地址

第三个参数count:复制数据的字节数

第四个参数标注数据传输方向，它只能取以下五个值：

cudaMemcpyHostToHost，表示从主机复制到主机；

cudaMemcpyHostToDevice，表示从主机复制到设备端；

cudaMemcpyDeviceToDevice，表示从设备复制到设备；

cudaMemcpyDeviceToHost，表示从设备复制到主机；

cudaMemcpyDefault，表示根据dst和src自动判断，这要求系统具有统一虚拟寻址的功能（要求64位的主机）

六：核函数中数据和线程的对应

这部分其实是进行cuda编程的核心，像上面，本来在cpp中的for循环，在核函数中，只用

    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    z[n] = x[n] +y[n];

两行就代替了，原因就是把每个数据对应不同的线程，以执行并行计算，这也是cuda编程加速的原理，这部分很重要，但这里就不详细展开了，感兴趣或者确实需要的还是建议买本书或者查查相应的资料；

七：核函数的要求

1：核函数的返回类型必须是void，在核函数中可以用return关键字，但不可以有返回值；

2：必须使用限定符号__global__，也可以加上一些c++中的其他限定词，如static等；

3：不支持可变数量的参数列表，即参数的个数必须确定；

4：可以向核函数传递非指针变量，其内容对每个线程可见，如上面add的最后参数：

5：除非使用统一内存编程机制，否则传给核函数的数组指针必须指向设备内存，即数据必须放到设备端才能进行计算；

6：核函数不可以称为类成员

7：以往，核函数之间不可以互相调用；但现在应该是都可以了

八：设备函数

核函数是主机调用设备，设备函数，是设备调用设备，

1：__global__修饰的函数称为核函数，一般由主机调用，在设备端执行，但也可以用核函数调用核函数；

2：用__device__修饰的函数称为设备函数，只能被核函数和其他设备函数调用，在设备端执行；

3：__host__修饰的函数就是主机端的普通C++函数，可以用__host__和__device__同时修饰同一个函数，是的该函数既可以作为C++中的普通函数调用，又可以作为设备函数，在设备中调用；

4：不能用__global__和__device__同时修饰同一个函数；

九：数据和线程之间的匹配

实际上线程和数据之间时没啥关系的，这里的匹配是人为设置的，比如我开了<<<grid_size, block_size>>>大小的线程数，此时核函数内部有两个变量是自动生成的，一个是gridDim.x,表示grid_size的大小，也就是有多少个线程块，比如定义了512个线程块，还有一个是blockDim.x，表示的是每一个线程块对的大小，比如每一个线程块有256个线程，所以总的线程就是512x256；

blockIdx.x：表明当前线程所属的线程块的id，比如属于第128块线程块；

threadIdx.x：表明当前线程在所属线程块中的id，比如第56个线程

所以，上面这个线程在总的线程里面的id就是blockIdx.x * blockDim.x + threadIdx.x；也就是128 * 512 + 56，再把这个id当作数组的id，就把线程和数据联系到了一起，所以，数据和线程其实是没关系的。

那么，说道这里有个问题，当线程数小于数据量的时候怎么办呢？

其实也好办，一个线程可以访问多个数据呀，比如总的线程数目是10001，但是现在的数据量是1000000，好了，现在有一个线程id = 1024，那么可以用一个循环来做，即

while(id <1000000):
       执行函数
      id += 10001

每一个线程除了处理当前id，还处理整个线程倍数的数据；

而当线程数大于数据量的时候，可以用

if n >N: return;直接返回，避免多余的线程访问不正确的数据，造成问题，这里n是线程号，N是数据量

ok，第一部分CUDA编程基本到这里结束了~~；

pytorch的c++/cuda扩展

一：三大命名空间Aten，c10，torch

上面，我们已经基本了解了cuda编程是怎么回事了，接下来我们来讲讲怎么把编写好的cuda核函数封装给python调用，最主要的是张量的运算等；

我们在做网络训练的过程中，数据基本都是以张量的形式在GPU上运行，也就是在设备端运行，这有一个好处在于数据已经在设备端了，但页存在一个问题，python中的张量类型数据，在C++中应该怎么接收呢？

实际上，pytorch底层就是由C++实现的，所以C++中肯定有关于张量等数据变量和计算操作函数的实现。C++中有关张量的的实现主要在于三个命名空间：Aten，c10，torch：

Aten：at(ATen)负责声明和定义Tensor运算，是最常用到的命名空间，我们在进行c++扩展的时候，一般都是用到这个命名空间；

c10：c10(Caffe Tensor Library)是 ATen 的基础，包含了PyTorch的核心抽象、Tensor和Storage数据结构的实际实现。这个怎么理解呢，我们知道张量的存储格式是分为头文件和数据文件的，这就是靠c10来实现的，所以本质上pytorch张量也是c++实现的。

torch：torch命名空间下定义的 Tensor 相比于ATen 增加自动求导功能，但 c++ 扩展中一般不常见）

命名空间的运用：命名空间：：变量/函数名，比如at：：Tensor，指的是Aten下的Tensor变量；

二：扩展流程

好了，了解了这些就可以开始扩展代码的编写了，整体的流程可总结如下：

1：核函数/设备函数实现（cu文件中实现）

2：cpp文件中主机调用核函数（但其实在cu文件中实现也可以）

3：通过PYBIND11,将函数接口暴露给python，以供python进行import调用

第一步：核函数实现

主要就是实现所需要的计算核心：一般在.cu文件中实现；

第二步：核函数调用：在一个正常的cpp函数中调用核函数

第三步：将函数通过PYBIND11暴露给python

之后，通过python的setup.py文件进行相关的设置，生成对应的动态库.so文件，就可以进行import了，这部分的话，后面再补充吧，写累了~~~~

xx_xjm

关注

5
点赞
踩
15

收藏

觉得还不错? 一键收藏
0
评论
pytorch的c++/cuda扩展，CUDA编程

我们知道C++，C这类的编程语言是为了让计算机执行我们的指令，确切一点是让计算机的cpu执行我们的执行，现在cuda编程则是要让显卡中的计算核心执行我们的指令；所以，cuda编程其实就是编写显卡中计算核心执行指令。为了区别于.cpp，.c这样的文件，我们取.cu后缀来指明当前的代码文件是给显卡用的；我们上面说了，.cu本质上还是写给CPU的，所以核函数其实也是C++函数的一种，只不过有一个特殊的限定词"__global__"，用以指明“这个C++函数，是用来调用显卡的！
复制链接

扫一扫