樟木头网站仿做,优化大师人工服务电话,孟村做网站,网站建设上传服务器步骤文章目录 前言一、使用的工具二、学习资源分享三、libtorch环境配置1.配置CUDA、nvcc、cudnn2.下载libtorch3.CLion配置libtorch4.CMake Application指定Environment variables5.测试libtorch 四、PyTorch CUDA CPP项目流程1.使用CLion结合torch extension编写可以调用cuda的C代… 文章目录 前言一、使用的工具二、学习资源分享三、libtorch环境配置1.配置CUDA、nvcc、cudnn2.下载libtorch3.CLion配置libtorch4.CMake Application指定Environment variables5.测试libtorch 四、PyTorch CUDA CPP项目流程1.使用CLion结合torch extension编写可以调用cuda的C代码2.编写CU设备代码在代码中调用CUDA核函数使用GPU运算3.创建header定义设备函数4.编写CPP主机代码在代码中使用PYBIND库连接pytorch和c代码使用主机函数调用设备函数5.在Pycharm中创建setup.py脚本将C编写的cuda代码编译为python的一个package6.使用pip安装自定义的cuda代码库7.安装好自定义的cuda自定义库之后编写python代码调用 五、总结 前言
这学期确定了研究方法具体为三维重建相关转而研究三维重建相关的知识。最近3D Gaussian Splatting方法效果十分的好并且开源了源代码因此十分值得对其源码进行研究源码中对于可微光栅化的实现是基于CUDA实现的因此想要后续对这块内容进行改进则必须了解CUDA程序是如何编写的作为一名DL农工学习CUDA编程也有助于自己对于整个DL流程的理解也能够进一步拓展自己的编程技术因此无论出于什么目的作为DL农工我认为有必要掌握这项技术。
本教程仅作为Pytorch CUDA CPP开发的简易教程我也是初学者。同时本教程的目的也是为了方便不熟悉CUDA项目结构的人快速了解一份新的程序的文件结构是怎样的。
一、使用的工具
Pytorch, libtorch, CUDA, C, CLion, Pycharm
二、学习资源分享
如果想要详细了解整个CUDA程序的工作原理可以观看这个详细教程CUDA编程基础入门系列
如果想功利点的学习如何使用Pytorch调用CUDA程序可以观看这个教程Pytorchcpp/cuda extension 教学
本简易教程也是基于第二个视频教程写的这个教程讲了一种通用的Pytorch、CUDA程序编写范式使用pytorchcppcuda开发的几乎都遵循这套代码编写流程因此学会怎么写基本上就能看明白其他人的CUDA程序文件是如何组织的尽管你还是有可能看不懂CUDA代码但是你学完后一定能看明白CUDA程序是如何组织的。
自定义 C 和 CUDA 扩展 PYTORCH C API
三、libtorch环境配置
注意我所有的操作都是在Windows上进行的Linux上的操作应该也大差不差 libtorch环境配置的教程有很多你可以参考其他人的博客进行配置这里我仅贴我使用CLion配置libtorch的流程
1.配置CUDA、nvcc、cudnn
该项不进行说明有许多攻略可供参考这里贴一篇博客 Windows安装CUDA及cuDNN【保姆级教程】 注意安装的CUDA版本不要太新了且要与libtorch提供的CUDA版本保持一致目前官网提供的libtorch CUDA版本为11.8因此你安装的CUDA最好也是11.8否则会有许多奇怪的错误 你都来看CUDA编程了不至于DL环境还不会配置吧
2.下载libtorch
前往pytorch官网下载libtorchPytorch 你可以选择CPU和CUDA版本这就和Pytorch的CPU和CUDA版是一个意思个人感觉对于Pytorch CUDA开发CPU和CUDA版没有什么区别libtorch就相当于一个C版本的Pytorch但是我们的目的仅仅是用C版的torch调用CUDA程序最后还是要用python代码调用C代码最后由C调用CUDA程序。为了避免出现奇怪的bug还是一开始就下载CUDA版的吧。
有release和debug两个版本可以选择debug版支持debug操作相应容量会大些release版不能debug容量会小很多。你自己选择一般选debug版。
3.CLion配置libtorch
创建CLion项目一定要创建 C Executable 项目不要创建CUDA Executable否则CMAKE时会不通过出现奇怪的BUG我也不知道为啥。
在CMakeLists.txt中填入一下内容 有#标识的表示需要修改成你自己的libtorch里对应的路径
cmake_minimum_required(VERSION 3.25)
project(xxx) # xxx为你自己的项目名称这部分为CLion自动生成的set(CMAKE_CXX_STANDARD 17)
set(Torch_DIR E:/CLion/libtorch/share/cmake/Torch) # your own path
find_package(Torch REQUIRED)include_directories(E:/CLion/libtorch/include) # your own path
include_directories(E:/CLion/libtorch/include/torch/csrc/api/include) # your own pathadd_executable(xxx main.cpp)
target_link_libraries(xxx ${TORCH_LIBRARIES})
set_property(TARGET xxx PROPERTY CXX_STANDARD 17)4.CMake Application指定Environment variables
经过实测如果不指定这个选项会出现CMake时找不到xxx.dll这些dll文件其实是存在的但就是找不到位置因此需要手动去指定参考我的这篇博客CLion配置libtorch找不到xxx.dll
5.测试libtorch
在main.cpp粘贴如下代码
#include iostream
#include torch/torch.hint main() {torch::Tensor tensor torch::rand({2, 3});std::cout tensor std::endl;std::cout torch::cuda::is_available() std::endl;std::cout Hello, World! std::endl;return 0;
}出现上面的输出libtorch配置成功
如果在CMAKE时出现Unicode相关字样的报错请参考我的这篇博客warning C4819: 该文件包含不能在当前代码页(936)中表示的字符。请将该文件保存为 Unicode 格式以防止数据丢失
四、PyTorch CUDA CPP项目流程 可以把C理解为沟通Pytorch代码和CUDA代码的桥梁
1.使用CLion结合torch extension编写可以调用cuda的C代码 创建interpolation.cpp你可以理解为该文件内存放着主机函数即在CPU上执行的函数用于调用设备函数即在GPU上执行的函数。
#include torch/extension.h
#include include/utils.h // 自定义的headertorch::Tensor trilinear_interpolation_fw(torch::Tensor feats, torch::Tensor points){/*使用torch库定义cuda函数* 相当于在主机函数中调用设备函数代码* 用于执行前向传播操作* */// 只要是传进来的参数是Tensor类型都要进行CHECK_INPUT// 如果传入的参数是一般的intfloat等则不需要CHECK_INPUTCHECK_INPUT(feats);CHECK_INPUT(points);// trilinear_fw_cu() 为设备函数// 可以理解为主机函数作为设备函数的入口这只是一种写法你也可以有别的写法return trilinear_fw_cu(feats, points);
}torch::Tensor trilinear_interpolation_bw(torch::Tensor dL_dfeat_interp, torch::Tensor feats, torch::Tensor points){/*使用torch库定义cuda函数* 相当于在主机函数中调用设备函数代码* 用于执行反向传播操作计算梯度* */// 只要是传进来的参数是Tensor类型都要进行CHECK_INPUT// 如果传入的参数是一般的intfloat等则不需要CHECK_INPUTCHECK_INPUT(dL_dfeat_interp);CHECK_INPUT(feats);CHECK_INPUT(points);return trilinear_bw_cu(dL_dfeat_interp, feats, points);
}PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {/*pybind库用于作为连接python和C的桥梁* name_: 表示python调用c函数的名称(名称可一致也可以不一致)f表示c函数的地址*/m.def(trilinear_interpolation_fw, trilinear_interpolation_fw);m.def(trilinear_interpolation_bw, trilinear_interpolation_bw);
}2.编写CU设备代码在代码中调用CUDA核函数使用GPU运算
注意设备函数卸载 .cu 文件中只有这样在编译时nvcc才能识别 创建interpolation_kernel.cu编写设备函数
#include torch/extension.htemplatetypename scalar_t
__global__ void trilinear_bw_kernal(const torch::PackedTensorAccessorscalar_t, 2, torch::RestrictPtrTraits, size_t dL_dfeats_interp,const torch::PackedTensorAccessorscalar_t, 3, torch::RestrictPtrTraits, size_t feats,const torch::PackedTensorAccessorscalar_t, 2, torch::RestrictPtrTraits, size_t points,torch::PackedTensorAccessorscalar_t, 3, torch::RestrictPtrTraits, size_t dL_dfeats
) {// 得到当前线程的在整个grid中的索引每个线程计算一组数据const int n blockIdx.x * blockDim.x threadIdx.x; // n表示每一个点的index是多少const int f blockIdx.y * blockDim.y threadIdx.y; // f表示每一个点的第f个特征// 因为实际定义的线程数量是数据量的因此会存在用不到某些线程的情况因此需要将这部分线程剔除掉// nfeats.size(0) ffeats.size(2) 因此只使用索引小于数据量的线程if (n feats.size(0) || f feats.size(2)) return;// 每个点的范围在[-1, 1]之间因此需要进行归一化 -[0, 1]const scalar_t u (points[n][0] 1) / 2;const scalar_t v (points[n][1] 1) / 2;const scalar_t w (points[n][2] 1) / 2;// 计算内插权重// 三线性插值可以分解为两组双线性插值加一组单线性插值// a b c d表示计算双线性插值时的4个顶点对应的插值权重const scalar_t a (1 - v) * (1 - w);const scalar_t b (1 - v) * w;const scalar_t c v * (1 - w);const scalar_t d 1 - a - b - c;// 计算偏导数我们需要实现根据公式计算出8个特征点的偏微分dL_dfeats[n][0][f] (1-u)*a*dL_dfeats_interp[n][f];dL_dfeats[n][1][f] (1-u)*b*dL_dfeats_interp[n][f];dL_dfeats[n][2][f] (1-u)*c*dL_dfeats_interp[n][f];dL_dfeats[n][3][f] (1-u)*d*dL_dfeats_interp[n][f];dL_dfeats[n][4][f] u*a*dL_dfeats_interp[n][f];dL_dfeats[n][5][f] u*b*dL_dfeats_interp[n][f];dL_dfeats[n][6][f] u*c*dL_dfeats_interp[n][f];dL_dfeats[n][7][f] u*d*dL_dfeats_interp[n][f];
}// template typename scalar_t 表示告诉函数函数输入参数的类型会根据scalar_t进行可变的对应
// 下面的参数的输入方法也是固定的之后可以复制后并简单修改后使用
templatetypename scalar_t
__global__ void trilinear_fw_kernal(const torch::PackedTensorAccessorscalar_t, 3, torch::RestrictPtrTraits, size_t feats,const torch::PackedTensorAccessorscalar_t, 2, torch::RestrictPtrTraits, size_t points,torch::PackedTensorAccessorscalar_t, 2, torch::RestrictPtrTraits, size_t feat_interp
) {// 得到当前线程的在整个grid中的索引每个线程计算一组数据const int n blockIdx.x * blockDim.x threadIdx.x; // n表示每一个点的index是多少const int f blockIdx.y * blockDim.y threadIdx.y; // f表示每一个点的第f个特征// 因为实际定义的线程数量是数据量的因此会存在用不到某些线程的情况因此需要将这部分线程剔除掉// nfeats.size(0) ffeats.size(2) 因此只使用索引小于数据量的线程if (n feats.size(0) || f feats.size(2)) return;// 每个点的范围在[-1, 1]之间因此需要进行归一化 -[0, 1]const scalar_t u (points[n][0] 1) / 2;const scalar_t v (points[n][1] 1) / 2;const scalar_t w (points[n][2] 1) / 2;// 计算内插权重// 三线性插值可以分解为两组双线性插值加一组单线性插值// a b c d表示计算双线性插值时的4个顶点对应的插值权重const scalar_t a (1 - v) * (1 - w);const scalar_t b (1 - v) * w;const scalar_t c v * (1 - w);const scalar_t d 1 - a - b - c;feat_interp[n][f] (1 - u) * (a * feats[n][0][f] b * feats[n][1][f] c * feats[n][2][f] d * feats[n][3][f]) u * (a * feats[n][4][f] b * feats[n][5][f] c * feats[n][6][f] d * feats[n][7][f]);
}/*如果想要在C中回传多个值那么可以直接回传一个vector类型* std::vectortorch::Tensor trilinear_fw_cu(){* return {a, b, c, ...}* }* */
torch::Tensor trilinear_fw_cu(torch::Tensor feats, torch::Tensor points) {/*使用torch库定义cuda函数* fw表示前向传播 bw表示反向传播 cu表示这是一个cuda程序* feats: [N, 8, F]表示N个元素每个元素有8个顶点以及对应的特征* points: [N, 3]表示N个元素每个元素有3个顶点坐标* */const int N feats.size(0), F feats.size(2);// 生成output变量并指定变量的数据类型并指定数据存放在哪个设备上// 将运算结果存储在output变量中torch::Tensor feat_interp torch::zeros({N, F}, torch::dtype(torch::kFloat32).device(feats.device()));// 定义grid和block大小
// 根据要同时运算的数据维度的大小建立相应大小的threadconst dim3 threads(16, 16); // 定义一个block的大小一个block内有多少个线程一般一个block定义为256个线程不会出错const dim3blocks((N threads.x - 1) / threads.x, (N threads.y - 1) / threads.y); // 定义一个grid的大小一个grid内有多少个block
/*假设输入的大小为N20F10而一个block的大小为16*16* 因此需要额外的一个block才能包含整个输入数据因此一个grid的大小应该为2*1* 线程的数量应该数据量 */// 启动一个kernel
// 下面的书写方法是固定的以后复制粘贴使用即可
/* AT_DISPATCH_FLOATING_TYPES: 表示启动的这个kernel用于浮点数运算包括float32和float64的运算* AT_DISPATCH_FLOATING_TYPES_HALF: 表示可以进行163264位的浮点数运算* AT_DISPATCH_INTEGRAL_TYPES: 表示进行32和64位整数运算* feats.type(): 表示feats的数据类型* trilinear_fw_cu: 表示启动的kernel的名字建议kernel的名称与函数的名称一致方便找错误* trilinear_fw_cuscalar_tblocks, threads: 丢出一个kernel指定kernel的名称用于在设备函数上进行相应的运算* scalar_t: 表示相当于一个placeholder因为不知道参数具体是什么类型因此先占个位置如果知道是什么类型直接写上对应的类型即可比如float* blocks, threads: 表示启动的kernel的线程数以及线程的block数* feats.packed_accessorscalar_t, 3, torch::RestrictPtrTraits, size_t()* feats.packed_accessorpoints.packed_accessorfeat_interp.packed_accessor 表示这个kernel函数传入的三个值* packed_accessor表示这个值是一个指针指向一个内存地址这个地址存储着数据并且这个地址可以被grid和block访问并且这个地址可以被cuda的线程访问* 只有Tensor参数才需要进行packed_accessor处理如果是一般的参数则直接写入参数列表即可* 比如trilinear_fw_cuscalar_tblocks, threads(a, feats.packed_accessorscalar_t, 3, torch::RestrictPtrTraits, size_t(), ...)* scalar_t: 表示参数类型* 3: 表示参数维度这里表示传入的参数维度为3* torch::RestrictPtrTraits: 表示这个参数是restrict指针表示这个参数不能被修改 (torch::RestrictPtrTraits和size_t这两个参数一般不用改)* size_t: 表示这个参数这种类型占用的内存大小用于访问内存* */AT_DISPATCH_FLOATING_TYPES(feats.type(), trilinear_fw_cu,([] {trilinear_fw_kernalscalar_tblocks, threads(feats.packed_accessorscalar_t, 3, torch::RestrictPtrTraits, size_t(),points.packed_accessorscalar_t, 2, torch::RestrictPtrTraits, size_t(),feat_interp.packed_accessorscalar_t, 2, torch::RestrictPtrTraits, size_t());}));return feat_interp;
}/*反向传播利用梯度更细参数的值 */
torch::Tensor trilinear_bw_cu(torch::Tensor dL_dfeat_interp,torch::Tensor feats,torch::Tensor points) {/* 用于反向传播计算特征的梯度* dL_dfeat_interp: 表示根据损失值计算出来的特征关于损失的梯度* */const int N feats.size(0), F feats.size(2);torch::Tensor dL_dfeats torch::zeros({N, 8, F}, feats.options()); // 微分后的值的size和原来是一样大的保存每个特征的偏微分值const dim3 threads(16, 16);const dim3 blocks((N threads.x - 1) / threads.x, (F threads.y - 1) / threads.y);AT_DISPATCH_FLOATING_TYPES(feats.type(), trilinear_bw_cu,([] {trilinear_bw_kernalscalar_tblocks, threads(dL_dfeat_interp.packed_accessorscalar_t, 2, torch::RestrictPtrTraits, size_t(),feats.packed_accessorscalar_t, 3, torch::RestrictPtrTraits, size_t(),points.packed_accessorscalar_t, 2, torch::RestrictPtrTraits, size_t(),dL_dfeats.packed_accessorscalar_t, 3, torch::RestrictPtrTraits, size_t());}));return dL_dfeats;
}3.创建header定义设备函数
创建include/utils.h头文件定义函数
#ifndef TORCH_CUDA_CPP_UTILS_H
#define TORCH_CUDA_CPP_UTILS_H#endif //TORCH_CUDA_CPP_UTILS_H#include torch/extension.h// 下面的define是必须要添加的作用类似于python的assert
// CHECK_CUDA 检查变量是否为一个gpu的tensor
// CHECK_CONTIGUOUS 检测每一个tensor在内存上是否是连续的
// CHECK_INPUT 表示检测上面两个函数
#define CHECK_CUDA(x) TORCH_CHECK(x.type().is_cuda(), #x must be a CUDA tensor)
#define CHECK_CONTIGUOUS(x) TORCH_CHECK(x.is_contiguous(), #x must be contiguous)
#define CHECK_INPUT(x) CHECK_CUDA(x); CHECK_CONTIGUOUS(x)torch::Tensor trilinear_fw_cu(torch::Tensor feats, torch::Tensor points);torch::Tensor trilinear_bw_cu(torch::Tensor dL_dfeat_interp,torch::Tensor feats,torch::Tensor points);4.编写CPP主机代码在代码中使用PYBIND库连接pytorch和c代码使用主机函数调用设备函数 PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {/*pybind库用于作为连接python和C的桥梁* name_: 表示python调用c函数的名称(名称可一致也可以不一致)f表示c函数的地址*/m.def(trilinear_interpolation_fw, trilinear_interpolation_fw);m.def(trilinear_interpolation_bw, trilinear_interpolation_bw);
}trilinear_interpolation_fw表示你在python代码中调用cuda函数时这个cuda函数的别名这个别名进行和CUDA函数保持一致报错了好找位置。trilinear_interpolation_bw同理
5.在Pycharm中创建setup.py脚本将C编写的cuda代码编译为python的一个package
一定是创建setup.py不能是别的名称
# -*- coding: utf-8 -*-
# Data: 2023/11/14 10:30
# Project: torch_cuda_cpp
# File Name: setup.py
# Author: KangPeilun
# Email: 374774222qq.com
# Description:import glob
import os.path as osp
from setuptools import setup
from torch.utils.cpp_extension import CppExtension, BuildExtension, CUDAExtensionROOT_DIR rE:\CLion\Projects\torch_cuda_cpp # 填入C项目的根目录
include_dirs [osp.join(ROOT_DIR, include)] # 获取头文件文件夹
sources glob.glob(osp.join(ROOT_DIR, *.cpp)) glob.glob(osp.join(ROOT_DIR, *.cu)) # 获取所有cpp和cu文件路径# 这个setup.py脚本的目的是为了编译一个C的扩展模块
# 这里使用的CppExtension编译的其实是CPU版本的代码并没有调用cuda
# setup(
# namecuda_ext, # 定义package的名称也就是python中import调出来的名称
# version1.0, # 然后你可以定义版本作者简介
# authorKang Peilun,
# author_email374774222qq.com,
# descriptioncudacpp example,
# long_description
# this is a cudacpp example,
# ext_modules[ # 指定需要build的c代码在哪里, 在sources参数中指定代码路径如果有多个路径则在list中一次添加即可
# CppExtension(namecuda_ext, sources[
# rE:\CLion\Projects\torch_cuda_cpp\interpolation.cpp # 可以指定C文件的绝对路径
# ])
# ],
# cmdclass{ # 告诉代码需要build这个东西
# build_ext: BuildExtension
# }
# )# 对于CUDA程序的编译需要使用CUDAExtension
setup(namecuda_ext, # 定义package的名称也就是python中import调出来的名称version1.0, # 然后你可以定义版本作者简介authorKang Peilun,author_email374774222qq.com,descriptioncudacpp example,long_descriptionthis is a cudacpp example,ext_modules[ # 指定需要build的c代码在哪里, 在sources参数中指定代码路径如果有多个路径则在list中一次添加即可CUDAExtension(namecuda_ext,sourcessources, # 导入cpp和cu文件的路径include_dirsinclude_dirs, # 导入header的所在的文件夹)],cmdclass{ # 告诉代码需要build这个东西build_ext: BuildExtension}
)6.使用pip安装自定义的cuda代码库
使用setup.py编译安装自定义的CUDA库 注意一定得是setup.py所在文件夹
pip install .
# . 表示setup.py所在的文件夹
# 当setup.py在别的文件夹中时把.替换为setup.py所在文件夹路径即可
# 注意一定得是setup.py所在文件夹出现下方字样即代表安装成功
7.安装好自定义的cuda自定义库之后编写python代码调用 # -*- coding: utf-8 -*-
# Data: 2023/11/14 15:24
# Project: torch_cuda_cpp
# File Name: interpolation.py
# Author: KangPeilun
# Email: 374774222qq.com
# Description:
import torch
import cuda_ext
from time import timedef trilinear_interpolation_py(feats, points)::param feats: [N, 8, F]:param points: [N, 3]:return:u (points[:, 0:1]1)/2v (points[:, 1:2]1)/2w (points[:, 2:3]1)/2a (1-v)*(1-w)b (1-v)*wc v*(1-w)d 1-a-b-c # d v*wfeats_interp (1-u)*(a*feats[:, 0] b*feats[:, 1] c*feats[:, 2] d*feats[:, 3]) \u*(a*feats[:, 4] b*feats[:, 5] c*feats[:, 6] d*feats[:, 7])return feats_interpclass Trilinear_interpolation_cuda(torch.autograd.Function):使用torch.autograd.Function包装fw和bw处理实现这个类需要手动实现forward和backward两个函数并且要用staticmethod修饰ctx: 负责存储反向传播用到的值这个参数不能省略forward的返回值的个数与backward输入参数的个数一致backward的返回值个数与forward的输入参数个数一致且返回计算过梯度的参数如果某个参数不需要计算梯度则对应返回Nonestaticmethoddef forward(ctx, feats, points):feat_interp cuda_ext.trilinear_interpolation_fw(feats, points)ctx.save_for_backward(feats, points)return feat_interpstaticmethoddef backward(ctx, dL_dfeat_interp):feats, points ctx.saved_tensorsdL_feats cuda_ext.trilinear_interpolation_bw(dL_dfeat_interp.contiguous(), feats, points)return dL_feats, Noneif __name__ __main__:N 65536F 256# feats torch.rand(N, 8, F, devicecuda).requires_grad_()rand torch.rand(N, 8, F, devicecuda)feats1 rand.clone().requires_grad_()feats2 rand.clone().requires_grad_()points torch.rand(N, 3, devicecuda)*2 - 1t time()# 自定义的CUDA函数是没有办法自动计算梯度的需要我们手动计算梯度# 通过torch.autograd.grad自动保存计算后的梯度# 使用 .apply 执行自定义的CUDA前向传播过程out_cuda Trilinear_interpolation_cuda.apply(feats1, points)torch.cuda.synchronize()print(\tCUDA time: , time()-t, out_cuda)# 1.使用Pytorch计算梯度t time()out_py trilinear_interpolation_py(feats2, points) # Pytorch可以自动计算梯度torch.cuda.synchronize()print(\tPytorch time: , time()-t, out_py)## print(torch.allclose(out_cuda, out_py)) # 判断两个tensor在误差范围内是否接近loss2 out_py.sum() # 简单的将他们的和作为loss仅用于测试loss2.backward() # 反向传播计算梯度# 2.使用自定义CUDA函数计算梯度loss1 out_cuda.sum()loss1.backward()print(bw all close, torch.allclose(feats1.grad, feats2.grad))五、总结
第四部分介绍了整个Pytorch CUDA CPP项目一般会包含哪些文件所有代码都是基于教程Pytorchcpp/cuda extension 教学 编写推荐有时间和能力的人把这个教程学完会受益匪浅的。
总体来说Pytorch CUDA CPP项目一般同时会包含.py、.cpp、.cu、.h文件cpp文件存放主机函数用于调用设备函数cu文件存放设备函数用于调用kernel核函数在GPU上进行运算setup.py文件用于将自定义的cuda程序构建为python的一个package其他py文件调用package实现python调用C代码C代码调用CUDA程序这一流程。
