yan31415的博客

原创 CUDA性能优化系列——Kmeans算法调优（三）

本篇对调度方式进行优化，实现内存拷贝和计算overlap。单流同步调用/*单流同步*/void CallKmeansSync(){ //TODO:init host memory float* h_Src/*[Coords][SrcCount]*/, * h_Clusters/*[ClusterCount][Coords]*/; int* h_MemberShip/*[kSrcCount]*/, * h_MemberCount/*[kClusterCount]*/; h_Src

原创 CUDA性能优化系列——Kmeans算法调优（二）

本篇介绍Kmeans算法中计算新的聚类中心部分。这部分主要逻辑：根据计算出的新的分类信息，对全部数据点依次对每个类别求出所属当前类别的数据点个数与坐标和。本质上就是进行规约运算。V1 Atomic实现全局规约由于最终生成16个聚类中心，因此这里的规约操作需要针对算法进行一定的调整。V1在实现逻辑为：先在共享内存上分别通过原子操作，对16个类进行规约，再通过原子操作进行设备内存的全局规约操作。/*V1 atomic规约*/__global__ void kKmeansSumAtomic(int Src

原创 CUDA性能优化系列——Kmeans算法调优（一）

__global__ void k_kmeans(float* d_src, int srcsize, int dim, float* d_cluster, int clustersize, float* d_dst){ extern __shared__ float sm_cluster[]; float regData[4] = { 0.0f,0.0f, 0.0f, 0.0f };//维度不确定的情况下how to use registers???????? int tid = threadId

原创 三种基于CUDA的归约计算

归约在并行计算中很常见，并且在实现上具有一定的套路。本文分别基于三种机制（Intrinsic，共享内存，atomic），实现三个版本的归约操作，完成一个warp大小的数据的归约求和计算。Intrinsic版本基于Intrinsic函数 __shfl_down_sync 实现，使一个warp内的线程通过读取相邻线程寄存器中数据，完成归约操作。实现如下：__global__ void kIntrinsicWarpReduce(int* d_src, int* d_dst){ int val = d_

原创 可分离卷积CUDA实现

记录代码#include "include/SepConv.cuh"#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#include <stdlib.h>#include <string>#include <cooperative_groups.h>#include <iostream>#include <ss

原创 CUDA 统计数组直方图

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdio.h>#include <stdlib.h>#include <string>#include <cooperative_groups.h>#include <iostream>#define HISTOGRAMGRID 36#define HISTOGRAMBLOCK 1.

原创 Nsight Compute内存访问常用Metrics含义理解

Memory L1 Transcations Global:实际全局内存加载至L1缓存的内存交换次数，粒度128bytesMemory L2 Transactions Global:实际全局内存加载至L2缓存的内存交换次数，粒度32bytes,该参数的值应该是Memory L1 Transcations Global 的4倍Memory Ideal L2 Transactions Global:理论需要从全局内存加载至L2缓存的内存交换次数，当数值比Memory L2 Transactions Glob

原创 Pinned Memory 多设备异步拷贝

void TestHostPinnedMem(){ //compare pinned and pagable memcpy const int size = 1024 * 1024 * 100; int* h1 = (int*)malloc(size * sizeof(int)); int* d1; cudaMalloc(&d1, size * sizeof(int)); cudaMemcpy(d1, h1, size * sizeof(int), cudaMemcpyHostToDe

原创 CUDA处理归约问题

归约问题由于计算操作的对称性，非常适合并行处理。本文以数组求和为例，通过CUDA先实基础版本，并基于基础版本尝试通过不同的优化手段实现几个方案，最后将所有优化手段集成到最终的实现。其中核函数执行时间是通过Nsight System工具抓取，命令参数如下nsys profile --stats true a.out 5基础版本实现核函数实现__global__ void kReduceWholeWarpSafe(int64_t* d_src, int64_t* d_sum){ int64_t*

原创 nvcc 编译时输出核函数寄存器使用情况

nvcc --ptxas-options=-v kernel.cu

原创 Nsight Compute与nvprof metrics 对照

NVIDIA 计算能力7.5及以上的GPU设备不再支持nvprof工具进行性能剖析，提示使用Nsight Compute作为替代品，如下图所示。Nsight Compute Cli（命令行）剖析的参数与nvprof不一样，当按照nvprof的参数抓取数据时，因为参数不识别，无法抓取希望得到的指标，如下图所示；同时，Nsight Compute Cli参数成千上万，虽然可以将这些参数全部专区，但是会对使用者筛选关注信息带来很大的麻烦。因此，非常有必要找到两个工具参数之间的对应关系。这里记录命令行使用N

转载 vi中 wq 、wq!、x、q、q!区别

上面的命令只是在vi编辑命令中使用wq：表示保存退出wq!：表示强制保存退出x：表示保存退出wq和wq!的区别如下：有些文件设置了只读，一般不是修改文件的，但是如果你是文件的owner或者root的话，通过wq!还是能保存文件退出如果文件设置为只读了的话，用 :wq命令是不能保存并退出的，但是最高权限者可通过wq!来进行文件的保存并退出文件。已设定选项 ‘readonly’ (请加 ! 强制执行)文件所有者通过 wq! 可以保存只读文件!是强制执行如果不带!碰上只读文件，会给提示会是只读

原创 Everything搜索网络路径

在 工具->选项->索引->文件夹->添加 中添加网络路径并保存，配置后如下：配置成功后可以在网络路径中搜索文件，如下：

原创 root用户 ssh远程登录 提示access denied

ssh远程登录 提示Access denied，如下图所示：需要修改SSH服务配置文件sshd_config来解决这个问题。文件路径为 /etc/ssh/sshd_config ，默认文件内容如下图：图中红框内容修改为 PermitRootLogin yes，如下图所示：修改好配置文件并保存，service ssh restart 重启SSH服务，ssh可以成功登陆，如下图：...

原创 Nsight Compute 使用

记录使用Nsight Compute 分析cuda性能的方法。1.单击菜单栏上的Connet，弹出如下界面，设置要剖析的执行程序路径等执行相关参数，选择Interactive Profile模式，可以对剖析流程进行控制，所有参数设置完成后，单击Launch开始性能分析。2.在API Stream执行到要剖析的核函数，然后单击菜单栏的 Profile Kernel对核函数进行剖析3.生成性能分析结果，可以通过切换Page选项查看关注的参数，如Source ，Detail等Detail信息：记录

原创 Nsight Compute Profile Kernel无法定位源码问题

Nsight Compute是NVIDIA提供的性能分析工具。本人在剖析核函数过程中遇到了无法定位源码的问题。在此记录原因与解决方法。如图，在Profile Kernel时，想要查看源码中某行代码的性能情况，发现只能查看汇编，无法查看源码。造成这个现象的原因是：编译时没添加调试信息，因此只能定位汇编，无法定位源码。处理方法：这里编译通过VS集成CUDA的编译器，设置调试信息方法如下：打开Project->属性->CUDA C/C+±>Device ->Generate重

原创 cuda 配置共享内存Bank模式API

记录GPU共享内存4Byte or 8Byte模式的cuda API。Access Mode Configuration对Kepler来说，默认情况是4-byte模式，可以用下面的API来查看：cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);返回结果放在pConfig中，其结果可以是下面两种：cudaSharedMemBankSizeFourBytecudaSharedMemBankSizeEigh

原创 NotePad++ 缩进显示XML

NotePad++打开XML文件，发现所有内容都在同一行，没有按照缩进的形式展示。可以通过安装插件XML Tool解决问题。安装打开 Plugins ->Plugins Admin，在Avalable页面中搜索xml，找到XML Tool插件并安装。因为我这里已经安装了该插件，因此无法显示搜索结果。使用快捷键 Ctrl + Alt + Shift +B，或者点击Plugins ->XML Tools->Pretty Print效果如下...

原创 C++性能优化系列——3D高斯核卷积计算(十二)Intrinsic实现

void GaussSmoothCPU3DOptZYXSplitZIntrinsic(float* pSrc, int iDim[3], float* pKernel, int kernelSize[3], float* pDst, float* pBuffer) { //计算结果正确 //16 thread dynamic GaussSmoothCPU3DOptZYXSplitZIntrinsic cost Time(ms) 145.8 int iSliceSize = iDim[1] *

原创 C++性能优化系列——3D高斯核卷积计算(十一)优化内存访问模式

void GaussSmoothCPU3DOptZYXSplitZIntrinsic(float* pSrc, int iDim[3], float* pKernel, int kernelSize[3], float* pDst, float* pBuffer) { //计算结果正确 //16 thread dynamic GaussSmoothCPU3DOptZYXSplitZIntrinsic cost Time(ms) 145.8 int iSliceSize = iDim[1] *

原创 C++性能优化系列——3D高斯核卷积计算(十)合并多线程并行区

void GaussSmoothCPU3DOptZYX(float* pSrc, int iDim[3], float* pKernel, int kernelSize[3], float* pDst, float* pBuffer) { //计算结果正确 //执行时间 16 dynamic GaussSmoothCPU3DOptZYX cost Time(ms) 338.4 //16 static GaussSmoothCPU3DOptZYX cost Time(ms) 590 //

原创 Intel优化手册 - PAUSE指令解释

PAUSE指令提升了自旋等待循环（spin-wait loop）的性能。当执行一个循环等待时，Intel P4或Intel Xeon处理器会因为检测到一个可能的内存顺序违规（memory order violation）而在退出循环时使性能大幅下降。PAUSE指令给处理器提了个醒：这段代码序列是个循环等待。处理器利用这个提示可以避免在大多数情况下的内存顺序违规，这将大幅提升性能。因为这个原因，所以推荐在循环等待中使用PAUSE指令。PAUSE的另一个功能就是降低Intel P4在执行循环等待时的耗电量。I

转载 Intel64及IA-32架构优化指南——3.4优化执行引擎前端

3.5 优化执行核心在最近微架构产品中的超标量、无序执行核心含有多个执行硬件资源，可以并行地执行多个微操作。这些操作一般确保微操作高效地执行并以固定的延迟执行下去。利用可用的并行性的通用准则有：●    遵循规则（3.4小节）以最大化有用的译码带宽以及前端吞吐。这些规则包括偏向使用单个微操作指令并利用微融合、栈指针追踪器以及宏融合。●    最大化重命名带宽。在本小节所讨论到的准则包含适当的对部分寄存器、ROB［译者注：重排序缓存］读端口以及对标志引起...

转载 Intel X86 优化指南阅读笔记--基础体系结构

Intel Sandy Bridge Microarchitecture 流水线 前端按顺序取指令和译码，将X86指令翻译成uop。通过分支预测来提前执行最可能的程序路径。带有超标量功能的执行引擎每时钟周期最多执行6条uop。带有乱序功能的执行引擎能够重排列uop执行顺序，只要源...

转载 Intel X86 优化指南阅读笔记--后端优化

OPTIMIZING THE EXECUTION CORE 通用的优化准则来充分利用超标量cpu的并行性： - 遵循前端优化准则，最大化译码带宽和前端吞吐量。 - 最大化寄存器重命名带宽。本文将讨论包括如何正确处理partial registers, ROB read port...

转载 Intel X86 优化指南阅读笔记--前端优化

PROCESSOR PERSPECTIVES 以下优化建议，在不同微架构下收益差别较大： 指令译码的吞吐量很重要。利用好decoded ICache，Loop Stream Detector和macro-fusion能进一步提高CPU前端性能。充分利用好4个译码器来产生代码。利用...

原创 C++性能优化系列——3D高斯核卷积计算(九)3D高斯卷积OpenMP Task优化内存访问

本篇基于 C++性能优化系列——3D高斯核卷积计算(八)3D高斯卷积中的代码实现的计算逻辑，通过OepnMP的Task特性，优化Z维度卷积计算中内存性能瓶颈。（本篇只优化内存访问问题，不保证总执行时间会提升）。代码实现void GaussSmoothCPU3DBase_Task(float* pSrc, int iDim[3], float* pKernel, int kernelSize[3], float* pDst, float* pBuffer) { int iSliceSize = iD

原创 超线程对多线程并行化的影响

本篇基于 C++性能优化系列——3D高斯核卷积计算(八)3D高斯卷积 中的代码实现，测试超线程对多线程并行性的影响。代码实现测试平台支持8核16线，具体信息参考 C++性能优化系列——百倍加速比的矩阵转置性能调优并行区开辟8线程void GaussSmoothCPU3DBase_OneParallelRegion(float* pSrc, int iDim[3], float* pKernel, int kernelSize[3], float* pDst, float* pBuffer) {

原创 C++性能优化系列——3D高斯核卷积计算(八)3D高斯卷积

本篇基于可分离卷积的性质，按照 X Y Z的顺序，依次计算每个维度的一维卷积。代码实现因为是按照X Y Z的计算顺序，因此只能够在计算X维度的卷积时，复用之前实现的一维卷积计算函数。Y维度的计算是将一个Z平面上的二维数据中每行与卷积核中一个点相乘，并将31个点的卷积核计算出的结果累加至一行，更新到中间缓存的目标位置。Z维度的计算是将一个Z平面的二维数据和卷积核中的一个点相乘，并将31个点的卷积核计算出的结果累加至一个二维平面，更新到结果的目标位置。这里对Y 和 Z维度的计算都是通过编译器ICC实现向量化

原创 C++性能优化系列——3D高斯核卷积计算(七)Intrinsic实现2D高斯卷积

本篇基于 C++性能优化系列——3D高斯核卷积计算(六)交换维度计算2D卷积 中2D高斯卷积的计算逻辑，通过Intrinsic函数实现相同的功能并对比性能差异。代码实现先用指令实现对x维度的计算void Conv2D_Fuse_InstructX(float* pSrcSlice, int iDim[2], float* pKernel, int iKernelSize, float* pBuffer, float* pDstSlice,float* pSimdKernel) { //combi

原创 C++性能优化系列——3D高斯核卷积计算(六)交换维度计算2D卷积

在 C++性能优化系列——3D高斯核卷积计算(五)2D卷积计算 中，基于可分离卷积的性质，先计算x维度卷积，再将y维度卷积计算过程打乱并重组，完成了两个维度的向量化计算。本篇以先计算y维度卷积，后计算x维度卷积的顺序计算2D高斯卷积。代码实现代码实现void Conv2D_Fuse(float* pSrcSlice, int iDim[2], float* pKernel, int iKernelSize, float* pBuffer, float* pDstSlice,float* pSimdKe

原创 C++性能优化系列——3D高斯核卷积计算(四)Intrinsic实现一维卷积与循环展开

基于 性能优化系列(CPU)——3D高斯核卷积计算(三)FMA向量化计算一维卷积 中实现，本文用Intrinsic实现向量化的一维卷积计算。代码实现void Conv1D_Ins_Cmb(float* pSrcLine, int iLength, float* pKernel, int kernelSize, float* pDst,float* pSimdKernel) { int iConvSize = iLength - kernelSize + 1; float* pDstTemp

原创 ICC编译器生成优化报告

记录两个操作系统上ICC编译器生成优化报告的设置步骤。WindowsICC手册设置说明VS属性页面设置如下：这里对ICC优化报告的选项参数进行简单说明。Optimization Diagnostic File：生成报告文件。注意，Windows操作系统上生成的报告后缀是 .rep。Optimization Diagnostic Level：报告详细程度。ICC手册说明如下。Optimization Diagnostic Phase：报告期项（优化类型）。LinuxCMakeLi

原创 C++性能优化系列——3D高斯核卷积计算(五)2D卷积分离计算

高斯卷积核具有可分离的性质，因此可以通过以下方法计算二维高斯卷积：构造一个一维高斯卷积核，将原始二维矩阵分别以行主序与列主序，与一维卷积核做卷积计算，得到的结果就是目标二维高斯卷积的结果。本篇按照上述描述的思路实现了可分离的二维高斯卷积计算，并在此基础上对计算的过程分解与重构，挖掘实现的并行性。基线版二维高斯卷积为了让运行时间更加稳定，增加函数的执行次数至1000#define CONV2DREPT 1000这里对代码实现功能做一下简单的说明。原始矩阵维度432 * 432，维度定义为x和y。高斯

转载 Nsight Visual Studio 调试

官方手册：https://docs.nvidia.com/nsight-visual-studio-edition/5.4/Nsight_Visual_Stud...

原创 C++性能优化系列——3D高斯核卷积计算(三)IPP计算一维卷积

本篇通过Intel高性能计算库IPP中的卷积计算API实现一维卷积，并记录程序耗时情况，同时与 性能优化系列(CPU)——3D高斯核卷积计算(三)FMA向量化计算一维卷积 中只通过编译器向量化版本对比性能。功能实现代码实现：void Conv1d_IPP(float* pSrc, int iLength, float* pKernel, int kernelSize, float* pDst) { IppEnum funCfg = (IppEnum)(ippAlgAuto); int buf

原创 CUDA C编程权威指南笔记 第五章 共享内存和常量内存

原创 CUDA C编程权威指南笔记 第七章 调整指令级原语

原创 CUDA C编程权威指南笔记 第六章 流和并发

原创 C++性能优化系列——3D高斯核卷积计算题目

计划写一个优化3D高斯核卷积计算的系列博客。3D高斯核卷积计算的需求来自于公司内部举办的性能优化竞赛，这里介绍题目信息。题目题目三3D卷积(相关)操作 （20分）【场景和性能问题描述】背景：在图像处理中，卷积（滤波）是一个基本、常用的操作，很多图像处理方法都要用到，但是卷积是非常耗时的操作，尤其是在卷积核比较大、图像比较大(3D)的情况。当前XXX配准35%时间用在高斯卷积上，所以对它进行性能研究有实际价值。题目：对一个432432457的图像做3D高斯卷积，卷积核大小为313131。【要求】