fb_help的专栏

cuda 版本cuda toolkit版本是用于开发cuda程序的 针对开发者cuda 驱动是用于运行程序的 针对用户例子对于开发者而言例如 pytorch v1.0 的开发者们使用了cuda 10的新特性。但他们的cuda 环境不会打包给你你一个开发者想使用pytorch 开发算法，这时除了安装 pytorch v1.0 外还需要配置代码里面使用的cuda的环境。这时，你只需要使用CUDA 10.0 以上的toolkit就可以了（兼容cuda 10的特性

跑新版CUDA为什么要更新驱动更新驱动是为了让现有的设备支持CUDA程序的新性能。设备是硬件，实际操纵硬件的是驱动程序，而驱动程序也给cuda程序提供了接口，实际上一串代码操纵GPU的流程是这样的：cuda程序 ----> 二进制文件 -----> 驱动解析机器码 ------>控制GPU运行。当cuda版本更新，如cuda 11多了新特性，那么驱动程序也必须识别这样的新特性，这时如果有程序使用了cuda 11新特性，驱动就必须更新才能在GPU上跑。为此，NVIDIA一直在维护主

cudaErrorLaunchOutOfResources错误代码在头文件中的定义如下：/*** This indicates that a launch did not occur because it did not have* appropriate resources. Although this error is similar to* ::cudaErrorInvalidConfiguration, this error usually indicates that the* use

CUDA：纹理内存的作用（Texture memory)1.容量很大的常量内存纹理内存是DRAM上的内存，可以申请很大的空间，相比常量内存只能申请64kb来说，是一种很大空间的常量内存，而常量内存的好处是可以广播，当多个swap访问同一位置时，广播机制可以减少全局内存的访问，来提速。如图2. 边界问题与插值问题纹理内存有边界与插值的优化，有句话讲纹理内存永不越界就是因为纹理内存的边界优化。...

GPU架构的三层“exploit”GPU架构相比CPU架构快的原因多人工作，去掉管理者不像CPU那样需要做自动的优化（流水线，层次缓存，分支预测等等），把这部分的资源消耗拿出来，多来一下实际工作的单元（多核心）SIMT单指令多线程，把一个工作单元内放置多个线程，这些线程共享相同的指令，相当于给一个单元“三头六臂”延迟隐藏当在等待数据的时候，去做其它的指令（到要求指令之间存在独立性）...

NSight使用–查看CPU和CUDA时间线NSight查看CPU时间线需要管理员权限，这时要退出NSight，使用管理员权限打开NSight，然后在VS中启动NSight。如图 勾选Trace Setting中的System和CUDA点击lanuch： TimeLine： ...

CUDA：延迟隐藏详解延迟隐藏线程束调度器的功能是调度线程束参与指令（流水线）的执行，例如运行指令和内存指令。 如果活跃的线程束有32个，线程束调度器有4个，线程束调度器每时钟周期会调度4个线程束进行指令执行，也就是说需要8个时钟周期，32个活跃的线程束都会执行。那么如果说一个指令的耗时或者说延迟是7个周期，当线程束调度器调度最后4给线程执行该指令的时候，第一组线程束的指令已经执行完...

Efficiency of CUDA vector types (float2, float3, float4)https://stackoverflow.com/questions/26676806/efficiency-of-cuda-vector-types-float2-float3-float4I’m expanding njuffa’s comment into a wor...

查询NVIDIA相关信息网址查询设备支持计算能力https://developer.nvidia.com/cuda-gpus一些CUDA技巧方面的谈论https://developer.nvidia.com/gpu-computing-webinarsCUDA开发者https://developer.nvidia.com/developer-program计算能力与CU...

CUDA:核函数中几种变量所属内存的类型局部变量：存在寄存器中 静态数组：存在本地内存（栈）中 动态变量：存在全局内存中 __shared__（静态）：存在共享内存中...

CUDA：NVCC编译过程和兼容性详解https://codeyarns.com/2014/03/03/how-to-specify-architecture-to-compile-cuda-code/ https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#supported-phasesCUDA：NVCC...

CUDA：寄存器详解前言1与CPU不同，GPU的每个SM（流多处理器）有上千个寄存器。CPU与GPU架构的一个主要区别就是CPU与GPU映射寄存器的方式。CPU通过使用寄存器重命名和栈来执行多线程。为了运行一个新任务，CPU需要进行上下文切换，将当前所有寄存器的状态保存到栈上，然后从栈中恢复当前需要执行的新线程上次的执行状态。这些操作需要花费上百个CPU时钟周期。如果在CPU上开启过多...

CUDA: Occupancy(占用率)详解占用率是指每个多处理器（Streaming Multiprocessor，SM）的活动线程束（warps）数量与实际的活动warps数量的比率。 高的占用率不一定能提升性能，但低的占用率会降低内存延迟隐藏的作用，Higher occupancy does not always equate to higher performance-there...

- -ptxas-options=-v 命令前言.ptx文件存放了kernel的相关执行情况，C/C++程序可通过调用.ptx的信息来运行核函数，完成cuda功能，但.ptx不会显式生成，通过编译时向编译器nvcc输入-ptx（生成ptx文件） 或 -keep（生成全部编译中间文件） 可生成.ptx文件1。 但.ptx文件不是今天讨论的重点，但它与我们要讲的--ptxas-options...

texture memory摘取前言摘取自cuda programming guide中texture memory相关部分Texture memoryWhy use texture memory The texture and surface memory spaces reside in device memory and are cached in texture ca...

CUDA：Texture所需参数详解前言texture在设置时，要选择几个必要的参数，此篇详细介绍这些参数Texture 定义时所需要的参数：read mode：cudaReadModeNormalizedFloat or cudaReadModeElementType读取像素的值是归一化或不归一化（原值）texture coordinates mode：norma...

CUDA：Texture 图像应用详解什么时候使用texture原则： 1.当对图片内存的读取效率不高时，即（memory reads do not follow the access patterns） 2.果涉及到插值，对影像数据的处理使用Texture。 第一种情况真的很少，因为现在缓存的颗粒都是32，所以在一般情况下，全局内存访问的效率都比较高。这时就没有必要使用textur...

CUDA texture surface应用纹理的优势使用纹理有几个性能优势。纹理可以暗示插值（即，使用浮点坐标从纹理读取）。任何需要这种数据插值的应用程序都可以受益于GPU上纹理单元内的HW插值引擎。在任意GPU代码中使用纹理最重要的另一个好处是纹理缓存，它可以备份存储在全局内存中的纹理。纹理是一种只读操作，但如果您有一个只读数据数组，纹理缓存可能会改善或扩展您快速访问数据的能力。这通常...

影像畸变矫正带扭曲参数s(skew)像素比例ρOpenCV自带纠正方法不能矫正扭曲参数s和ρ。因此自己写代码来完成纠正。包括OpenCV的畸变模型，和Smart3D的畸变模型，其原理见影像畸变纠正详解 Mat src_cpu = imread(in_filename); cv::Mat mapx, mapy; mapx.create(src_cpu.size(), CV_32FC1);...

OpenCV实现影像畸变矫正GPUOpenCV实现影像矫正使用的是initUndistortRectifyMap()计算畸变的映射remap()计算映射，其详解见：OpenCV函数remap详解 Mat src_cpu = imread(in_filename); cv::cuda::GpuMat src(src_cpu); cv::cuda::GpuMat distortion(sr...

CUDA编程注意传给CUDA编译器编译的文件里不能包含boost的头文件，会报错。例如xxCUDA.cuh中最好不要包含boost的头文件。CUDA编程中核函数一般写在.cu文件中，也可以使用.cu生成的ptx文件（起到了类似OpenGL中的着色器的作用）添加到C++的程序中，cuda给了一套使用ptx编程的接口，这使得CUDA程序不需要.cu文件。详情见https://www.cnblogs...

可分页内存，固定内存，零拷贝内存可分页内存是cuda中最长使用的其申请使用cudaMalloc()，可分页的意思是因页面错误导致的操作，该操作按照操作系统的要求将主机虚拟内存上的数据移动到不同的物理位置，我的理解，可分页内存是可以被硬盘记录的内存（如内存超限，或电脑睡眠，内存可以和硬盘进行交换）， 固定内存 GPU不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数...

openacc使用cuda申请的内存https://stackoverflow.com/questions/16349400/openacc-memory-management

使用CUDA加速CPU程序的步骤：通过性能分析工具（如vs）找到CPU程序最耗时的多个地方，并确定耗时程序的入口函数将CPU函数进行清理 1.将循环部分的代码找出来。 2.将函数内所用到的数据从C++类结构变成C的结构体。 3.标准化输入输出，保证其为C结构，并与原程序的数据进行无缝对接。 4.将循环内部的函数也做相同处理，最终得到C版本的且输入输出与原程序对接的CPU程序。 5...

CUDA寄存器资源测试实验前言最近在编码中发现了错误：an illegal memory access was encountered. 在我过度使用本地内存（局部内存，kernel中的内存）时。 首先：该错误产生的原因是内存越界或多个内存请求访问同一内存。为此我做实验验证产生此错误原因是寄存器资源不足： 设备：GTX1050Ti实验代码/** test the r...

cuda程序的组织（与C++融合）前言：我们都希望有一个好的cuda程序组织方式，让我们的代码或程序变得干净整洁，不容易出错，可以移植性能强。今天介绍一种最基础CUDA程序组织方式。我们都知道面向对象程序的好处，或者说，我们现在有一个C++程序需要并行化。我们不得不让CUDA程序与C++相融合，但cuda是基于c语言的，怎样组织cuda程序就尤为关键了。问题：由于有关CUDA的...

CUDA优化实例（三）共享内存前言经过前面的实验发现，共享内存是优化CUDA程序的核心方法。共享内存可以通过对全局内存数据进行合并访问，让kernel内交错的内存需求去访问共享内存。如：矩阵转置问题，将二维内存的行写入二维内存的列。对列的写入就是一个内存交错访问的例子。可以用合并的方式将块要操作的数据写入共享内存，让复杂的内存交错访问访问共享内存，然后将结果以合并的方式写入全局内存。 ...

有关CUDA nvprof 调试的metrics(指标)nvprof --metrics achieved_occupancy,gld_throughput,gst_throughput,gld_efficiency,gst_efficiency,gld_transactions,gst_transactions,gld_transactions_per_request,gst_transac...

笔记–关于CUDAWhy GPU is faster than CPU?One well known reason is that GPU has more threads and it can parallel process rather than serial. Another significant reason is the GPU programming Model – ...

CMake构建OpenCV，GDAL项目本文举例的开源库是OpenCV和GDAL，其他块源库类似使用本篇方法的前提是make install，即开源库是install的。 当然，没有install的也是可以用的，只需要在findpage前找到DIR即可。 如找OpenCV： set(OpenCV_DIR /home/jie/third_party/opencv-3.4.0/build...

CUDA优化实例（二）对齐与合并本篇主要通过反正的方式进行实验，即说明不合并的内存访问方式慢，以此来说明对全局内存的访问一定要保证合并。引言关于全局内存的对齐与合并问题，前面的文章1 前面的文章2也介绍了，我在做有关对齐的试验时发现许多不可解释的问题，主要是对齐的问题，我发现这与我在书中学的不一样，为此我去官方文档中寻找线索，发现，现在的GPU对非对齐问题都进行了优化，不管对齐不...

CUDA优化实例（四）纹理内存本文参考：http://www.cnblogs.com/cuancuancuanhao/p/7809713.html 本节的内容可能和标题不服，本节主要将纹理内存的使用，它到底有什么有什么速度的提升，优化体现在哪里，我下节会写，本节主要写一个纹理内存的例子。引言纹理内存与全局内存一样，都在DRAM上，所以纹理内存的容量是很大的纹理内存是专为图像而...

有关展开的问题理论：在相同情况下，内存效率与内存事物数成反比 低的内存效率意味着需要更多的内存事物。 在实验中可以证实这一点，当平移数据字节时（非对齐访问），内存效率会下降，且它们乘积是定值都是处理数据所需要的内存事物数。问题：但在实验中也发现问题： 展开是 1unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;...

CUDA编译（二）—用CMake混合编译C++与cuda引言 示例引言许多c/c++的项目会使用cuda加速其算法，c/c++有其编译器：gcc/g++，cuda有其编译器nvcc。为了实现我们的目的，我们一般采用gcc/g++编译c/c++部分代码，nvcc编译cuda代码部分，即分离式编译。如此博文nvcc gcc g++混合编译器编程。 分离式编译的基本思路： 将c...

CMake编译CUDA首先要有CUDA环境。 下面简单介绍一下CMakeLists.txt：# CMakeLists.txt for G4CU projectproject(project)# required cmake versioncmake_minimum_required(VERSION 2.8)# packagesfind_package(CUDA)...

内存事物颗粒内存事物数×颗粒大小×效率 = 处理数据量每个请求数量的事物数量×颗粒大小×效率 = 一个请求处理的事物量事物量/每个请求事物量 = 请求数GTX1050Ti经实验得： 以此求GTX1050Ti的内存事物颗粒大小： 加载颗粒： 在请求内存量<32字节时，加载颗粒是32字节。 在请求内存量>32字节时，加载颗粒是16字节。 存储颗粒： 无论...

nvcc编译器参数参考： http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html https://www.cnblogs.com/shrimp-can/p/5567518.htmnvcc编译器参数示例nvcc编译器参数nvcc是cuda程序的编译器，了解它的一些关键参数有利于我们更加从容，更加准...

CUDA优化The method of CUDA to improve performance:1.block size to increase occupancy 2.对其和合并 increase the 内存事物使用率. 3.减少分支化 4.展开 unrolling（一般是最有效的方法） 5.尽量使用共享内存，但要避免共享内存的冲突（正确使用共享内存）可扩展性...

CUDA中内存访问越界问题在核函数中对非法的内存访问不一定会报错，那我们加CHECK（error）会帮我们报错吗？经实验发现，有的非法访问很致命，函数不会运行，这种错误，可以CHECK到，但有些非法内存访问不致命，可能是访问到其他设备内存上了，这时CHECK不会报错，但是会出现逻辑上的错误，因为访问的值是错误的。 所以，我们在做内存访问时，一定要自己做好逻辑上的判断，保证内存访问不错位，不错...

使用Padding（cudaMallocPitch）的二维数组有关为什么使用Padding的二维数组，在此篇已说明（为了对齐访问） 本篇就已实验来验证，Padding的必要性。 代码附在后面。前言本文的内容：介绍CUDA APIcudaMallocPitch和cudaMemcpy2D。实例代码实现cudaMallocPitch和cudaMemcpy2D。实验验证P...