anluo8

摘要： 1. TensorCore 简介： 硬件层面支持半精度浮点矩阵乘法，与昇腾NPU的 cube 核类似，最小只能计算规定尺寸的矩阵乘法。 wmma API 封装在 nvcuda 命名空间 2. naive : 一个block 1 个warp，wmmaTile 16*16 点击查看代码 //naive 阅读全文

摘要： 1. 介绍 参考文章：https://www.cnblogs.com/cuancuancuanhao/p/7763256.html ，本文对 cublas gemm 的接口参数进行原理解释。 2. 接口 cuBLAS中用于运算矩阵乘法的函数有4个，分别是 cublasSgemm（单精度实数）、cub 阅读全文

摘要： 1.介绍 对 2048 * 512 矩阵转置，使用NCU进行性能分析，并进行性能优化。测试环境 CUDA 12.8,显卡 5070。 2. Native: 二维 Block 二维block,一个线程处理一个元素 点击查看代码 //native：二维block,一个线程处理一个元素 //矩阵 M * 阅读全文

摘要： 1. 背景 最开始在学习 cuda 编程时，只知道 warp scheduler 线程束调度的概念，但是不清楚调度的细节。现在查看 CUDA Pragramming Guide 性能优化篇看到了关于 warp 调度更清晰的细节。 2. 概念 延时：一个线程束准备好执行下一条指令的时钟周期，受数据依赖 阅读全文

摘要： 1. 介绍： 基于最近对大模型 KV_cache，及 Attention 变种学习中遇到的问题和理解记录下来，帮助大家解决一点疑惑。 2. kv_cache 显存对比： 参数说明 batch_size：B seq_len：L head_num：H head_dim：D layer_num：N gro 阅读全文

摘要： 1. 项目介绍 源项目仓：https://gitee.com/ascend/samples/tree/master/operator/ascendc/0_introduction/20_mmad_kernellaunch 目标项目仓：https://gitee.com/ascend/cann-ops 阅读全文

摘要： 1. 介绍 矩阵向量乘法: A * X = Y, A(M,K) X(K,1) Y(M,1); 实现多种并行算法及优化方法和 cublas 库 sgemv 的效率对比。 2. gpu 并行算法介绍 并行算法一：一个线程计算一个结果元素； 并行算法二：使用合并访存,需要将输入数据转置； 并行算法三: 合 阅读全文

摘要： 1. 介绍 矩阵A（MK） B（KN）单精度浮点数进行矩阵乘法； 分别实现CPU串行，GPU多种并行计算算法，与 cublas 库 sgemm 函数效率对比。 2. gpu 并行算法简介 并行算法一：二维block,一个线程程计算一个C元素,缺点:访存次数过多 并行算法二: 优化一: 使用线程块 t 阅读全文

摘要： 1. 扫描概念 对数组arr[N]扫描就是得到数组prefix[N]，每个元素是之前arr元素的求和. 开扫描定义:prefix1[N] = { arr[0], arr[0]+arr[1], ..., arr[0]+arr[1]+arr[N-1] } 闭扫描定义: prefix2[N] = { 0, 阅读全文

摘要： 1. 目标：对数组进行求和，并做优化对比 2. baseline 代码 相邻求和: 根据blockSize对数据分块,并将数据放在共享内存,以线程块为单位,块内线程数量=数据个数,相邻配对,用其中第一个元素索引为ID的线程进行计算,计算结果放在第一个元素位置,循环进行下一轮计算,最后块求和计算结果赋 阅读全文

摘要： 1. 延迟隐藏作用：可以最大化SM的使用效率，提高计算性能 2. 延迟隐藏概念 （1）指令延迟是指令发出到完成之间的时钟周期间隔； （2）指令可以分为两种：算数指令和访存指令。 3. 算数指令延迟隐藏实现 （1）假如一个算数指令的指令延迟是 4 个时钟周期 （2）查询cuda文档得到如下计算能力的设 阅读全文

摘要： 1. 简介 (1) 使用CPU对向量点乘进行串行计算 (2) 对数据进行分块,使用单进程多卡(多流)并行计算 (3) 使用不同数据规模,比较加速比的变化 2. 代码 #include <stdio.h> #include <sys/time.h> #include <stdlib.h> #defin 阅读全文

摘要： 1. 介绍： (1) 用CUDA计算 pow(sin(id),2)+ pow(cos(id),2)的结果 (2) 对比单流(同步传输、异步传输)、多流深度优先调度、多流广度优先调度的效率(包含数据传输和计算) 核心代码 1. 用CUDA计算 pow(sin(id),2)+ pow(cos(id),2 阅读全文

摘要： 1. 目标：对 16384*16384 规模的矩阵进行加法运算，对比 CPU 和 GPU 计算的效率，还有不同线程块大小规模下对效率的影响；并做可能的优化测试。 2. 核心代码 /* 用GPU对二维矩阵做加法，分析不同线程块规模下的性能变化 */ #include <stdio.h> #includ 阅读全文

摘要： 1. 目标：使用CPU和GPU对一千万数量级的一维信号进行均值滤波，并且根据GPU存储模型对数据存储进行优化，最终对比计算结果并计算加速比。 2. 代码 /* cuda实现对一维信号卷积平滑滤波处理，并于串行计算对比结果和加速比,卷积核大小为5 */ #include <stdio.h> #incl 阅读全文

摘要： 1. GPU的内存模型 GPU编程数据需要从CPU主存拷贝到GPU全局存储器，所有线程共享全局存储。开辟的全局存储器空间指针在CPU代码中不能解引用使用，应在计算完结果后再拷贝回CPU主存空间。线程块内共享存储。 (1) 线程私有的存储有寄存器、本地内存 (2) 线程块内有块内线程共享的共享内存，在 阅读全文

摘要： 1. 简介 (1) Intel® Integrated Performance Primitives，即英特尔集成性能基元（简称IPP），为信号、数据和图像处理特定应用领域，提供simd优化的一组全面的函数库。 (2) 本项目将对 exp、cos、sin、tone、Triangle函数用NEON向量 阅读全文

摘要： 1. 源码为对粒子移动状态模拟的项目。要求使用多种优化方法，对比串行优化、多线程优化、全部优化下的加速比。 2. 代码 项目代码地址：https://github.com/libo-0379/StellarSim_Optimize 以下为核心优化代码及分析 #include <stdlib.h> # 阅读全文

摘要： 1. 目标：使用 NEON intrinsic 函数，对512*512 png 四通道图像顺时针旋转90度。 思路: 像素分块，对块内转置；再水平镜像。图像库使用 stb img 2. 代码 #include <stdio.h> #include <arm_neon.h> #include <std 阅读全文

摘要： 1. 对寄存器数据重排 /* 两个向量，每两个通道一组，第一个向量每组的后一个元素与第二个向量每组的第一个元素一次彼此交换 */ #include <stdio.h> #include <arm_neon.h> void main() { int arrc[8]={0}; int arrd[4]={ 阅读全文