随笔分类 - 高性能计算
摘要:1. 介绍: PTX 指令集中 WMMA 矩阵计算从共享内存加载数据到 fragment 片段使用的封装API是 load__matrix_sync,其底层 PTX指令与mma 一致,并且 fragment 布局一致。本文介绍底层 ldamatrix、stmatrix 指令的行为,并且用代码进行验证
阅读全文
摘要: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]={
阅读全文
摘要:1. 目标:矩阵向量乘法 y = A * x (列向量 = 矩阵 *列向量),进行串行,循环展开+simd, simd+omp的效率对比。 2. 源码 #include <iostream> #include <ctime> #include <arm_neon.h> #include <omp.h
阅读全文
摘要:1. 这是一个对图像通道 bgra 转换为 rgb的示例程序。转换方式有普通写法、openmp-simd编译指导语句、neon intrinsic函数三种实现方式 2. 源码 #include <stdio.h> #include <stdlib.h> #include <iostream> #in
阅读全文
摘要:1. 目标 对数组求和,对比ARM(neon)向量化优化(SIMD)和循环展开,还有 O0 O1优化的效率对比。 2. 测试代码 #include <arm_neon.h> #include <stdio.h> #include <stdlib.h> #include <time.h> #defin
阅读全文
摘要:1. 目标:分析循环分块优化技术,并分析cache 命中情况 假设每个cacheline可以存储b个数据元素。 2. 源代码分析 for( int i=0;i<N;i++) { for(int j=0;j<M;j++) { A[i] += B[j]; } } cache miss分析: 对A总访问次
阅读全文
摘要:1. 目标:使用openmp实现图像滤波算法 给定一个输入图像,你的任务是将其划分为多个块,并使用多个进程对每个块进行滤波操作。 要求: 1、使用分段并行结构和指令section来并行执行不同的计算步骤。 2、使用单线程指令single或master确保某些代码段只在单个线程中执行。 3、使用线程同
阅读全文

浙公网安备 33010602011771号