【CUDA学习笔记】OneFlow公众号CUDA算子优化文章学习笔记

已于 2022-12-08 22:06:25 修改
阅读量1.6k 收藏 15
点赞数 1
文章标签： 学习 oneflow 人工智能
于 2022-12-07 16:09:25 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/songyuc/article/details/128222051
版权

1 CUDA学习资料合集

【OneFlow】岁末年初，为你打包了一份技术合订本

2 GPU概念介绍

《GPU的硬件结构与执行原理 —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

2.1 内存模型

2.1.1 Bank介绍

《GPU硬件结构之bank —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化》

3 算子优化

3.1 Conv

3.1.1 Img2col：卷积优化算法

博文《基于OneFlow实现Unfold、Fold算子》（以下简称为“《Fold优化》”）
《基于OneFlow实现Unfold、Fold算子》：理解img2col

3.1.2 Unfold & Fold

《基于OneFlow实现Unfold、Fold算子》：Unfold、Fold算子是卷积优化的基础操作

为什么这里out在索引时设计成6维的方式来进行操作呢？

在阅读《Fold优化》时，我们发现out采用6维的形式来进行操作，
在这里插入图片描述
这样是为了CUDA编程时，索引可以直接对应上去，这样代码写作起来会更加简单；

3.2 Norm

3.2.2 LayerNorm

《CUDA优化之LayerNorm性能优化实践》

3.3 Elementwise operator

《高效、易用、可拓展我全都要：OneFlow CUDA Elementwise模板库的设计优化思路》

3.4 Softmax

《如何实现一个高效的 Softmax CUDA kernel？——OneFlow 性能优化分享》

（1）OneFlow为什么在Softmax实现时会使用ReduceMax操作呢？

这个问题的来源是这样的，博文《如何实现一个高效的 Softmax CUDA kernel？——OneFlow 性能优化分享》（以下简称为“《Softmax优化》”）在描述Softmax的CUDA实现时表示使用了ReduceMax操作，（也就是求某个维度上的最大值），但是根据Softmax的公式，这个操作在数学上其实是没有必要的，那为什么OneFlow会在CUDA实现时使用ReduceMax操作呢？
关于这一点，我们请教了晓雨哥，

【晓雨哥】：
应该是防溢出吧。

于是我们可以知道，先进行ReduceMax操作的目的是为了减去最大值，从而减小每个元素的绝对值，防止指数操作可能产生的数值溢出。

3.5 Dim transformation

《如何实现比PyTorch快6倍的Permute/Transpose算子？》

4 性能优化方法

4.1 CUDA Kernel中 grid_size 和 block_size 的设置优化

《如何设置CUDA Kernel中的grid_size和block_size？》

4.2 访存优化

《OneFlow GPU性能优化方法一：减少全局内存的访问 —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

4.2.1 Kernel fusion：核函数融合，一次访存，多次计算

（1）Element-wise kernel fusion

《1. Element-wise kernel fusion —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

（2）借助 shared memory 合并带有Reduce计算的Kernel

《2. 借助shared memory合并带有Reduce计算的Kernel —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

（3）使用bitset优化mask计算

《3. 减少实际需要的访存大小（以ReLU为例） —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

4.1.1 Memory access merging：内存访问合并

《OneFlow GPU性能优化方法二：确保全局内存访问合并 —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

以 bitset mask 生成为例讲解访存合并

《如何生成这个 bitset mask 加速Kernel的访存 —— 开源100天，OneFlow送上“百天大礼包”：深度学习框架如何进行性能优化 》

4.2 计算优化

《OneFlow GPU性能优化方法三：优化Kernel计算量 》

4.3 延迟隐藏

《OneFlow GPU性能优化方法四：延迟隐藏 》

4.4 优化小技巧

《OneFlow 其他GPU性能优化小技巧 》

5 CUDA高性能计算经典优化问题

① 归约

《CUDA高性能计算经典问题①：归约》

② 前缀和

《CUDA高性能计算经典问题②：前缀和》

③ 矩阵乘法

《关于ChatGPT的一切；CUDA入门之矩阵乘；PyTorch 2.0发布｜AI系统前沿动态》

songyuc
关注
  • 1
    点赞
  • 踩
  • 15
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
高效、易用、可拓展我全都要：OneFlow CUDA Elementwise模板库的设计优化思路
OneFlow深度学习框架
12-21 694
OneFlow 针对这种 Elementwise 操作抽象出一套 CUDA 模板，开发者只需把计算逻辑封装到一个结构体内，即可获得一个 CUDA Elementwise 算子
OneFlow微信公众号二维码
OneFlow深度学习框架
08-20 836
CUDA——reduce算子及优化（Ⅰ）
最新发布
weixin_46706926的博客
04-12 976
在以上for循环中，index = 1时，if(tid%2==0) 偶数id的线程，执行smem[tid] += smem[tid + 1]，例如tid=0时，smem[0]+=smem[0+1];将1号线程的值加到0号线程上，并保存到0号线程。在优化（Ⅰ）中，我们在shared memory上做并行规约计算，然后把最终结果先存在shared memory[0]，再把结果写回显存。可以看到，对8个数据而言，串行和并行（2、3方式）的实现效率已经对比很明显了，时间复杂度O（N）与O（log(N))。
DeepRoute Lab | CUDA算子优化：转置篇
DeepRoute_Lab的博客
02-21 1456
转置是深度模型中应用最广泛的算子之一。各种深度学习框架崛起的今天，它被“隔离”到了框架的底层，不再显式的活跃于用户面前。谈及转置，就不得不先聊聊卷积的优化算法以及它们所需要的图像数据格式。
cuda编程-卷积优化
weixin_34405332的博客
01-16 617
CUDA Convolution https://www.evl.uic.edu/sjames/cs525/final.html Improve Image Processing Using GPU Computing on Mali™ https://www.arm.com/files/event/2014_ARM_Multimedia_Seminar_Arcsoft.pdf GPU ...
CUDA优化之LayerNorm性能优化实践
OneFlow深度学习框架
12-09 1139
撰文 | 郭冉、姚迟、郑泽康、柳俊丞2020年末，OneFlow 发布了《OneFlow 性能优化分享：如何实现一个高效的 Softmax CUDA kernel？》 ，其中介绍了OneF...
【yolov4】学习笔记一
12-21
这位大神厉害啊，这么快就复现了,好好学习膜拜中！ yolov4.weights权重百度网盘下载地址: https://pan.baidu.com/s/1dAGEW8cm-dqK14TbhhVetA 提取码:dm5b 下载后用pycharm打开工程。 个人更喜欢把参数写在代码中，...
CUDA学习笔记，GPU编程，C++编程
06-08
本篇笔记将深入探讨CUDA编程的基础概念，包括GPU架构、并行计算原理以及CUDA编程的关键要素。 首先，我们需要理解为什么需要GPU进行并行计算。传统的CPU（中央处理器）擅长于串行处理任务，但对于大规模并行计算，...
我的CUDA学习笔记
08-29
"CUDA学习笔记" 本资源摘要信息中，我们将对CUDA学习笔记进行详细的知识点梳理，涵盖了CUDA程序架构、Thread管理、存储器体系、混合编程、CUDA程序接口描述等方面的知识点。 1. CUDA 程序架构 CUDA 程序架构中...
CUDA-:CUDA C学习笔记
03-16
这篇学习笔记将深入探讨CUDA C的基础知识和关键概念。 1. **CUDA架构** - **CUDA核心**: GPU内部包含大量的流处理器（Streaming Multiprocessors, SM），每个SM由多个CUDA核心组成，它们执行并行计算任务。 - **...
北大tensorflow学习笔记
11-09
在"北大tensorflow学习笔记"中，我们可以深入探索这个广泛应用于学术研究和工业实践的平台。 一、TensorFlow基础知识 1.1 **张量（Tensor）**：TensorFlow的核心概念，是多维数据数组，可以是标量、向量、矩阵甚至...
CUDA卷积详解
11-04
描述了GPU架构下CUDA编程处理卷积的方法
cuda 优化步骤
08-04
cuda概述，cuda架构，cuda内核优化详细步骤说明
CUDA卷积计算及其优化——以一维卷积为例
qq_40491305的博客
04-28 3486
CUDA 卷积计算及其优化——以一维卷积为例 CUDA C 编程学习，关注专栏 CUDA C 建议阅读：在卷积优化前，熟悉核函数的组织形式有利于更好的位置映射-CUDA编程入门（一）：以图片运算看线程的组织和核函数的使用 对于输入数据为N[Width]，卷积核大小为M[Mask_Width]的卷积运算，进行不同程度的优化（这里的卷积指滤波/内积，而不需要旋转），输出为P[Width]。 1.常规的一维卷积 常规的一维卷积比较简单，线程数为Width，每个线程负责一个输出值得Mask_Width宽度的卷积
【BBuf的CUDA笔记】一，解析OneFlow Element-Wise 算子实现
I good vegetable a!
12-11 1105
以上就是我对 OneFlow Element-Wise 系列 CUDA 算子实现的解析，后续有空会持续更新学习到的新知识。
【BBuf的CUDA笔记】四，介绍三个高效实用的CUDA算法实现（OneFlow ElementWise模板，FastAtomicAdd模板，OneFlow UpsampleNearest2d模板）
I good vegetable a!
01-08 262
本篇文章推荐和讲解一下OneFlow ElementWise模板，FastAtomicAdd，OneFlow UpsampleNearest2d模板的用法以及原理，并将其整理为最小的可以白嫖的头文件。相关代码请访问 https://github.com/BBuf/how-to-optim-algorithm-in-cuda 这里获得。
CUDA加速的卷积计算
禅与计算机程序设计艺术
04-07 651
CUDA加速的卷积计算 作者：禅与计算机程序设计艺术 1. 背景介绍 卷积运算是深度学习和图像处理领域中最基础和最重要的操作之一。传统的CPU实现卷积计算效率较低,尤其是在处理大规模数据时,计算速度很难满足实际需求。为了解决这一问题,CUDA作为一种通用GPU加速计算平台应运而生,可以
Oneflow概要
smartcat2010的博客
01-10 336
厚Compile阶段；薄Runtime阶段； 去中心化调度：以Op（叫做Actor）为核心，每个Actor有输入buffer和输出buffer；输入buffer有可用的且输出buffer有空闲的，来驱动Actor执行 天然支持pipeline并行：各个Actor并行执行； “整个数据流的执行像一个网络，数据在网络中的流动就完成了计算，如何避免生产者生产太快，消费者消费不及，以及如何避免生产者生产太慢，消费者感到饥饿的问题，这涉及到对计算、内存、传输带宽的规划，尽可能使系统的瓶颈之处最宽，...
如何在 GPU 上优化卷积
吴建明wujianming_110117
10-30 678
如何在 GPU 上优化卷积 将演示如何在 TVM 中编写高性能卷积实现。正方形大小的输入张量和过滤器为例，假设卷积的输入具有大batch批量。在这个例子中，使用不同的布局存储数据，实现更好的数据局部性。缓冲区布局是 HWCN，代表高度、宽度、通道、批次。 准备和算法 对具有 256 个通道和 14 x 14 维度的输入张量使用固定大小。批量大小为 256。卷积过滤器包含 512 个大小为 3 x 3 的过滤器。使用步幅大小 1 和padding大小 1 进行卷积。以下代码定义了 TVM 中的卷积算法。 im
如何用cuda c写卷积算子
03-29
以下是一个简单的使用CUDA C编写卷积操作的示例代码： ```c #include <stdio.h> #include <stdlib.h> #include <cuda_runtime.h> #define BLOCK_SIZE 16 __global__ void convolve(float *input, float *output, float *kernel, int input_width, int input_height, int kernel_size) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; float sum = 0.0f; for (int i = 0; i < kernel_size; ++i) { for (int j = 0; j < kernel_size; ++j) { int input_row = row + i - kernel_size / 2; int input_col = col + j - kernel_size / 2; if (input_row >= 0 && input_row < input_height && input_col >= 0 && input_col < input_width) { sum += input[input_row * input_width + input_col] * kernel[i * kernel_size + j]; } } } output[row * input_width + col] = sum; } int main() { int input_width = 512; int input_height = 512; int kernel_size = 5; float *input = (float*)malloc(input_width * input_height * sizeof(float)); for (int i = 0; i < input_width * input_height; ++i) { input[i] = rand() % 256 / 255.0f; } float *kernel = (float*)malloc(kernel_size * kernel_size * sizeof(float)); for (int i = 0; i < kernel_size * kernel_size; ++i) { kernel[i] = rand() % 256 / 255.0f; } float *output = (float*)malloc(input_width * input_height * sizeof(float)); float *d_input, *d_kernel, *d_output; cudaMalloc(&d_input, input_width * input_height * sizeof(float)); cudaMalloc(&d_kernel, kernel_size * kernel_size * sizeof(float)); cudaMalloc(&d_output, input_width * input_height * sizeof(float)); cudaMemcpy(d_input, input, input_width * input_height * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_kernel, kernel, kernel_size * kernel_size * sizeof(float), cudaMemcpyHostToDevice); dim3 block_size(BLOCK_SIZE, BLOCK_SIZE); dim3 grid_size((input_width + BLOCK_SIZE - 1) / BLOCK_SIZE, (input_height + BLOCK_SIZE - 1) / BLOCK_SIZE); convolve<<<grid_size, block_size>>>(d_input, d_output, d_kernel, input_width, input_height, kernel_size); cudaMemcpy(output, d_output, input_width * input_height * sizeof(float), cudaMemcpyDeviceToHost); cudaFree(d_input); cudaFree(d_kernel); cudaFree(d_output); free(input); free(kernel); free(output); return 0; } ``` 该示例代码中，`convolve()`函数是卷积操作的核函数，它接收输入图像、输出图像和卷积核作为输入参数，以及输入图像尺寸和卷积核尺寸。在核函数中，每个线程负责计算输出图像中的一个像素值，它将卷积核与输入图像中对应像素的值相乘，并将结果累加到一个变量中。最后，输出图像中对应像素的值被赋为累加的结果。 在主函数中，我们首先生成随机的输入图像和卷积核，然后在GPU上分配内存，将输入图像和卷积核从主机内存复制到设备内存中，调用卷积核函数进行卷积操作，最后将输出图像从设备内存复制到主机内存中，并释放分配的内存。 在实际使用中，您需要根据自己的需求修改输入图像、卷积核和卷积核函数，以适应不同的场景。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值