Cython与CUDA之BatchGather

最新推荐文章于 2025-03-10 16:18:20 发布
最新推荐文章于 2025-03-10 16:18:20 发布
阅读量21 收藏
点赞数
合集 - Cython专辑(8)
1.Cython与C函数的结合 2.使用Cython调用CUDA Kernel函数 3.使用Cython加速谐振势的计算 4.Cython编译报错“numpy/arrayobject.h: No such file or directory”解决方案 5.Cython将Numpy数组转为自定义结构体 6.Cython与CUDA之Gather 7.CUDA时长统计
8. Cython与CUDA之BatchGather
收起

技术背景

在前面一篇文章中，我们介绍过Cython+CUDA框架下实现一个简单的Gather算子的方法。这里演示Gather算子的升级版本实现——BatchGather算子。不过这里只是加了一个Batch维度，并没有添加其他的维度，例如Dimension维度，在这里暂不考虑。

CUDA头文件

这里我们保留了原本的Gather部分，只添加一个BatchGather的运算，以下为cuda_index.cuh的内容：

#include <stdio.h>

extern "C" float Gather(float *source, int *index, float *res, int N, int M);
extern "C" float BatchGather(float *source, int *index, float *res, int N, int M, int B);

BatchGather只是在Gather的基础上加了一个B的维度。除了CUDA算子本身的头文件之外，这里我们还使用到了异常捕获头文件error.cuh：

#pragma once
#include <stdio.h>

#define CHECK(call) do{const cudaError_t error_code = call; if (error_code != cudaSuccess){printf("CUDA Error:\n"); printf("    File:   %s\n", __FILE__); printf("    Line:   %d\n", __LINE__); printf("    Error code: %d\n", error_code); printf("    Error text: %s\n", cudaGetErrorString(error_code)); exit(1);}} while (0)

其中的宏可用于检测CUDA函数所抛出的异常。另外还有一个用于统计CUDA函数运行时长的头文件：

#pragma once
#include <stdio.h>
#include <cuda_runtime.h>

// 宏定义，用于测量CUDA函数的执行时间
#define TIME_CUDA_FUNCTION(func) \
    do { \
        cudaEvent_t start, stop; \
        float elapsedTime; \
        cudaEventCreate(&start); \
        cudaEventCreate(&stop); \
        cudaEventRecord(start, NULL); \
        \
        func; \
        \
        cudaEventRecord(stop, NULL); \
        cudaEventSynchronize(stop); \
        cudaEventElapsedTime(&elapsedTime, start, stop); \
        printf("Time taken by function %s is: %f ms\n", #func, elapsedTime); \
        \
        cudaEventDestroy(start); \
        cudaEventDestroy(stop); \
    } while (0)

// 宏定义，用于测量CUDA函数的执行时间并返回该时间
#define GET_CUDA_TIME(func) \
    ({ \
        cudaEvent_t start, stop; \
        float elapsedTime = 0.0f; \
        cudaEventCreate(&start); \
        cudaEventCreate(&stop); \
        cudaEventRecord(start, NULL); \
        \
        func; \
        \
        cudaEventRecord(stop, NULL); \
        cudaEventSynchronize(stop); \
        cudaEventElapsedTime(&elapsedTime, start, stop); \
        \
        cudaEventDestroy(start); \
        cudaEventDestroy(stop); \
        \
        elapsedTime; \
    })

可选择直接打印时长，也可以选择返回时长的float值。

CUDA文件

接下来就是正式的CUDA函数内容cuda_index.cu：

// nvcc -shared ./cuda_index.cu -Xcompiler -fPIC -o ./libcuindex.so
#include <stdio.h>
#include "cuda_index.cuh"
#include "error.cuh"
#include "record.cuh"

__global__ void GatherKernel(float *source, int *index, float *res, int N){
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N){
        res[idx] = source[index[idx]];
    }
}

extern "C" float Gather(float *source, int *index, float *res, int N, int M){
    float *souce_device, *res_device;
    int *index_device;
    CHECK(cudaMalloc((void **)&souce_device, M * sizeof(float)));
    CHECK(cudaMalloc((void **)&res_device, N * sizeof(float)));
    CHECK(cudaMalloc((void **)&index_device, N * sizeof(int)));
    CHECK(cudaMemcpy(souce_device, source, M * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(res_device, res, N * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(index_device, index, N * sizeof(int), cudaMemcpyHostToDevice));
    int block_size = 1024;
    int grid_size = (N + block_size - 1) / block_size;
    float timeTaken = GET_CUDA_TIME((GatherKernel<<<grid_size, block_size>>>(souce_device, index_device, res_device, N)));
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaMemcpy(res, res_device, N * sizeof(float), cudaMemcpyDeviceToHost));
    CHECK(cudaFree(souce_device));
    CHECK(cudaFree(index_device));
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaFree(res_device));
    CHECK(cudaDeviceReset());
    return timeTaken;
}

__global__ void BatchGatherKernel(float *source, int *index, float *res, int N, int M, int B){
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N*B){
        int batch_idx = idx / N;
        int source_idx = batch_idx * M + index[idx];
        res[idx] = source[source_idx];
    }
}

extern "C" float BatchGather(float *source, int *index, float *res, int N, int M, int B){
    float *souce_device, *res_device;
    int *index_device;
    CHECK(cudaMalloc((void **)&souce_device, B * M * sizeof(float)));
    CHECK(cudaMalloc((void **)&res_device, B * N * sizeof(float)));
    CHECK(cudaMalloc((void **)&index_device, B * N * sizeof(int)));
    CHECK(cudaMemcpy(souce_device, source, B * M * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(res_device, res, B * N * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(index_device, index, B * N * sizeof(int), cudaMemcpyHostToDevice));
    int block_size = 1024;
    int grid_size = (B * N + block_size - 1) / block_size;
    float timeTaken = GET_CUDA_TIME((BatchGatherKernel<<<grid_size, block_size>>>(souce_device, index_device, res_device, N, M, B)));
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaMemcpy(res, res_device, B * N * sizeof(float), cudaMemcpyDeviceToHost));
    CHECK(cudaFree(souce_device));
    CHECK(cudaFree(index_device));
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaFree(res_device));
    CHECK(cudaDeviceReset());
    return timeTaken;
}

这里传入到CUDA之前，我们需要在Cython或者Python中把相关的数据压缩为一维，所以传入CUDA函数的是一个一维的指针。相比于单一的Gather操作，BatchGather中的几个输入含义有所变化，例如N表示的是单Batch的Index长度，M表示的是单Batch的源数组长度。

Cython文件

对于一个新的Batch函数来说，我们需要构建一个新的Cython调用函数wrapper.pyx：

# cythonize -i -f wrapper.pyx

import numpy as np
cimport numpy as np
cimport cython

cdef extern from "<dlfcn.h>" nogil:
    void *dlopen(const char *, int)
    char *dlerror()
    void *dlsym(void *, const char *)
    int dlclose(void *)
    enum:
        RTLD_LAZY

ctypedef float (*GatherFunc)(float *source, int *index, float *res, int N, int M) noexcept nogil
ctypedef float (*BatchGatherFunc)(float *source, int *index, float *res, int N, int M, int B) noexcept nogil

cdef void* handle = dlopen('/path/to/libcuindex.so', RTLD_LAZY)

@cython.boundscheck(False)
@cython.wraparound(False)
cpdef float[:] cuda_gather(float[:] x, int[:] idx):
    cdef:
        GatherFunc Gather
        float timeTaken
        int N = idx.shape[0]
        int M = x.shape[0]
        float[:] res = np.zeros((N, ), dtype=np.float32)
    Gather = <GatherFunc>dlsym(handle, "Gather")
    timeTaken = Gather(&x[0], &idx[0], &res[0], N, M)
    print (timeTaken)
    return res

@cython.boundscheck(False)
@cython.wraparound(False)
cpdef float[:] batch_cuda_gather(float[:] x, int[:] idx, int B):
    cdef:
        BatchGatherFunc BatchGather
        float timeTaken
        int N = idx.shape[0] // B
        int M = x.shape[0] // B
        float[:] res = np.zeros((B*N, ), dtype=np.float32)
    BatchGather = <BatchGatherFunc>dlsym(handle, "BatchGather")
    timeTaken = BatchGather(&x[0], &idx[0], &res[0], N, M, B)
    print (timeTaken)
    return res

while not True:
    dlclose(handle)

这里我们还是接受一维的数组，多引入一个Batch维度的参数B，其他的都是一样的。

Python调用文件

最后是用来调用的最上层Python端的代码test_gather.py：

import numpy as np
np.random.seed(0)
from wrapper import batch_cuda_gather

B = 2
M = 1024 * 1024 * 128
N = 1024 * 1024

x = np.random.random((M*B,)).astype(np.float32)
idx = np.random.randint(0, M, (N*B,)).astype(np.int32)

np_res = np.zeros((B, N), dtype=np.float32)
for i in range(B):
    np_res[i] = x.reshape((B,-1))[i][idx.reshape((B, -1))[i]]
np_res = np_res.reshape(-1)

res = np.asarray(batch_cuda_gather(x, idx, B))
print (res.shape)
print ((res==np_res).sum())

为了方便处理，在构建数据的时候，我们直接在生成数据阶段就生成一维的数据，然后直接调用Cython函数进行CUDA相关运算。

运行方法

首先将CUDA文件编译成动态链接库，使其可以在Cython中被调用。然后将Cython文件编译成动态链接库，使其可以在Python中被调用。最后运行Python代码即可：

$ nvcc -shared ./cuda_index.cu -Xcompiler -fPIC -o ./libcuindex.so
$ cythonize -i -f wrapper.pyx
$ python3 test_gather.py

运行结果如下：

0.9606080055236816
(2097152,)
2097152

这表示CUDA核函数部分的运行时长为0.96ms，输入的数组总长度为2097152，跟numpy版本的数组索引实现对比之后，得到2097152个相同的元素。也就是说，计算结果跟numpy的计算结果是一致的，以此来校验CUDA部分的运算结果。

总结概要

以学习CUDA为目的，接上一篇关于Cython与CUDA架构下的Gather算子实现，这里我们加一个Batch的维度，做一个BatchGather的简单实现。

版权声明

本文首发链接为：https://www.cnblogs.com/dechinphy/p/cython-cuda-batchgather.html

作者ID：DechinPhy

更多原著文章：https://www.cnblogs.com/dechinphy/

请博主喝咖啡：https://www.cnblogs.com/dechinphy/gallery/image/379634.html

原创作者: dechinphy 转载于: https://www.cnblogs.com/dechinphy/p/18748323/cython-cuda-batchgather
李Moon
关注
Python硬件加速：Cython与CUDA混合编程.pdf
04-20
文档内所有文字、图表、函数、目录等元素均显示正常，无任何异常情况，敬请您放心查阅与使用。文档仅供学习参考，请勿用作商业用途。 想轻松敲开编程大门吗？Python 就是你的不二之选！它作为当今最热门的编程语言，...
CyFES是一款基于Python/Cython和CUDA混合开发的高性能自由能计算工具，面向用户开放Python API接口
09-04
这是一个使用Cython+CUDA+Python编写的高性能FES计算软件，可以用于加速数据透视的计算。CyFES是一款基于Python/Cython和CUDA混合开发的高性能自由能计算工具，面向用户开放Python API接口。（源码） 这是一个使用...
Cython与CUDA之Add
pantouyuchiyu的博客
03-10 782
本文介绍了使用CUDA和Cython来实现一个CUDA加法算子的方法，并介绍了使用CUDA参数来估算性能极限的算法。经过实际测试，核函数部分的算法性能优化空间已经不是很大了，更多时候可以考虑使用Stream来优化Host和Device之间的数据传输。
CUDA与Cython之BatchGather
fenfenfen520a的博客
03-03 18
技术背景 在前面一篇文章中，我们介绍过Cython+CUDA框架下实现一个简单的Gather算子的方法。这里演示Gather算子的升级版本实现——BatchGather算子。不过这里只是加了一个Batch维度，并没有添加其他的维度，例如Dimension维度，在这里暂不考虑。 CUDA头文件 这里我
Cython与CUDA之Gather
weixin_40325199的博客
03-07 33
合集 - Cython专辑(9)1.Cython与C函数的结合2024-07-252.使用Cython调用CUDA Kernel函数2024-08-023.使用Cython加速谐振势的计算2024-07-244.Cython编译报错“numpy/arrayobject.h: No such file or directory”解决方案2024-06-265.Cython将Numpy数组转为自定义结...
CUDA时长统计
weixin_35884307的博客
03-05 24
合集 - Cython专辑(8)1.Cython与C函数的结合2024-07-252.使用Cython调用CUDA Kernel函数2024-08-023.使用Cython加速谐振势的计算2024-07-244.Cython编译报错“numpy/arrayobject.h: No such file or directory”解决方案2024-06-265.Cython将Numpy数组转为自定义结...
使用Cython+CUDA+Python开发的高性能FES计算软件，用于加速数据透视，对用户开放Python API接口（源码）
08-27
使用Cython+CUDA+Python编写的高性能FES计算软件，可以用于加速数据透视的计算。 pip安装 本项目可以直接使用pip进行安装： $ python3 -m pip install cyfes --user --upgrade -i https://pypi.org/simple 源码...
基于Cython和CUDA的Python高性能FES计算软件发布
通过将Cython与CUDA结合，可以实现将Python代码中需要高性能计算的部分转移到GPU上执行，显著提高数据处理速度。 FES（Finite Element Software）计算软件是一种用于进行有限元分析的工具，广泛应用于工程计算、...
cython_bbox-0.1.3
03-26
Cython 的使用可能意味着这个包提供了高效的边界框操作，比如计算、比较或调整边界框，以优化与这些操作相关的性能。 该包的文件名列表只给出了 "cython_bbox-0.1.3"，这通常是软件包的根目录，包含了该包的源码、...
ffmpeg 的视频格式转换 c# win10
最新发布
06-02
ffmpeg 的视频格式转换 c# win10
基于51单片机利用MPU6050 DMP实现姿态角获取的完整例程
06-02
标题“51单片机通过MPU6050-DMP获取姿态角例程”解析 “51单片机通过MPU6050-DMP获取姿态角例程”是一个基于51系列单片机（一种常见的8位微控制器）的程序示例，用于读取MPU6050传感器的数据，并通过其内置的数字运动处理器（DMP）计算设备的姿态角（如倾斜角度、旋转角度等）。MPU6050是一款集成三轴加速度计和三轴陀螺仪的六自由度传感器，广泛应用于运动控制和姿态检测领域。该例程利用MPU6050的DMP功能，由DMP处理复杂的运动学算法，例如姿态融合，将加速度计和陀螺仪的数据进行整合，从而提供稳定且实时的姿态估计，减轻主控MCU的计算负担。最终，姿态角数据通过LCD1602显示屏以字符形式可视化展示，为用户提供直观的反馈。 从标签“51单片机 6050”可知，该项目主要涉及51单片机和MPU6050传感器这两个关键硬件组件。51单片机基于8051内核，因编程简单、成本低而被广泛应用；MPU6050作为惯性测量单元（IMU），可测量设备的线性和角速度。文件名“51-DMP-NET”可能表示这是一个与51单片机及DMP相关的网络资源或代码库，其中可能包含C语言等适合51单片机的编程语言的源代码、配置文件、用户手册、示例程序，以及可能的调试工具或IDE项目文件。 实现该项目需以下步骤：首先是硬件连接，将51单片机与MPU6050通过I2C接口正确连接，同时将LCD1602连接到51单片机的串行数据线和控制线上；接着是初始化设置，配置51单片机的I/O端口，初始化I2C通信协议，设置MPU6050的工作模式和数据输出速率；然后是DMP配置，启用MPU6050的DMP功能，加载预编译的DMP固件，并设置DMP输出数据的中断；之后是数据读取，通过中断服务程序从DMP接收姿态角数据，数据通常以四元数或欧拉角形式呈现；再接着是数据显示，将姿态角数据转换为可读的度数格
TSP（旅行商问题）-使用人工蜂群求解源码
06-02
程序100%可运行，可随机增加或减少城市，带有代码注释；也可转换为其他语言的代码
金昌EX9000教程学习资料同步实体书
06-01
同步当年出版书，详细教材，包含三维试衣毛毛虫使用方法
2021年MathorCup高校数学建模挑战赛D题
06-02
MathorCup高校数学建模挑战赛是一项旨在提升学生数学应用、创新和团队协作能力的年度竞赛。参赛团队需在规定时间内解决实际问题，运用数学建模方法进行分析并提出解决方案。2021年第十一届比赛的D题就是一个典型例子。 MATLAB是解决这类问题的常用工具。它是一款强大的数值计算和编程软件，广泛应用于数学建模、数据分析和科学计算。MATLAB拥有丰富的函数库，涵盖线性代数、统计分析、优化算法、信号处理等多种数学操作，方便参赛者构建模型和实现算法。 在提供的文件列表中，有几个关键文件： d题论文(1).docx：这可能是参赛队伍对D题的解答报告，详细记录了他们对问题的理解、建模过程、求解方法和结果分析。 D_1.m、ratio.m、importfile.m、Untitled.m、changf.m、pailiezuhe.m、huitu.m：这些是MATLAB源代码文件，每个文件可能对应一个特定的计算步骤或功能。例如： D_1.m 可能是主要的建模代码； ratio.m 可能用于计算某种比例或比率； importfile.m 可能用于导入数据； Untitled.m 可能是未命名的脚本，包含临时或测试代码； changf.m 可能涉及函数变换； pailiezuhe.m 可能与矩阵的排列组合相关； huitu.m 可能用于绘制回路图或流程图。 matlab111.mat：这是一个MATLAB数据文件，存储了变量或矩阵等数据，可能用于后续计算或分析。 D-date.mat：这个文件可能包含与D题相关的特定日期数据，或是模拟过程中用到的时间序列数据。 从这些文件可以推测，参赛队伍可能利用MATLAB完成了数据预处理、模型构建、数值模拟和结果可视化等一系列工作。然而，具体的建模细节和解决方案需要查看解压后的文件内容才能深入了解。 在数学建模过程中，团队需深入理解问题本质，选择合适的数学模
yolo11-pyqt5-gui检测施工 site 中的安全隐患-保障 construction workers 的安全+数据集+训练好的模型+pyqt5可视化界面.zip
06-02
yolo11-pyqt5-gui检测施工 site 中的安全隐患-保障 construction workers 的安全+数据集+训练好的模型+pyqt5可视化界面包含pyqt可视化界面，有使用教程 1. 内部包含标注好的目标检测数据集，分别有yolo格式(txt文件)和voc格式标签(xml文件), 共1206张图像， 已划分好数据集train，val, test，并附有data.yaml文件可直接用于yolov5,v8,v9,v10,v11,v12等算法的训练； 2. yolo目标检测数据集类别名：construction-safety（施工安全），包括 helmet（安全帽）、no-helmet（无安全帽）、no-vest（无背心）、person（人员）、vest（背心）等 3. yolo项目用途：检测施工 site 中的安全隐患，保障 construction workers 的安全 4. 可视化参考链接：https://blog.youkuaiyun.com/weixin_51154380/article/details/126395695?spm=1001.2014.3001.5502
等高线图绘制算法源码重新编写
06-02
以下是关于三种绘制云图或等高线图算法的介绍： 一、点距离反比插值算法 该算法的核心思想是基于已知数据点的值，计算未知点的值。它认为未知点的值与周围已知点的值相关，且这种关系与距离呈反比。即距离未知点越近的已知点，对未知点值的影响越大。具体来说，先确定未知点周围若干个已知数据点，计算这些已知点到未知点的距离，然后根据距离的倒数对已知点的值进行加权求和，最终得到未知点的值。这种方法简单直观，适用于数据点分布相对均匀的情况，能较好地反映数据在空间上的变化趋势。 二、双线性插值算法 这种算法主要用于处理二维数据的插值问题。它首先将数据点所在的区域划分为一个个小的矩形单元。当需要计算某个未知点的值时，先找到该点所在的矩形单元，然后利用矩形单元四个顶点的已知值进行插值计算。具体过程是先在矩形单元的一对对边上分别进行线性插值，得到两个中间值，再对这两个中间值进行线性插值，最终得到未知点的值。双线性插值能够较为平滑地过渡数据值，特别适合处理图像缩放、地理数据等二维场景中的插值问题，能有效避免插值结果出现明显的突变。 三、面距离反比 + 双线性插值算法 这是一种结合了面距离反比和双线性插值两种方法的算法。它既考虑了数据点所在平面区域对未知点值的影响，又利用了双线性插值的平滑特性。在计算未知点的值时，先根据面距离反比的思想，确定与未知点所在平面区域相关的已知数据点集合，这些点对该平面区域的值有较大影响。然后在这些已知点构成的区域内，采用双线性插值的方法进行进一步的插值计算。这种方法综合了两种算法的优点，既能够较好地反映数据在空间上的整体分布情况，又能保证插值结果的平滑性，适用于对插值精度和数据平滑性要求较高的复杂场景。
继电器驱动电路设计合集
06-02
合集包含内容目录： 1.继电器原理 2.继电器驱动电路设计指引（美的） 3.继电器驱动电路 4.大功率驱动电路 5.单片机驱动继电器电路详解 6.继电器控制 7.继电器生产原理
基于MATLAB实现的空间前方交会算法代码
06-02
在摄影测量领域，通过 MATLAB 编写的空间前方交会程序能够实现待定点三维坐标的计算。该程序基于点投影系数法，利用空间后方交会所求得的双像外方位元素，进而推导出待定点的三维坐标。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值