精华内容
下载资源
问答
  • CUDA原理+实战

    2018-09-10 10:49:13
    CUDA原理+实战,快速上手CUDA必备
  • NVIDIA CUDA原理和基础知识

    千次阅读 2019-01-21 17:45:08
    为什么现在使用GPU(Graphics Processing Unit)编程越来越多,主要是因为GPU相对于CPU的运算速度,内存带宽均有较大的优势,下面是摘自《CUDA C PROGRAMMING GUIDE》中的图片: 浮点数运算速度: 内存带宽: ...

    1. 为什么需要使用GPU

    为什么GPU(Graphics Processing Unit)编程越来越流行,主要是因为GPU相对于CPU的运算速度,内存带宽均有较大的优势,下面是摘自《CUDA C PROGRAMMING GUIDE》中的图片:

    浮点数运算速度:
    在这里插入图片描述
    内存带宽:
    在这里插入图片描述

    2. GPU为什么性能高

    这是因为GPU中硬件更多的用于data processing而不是data caching 或 flow control
    在这里插入图片描述
    NVIDIA GPU 更是采用了SIMT (Single-Instruction, Multiple-Thread)和Hardware Multithreading 技术来进行计算加速:

    • SIMT 相对于SIMD(Single Instruction, Multiple Data),前者主要采用线程并行的方式,后者主要采用数据并行的方式。
      下面是一个采用SIMD进行运算的例子:

      void add(uint32_t *a, uint32_t *b, uint32_t *c, int n) {
        for(int i=0; i<n; i+=4) {
          //compute c[i], c[i+1], c[i+2], c[i+3]
          uint32x4_t a4 = vld1q_u32(a+i);
          uint32x4_t b4 = vld1q_u32(b+i);
          uint32x4_t c4 = vaddq_u32(a4,b4);
          vst1q_u32(c+i,c4);
        }
      }
      

      下面是一个SIMT的例子:

      __global__ void add(float *a, float *b, float *c) {
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        a[i]=b[i]+c[i]; //no loop!
      }
      
    • Hardware Multithreading技术主要是将进程的运行上下文一直保存在硬件上,因而不存在运行上下文切换带来开销的问题(传统的CPU多进程是将进程运行上下文保存在内存中,进程切换时涉及到内存的读取,因而开销较大)

    3. 如何运用GPU进行编程

    既然GPU有这么多的优势,那么如何使用GPU进行编程呢?由于GPU种类很多,不同的GPU都有不同的硬件实现以及相应的软件接口。目前比较流行的是NVIDIA GPU, 这主要是因为其提供了一套易用的软件接口CUDA, CUDA(Compute Unified Device Architecture)是NVIDIA公司基于其生产的图形处理器GPU开发的一个并行计算平台和编程模型。

    在这里插入图片描述

    3.1 NVIDIA GPU Architecture

    NVIDIA GPU的硬件架构一般如下,以GeForce8600 为例:
    在这里插入图片描述
    每个GPU中都有多个多流处理器Streaming Multiprocessors(简称SM,有时也直接叫做Multiprocessor), 每个Multiprocessors中有多个core,线程最终就是在这些core上运行的。
    这些硬件信息可以通过CUDA Runtime API 获取,例如,我的Lenovo T440P上的GPU硬件信息如下:

     CUDA Device Query (Runtime API) version (CUDART static linking)
    
    Detected 1 CUDA Capable device(s)
    
    Device 0: "GeForce GT 730M"
      CUDA Driver Version / Runtime Version          10.0 / 10.0
      CUDA Capability Major/Minor version number:    3.5
      Total amount of global memory:                 984 MBytes (1031405568 bytes)
      ( 2) Multiprocessors, (192) CUDA Cores/MP:     384 CUDA Cores
      GPU Max Clock rate:                            758 MHz (0.76 GHz)
      Memory Clock rate:                             1001 Mhz
      Memory Bus Width:                              64-bit
      L2 Cache Size:                                 524288 bytes
      Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
      Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
      Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
      Total amount of constant memory:               65536 bytes
      Total amount of shared memory per block:       49152 bytes
      Total number of registers available per block: 65536
      Warp size:                                     32
      Maximum number of threads per multiprocessor:  2048
      Maximum number of threads per block:           1024
      Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
      Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
      Maximum memory pitch:                          2147483647 bytes
      Texture alignment:                             512 bytes
      Concurrent copy and kernel execution:          Yes with 1 copy engine(s)
      Run time limit on kernels:                     Yes
      Integrated GPU sharing Host Memory:            No
      Support host page-locked memory mapping:       Yes
      Alignment requirement for Surfaces:            Yes
      Device has ECC support:                        Disabled
      Device supports Unified Addressing (UVA):      Yes
      Device supports Compute Preemption:            No
      Supports Cooperative Kernel Launch:            No
      Supports MultiDevice Co-op Kernel Launch:      No
      Device PCI Domain ID / Bus ID / location ID:   0 / 2 / 0
      Compute Mode:
         < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
    
    deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 10.0, CUDA Runtime Version = 10.0, NumDevs = 1
    Result = PASS
    

    该GPU有2个Multiprocessor, 每个multiprocessor有192个core,总计384个core. 对于现在的Tesla型号的GPU,其core数为3584(56 * 64), 每个core都有其相对独立的寄存器等,这是GPU高性能的基础。

    3.2 Thread Hierarchy

    在NVIDIA GPU编程中,一个多线程的程序会采用分组的方式在GPU上运行,每个组称为一个block,每个block中含有若干个线程。每个thread block在一个Multiprocessor上运行;多个thread blocks可以在一个或多个Multiprocessor上运行。这样做的好处是当增加GPU中Multiprocessor的个数时,程序性能可以随之提高。

    在这里插入图片描述

    Block在Grid中的排列形式可以是1D或2D(没有3D的block),每个block中有若干线程,这些线程在block中的排列方式可以是1D/2D/3D,如下图:

    在这里插入图片描述
    在GPU编程中,相应的概念均可以找到具体的物理实体:

    • Grid 对应于GPU,一个GPU就是一个Grid,在多GPU的机器上,将会有多个Grid。
    • Block对应(从属于)MultiProcessors这个物理实体
    • Thread对应于MultiProcessors下面的core这个物理实体,thread 运行在core上

    具体的,当一个block运行在multiprocessor时,multiprocessor是以wrap为单位来调度block中的线程的,一个wrap一般是32个线程,这也就是我们为什么说NVIDIA GPU采用SIMT的原因。wrap是来源于实际生活中的概念(织布中用的经,经纱),下图中所有的竖线即为一个wrap:

    在这里插入图片描述
    对应于上面硬件GeForce GT 730M,其线程相关参数如下:

    • 每个Multiprocessor 最多可支持2048个线程;
    • 每个thread block中最多可支持1024个线程;
    • 每个thread block中维数方面x,y,z分别最多为1024,1024,64
    • 每个grid中维数方面x,y,z分别最多为2147483647, 65535, 65535

    Maximum number of threads per multiprocessor: 2048
    Maximum number of threads per block: 1024
    Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
    Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)

    注意这里grid size的z方向虽然最大可以是65535,但是在CUDA的实际编程接口中只能是1.

    3.3 Execution Model

    采用CUDA编程时,程序的运行步骤一般如下:

    1.准备GPU计算数据: 将数据从host内存拷贝到GPU内存
    2.在GPU中运行程序
    3.将计算结果从GPU内存拷贝到CPU内存

    在这里插入图片描述在GPU和CPU混合编程中,通常将GPU叫做device, 将CPU叫做host。如上步骤2中能够在Host端被调用,在device端执行的函数叫kernel function。

    3.4 kernel function

    对于运行在device端的函数,一般以__global__和以__device__ 作为标志。以__device__作为标志的函数只能在device上被调用;以__global__作为标志的函数可以在host端调用,也可以在device端调用,一般称为kernel function, 调用kernel function时我们需要提供两个参数:

    1. 以block为单位的,在grid内部block在x,y方向(不支持z方向)的维数B
    2. 以thread为单位的,在block内线程在x,y,z方向的维数T

    kernel function调用的一般形式为:

    myKernel<<< B, T >>>(arg1, … );
    

    B,T在CUDA中采用如下类似的数据结构dim3:

    struct dim3 {x; y; z;};
    

    其提供了int到dim3的隐式类型转换:

    myKernel<<< 2, 3 >>>(arg1, … );
    

    上面的参数等价于dim3 b(2,1,1) T(3,1,1)。CUDA为所有在device内运行的function提供了如下两个内置变量gridDim和blockDim:

    dim3 gridDim
    dim3 blockDim
    
    • 通过gridDim.x,gridDim.y,gridDim.z,获取grid在x,y,z方向的维数,也就是block在grid内部x,y,z方向的个数,gridDim.z始终为1
    • 通过blockDim.x,blockDim.y,blockDim.z,获取block在x,y,z方向的维数,也就是线程在block内部x,y,z方向的个数

    那么程序中使用到的block数和单个block内部线程总数将分别是:

    gridDim.x * gridDim.y*gridDim.z // girdDim.z = 1
    blockDim.x * blockDim.y * blockDim.z
    

    对于kernel function的调用,采用的是SIMT的方式,也就是说同一个function的函数指令将会运行在多个线程中,而线程又属于某个block,我们怎么获取这些线程的索引(index)呢? CUDA 提供了两个可以在kernel function内部使用的变量:

    uint3 blockIdx
    uint3 threadIdx
    
    • 通过blockIdx.x, blockIdx.y获取到当前block在grid内部x,y方向的索引
    • 通过threadIdx.x, threadIdx.y, threadIdx.z获取thread在block内部x,y,z方向的索引

    对于2D Grid和2D block,线程在x,y方向的全局唯一ID就可以通过如下计算得到:

    • x = blockIdx.x * blockDim.x + threadIdx.x;
    • y = blockIdx.y * blockDim.y + threadIdx.y;

    下面是一个2D Grid和2D block的示意图,:
    在这里插入图片描述
    对于2D Grid和3D block的情形,有类似:

    • x = blockIdx.x * blockDim.x + threadIdx.x;
    • y = blockIdx.y * blockDim.y + threadIdx.y;
    • z = blockIdx.z * blockDim.z + threadIdx.z;

    注意前面提到过Grid的排列形式没有3D的,只有2D的,也就是说blockIdx.z = 0;

    4. An example: Matrix Multiplication

    下面通过矩阵相乘的例子来说明采用如何使用GPU进行编程,回忆一下,对于矩阵A,B,矩阵向乘的结果C中的元素是通过如下公式得到:

    在这里插入图片描述
    具体计算过程如下:

    在这里插入图片描述
    在C中,一般的实现如下:

    void matrixMult (int a[N][N], int b[N][N], int c[N][N], int width)
    {
    	for (int i = 0; i < width; i++) {
    		for (int j = 0; j < width; j++) {
    			int sum = 0;
    			for (int k = 0; k < width; k++) {
    				int m = a[i][k];
    				int n = b[k][j];
    				sum += m * n;
    			}
    			c[i][j] = sum;
    		}
    	}
    }
    

    其中,矩阵width是矩阵A的列数,显然,上面算法的复杂度是O(N^3)。采用GPU编程只需将上面的方法写成kernel function的形式:

    __global__ void matrixMult (int *a, int *b, int *c, int width) {
    	int k, sum = 0;
    	int col = threadIdx.x + blockDim.x * blockIdx.x;
    	int row = threadIdx.y + blockDim.y * blockIdx.y;
    	if(col < width && row < width) {
    		for (k = 0; k < width; k++) {
    			sum += a[row * width + k] * b[k * width + col];
    		}
    		c[row * width + col] = sum;
    	}
    }
    

    对比一下C和GPU实现的线程数量和时间复杂度:

    线程数量时间复杂度
    C1N^3
    GPUN^2N

    较完整的GPU实现代码如下:

    #define N 16
    #include <stdio.h>
    __global__ void matrixMult (int *a, int *b, int *c, int width) {
    	int col = threadIdx.x + blockDim.x * blockIdx.x;
    	int row = threadIdx.y + blockDim.y * blockIdx.y;
    	if(col < width && row < width) {
    		for (k = 0; k < width; k++) {
    			sum += a[row * width + k] * b[k * width + col];
    		}
    		c[row * width + col] = sum;
    }
    
    int main()  {
    	int a[N][N], b[N][N], c[N][N];
    	int *dev_a, *dev_b, *dev_c;
    	// initialize matrices a and b with appropriate values
    	int size = N * N * sizeof(int);
    	cudaMalloc((void **) &dev_a, size);
    	cudaMalloc((void **) &dev_b, size);
    	cudaMalloc((void **) &dev_c, size);
    	cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice);
    	cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice);
    	dim3 dimGrid(1, 1);
    	dim3 dimBlock(N, N);
    	matrixMult<<<dimGrid, dimBlock>>>(dev_a, dev_b, dev_c, N);
    	cudaMemcpy(c, dev_c, size, cudaMemcpyDeviceToHost);
    	cudaFree(dev_a); 
    	cudaFree(dev_b); 
    	cudaFree(dev_c);
    
    }
    
    展开全文
  • 由于很多同学需要使用Windows系统进行GPU上的计算,比如运行TensorFlow或者Pytorch等,在Windows上正确安装CUDA和CUDNN则成了一个关键的问题。与在Linux上安装的整体流程类似,但是还有一些细节要注意,本文以...
    由于很多同学需要使用Windows 系统进行GPU 上的计算,比如运行TensorFlow 或者Pytorch 等,在Windows 上正确安装CUDA 和CUDNN 则成了一个关键的问题。与在Linux 上安装的整体流程类似,但是还有一些细节要注意,本文以Windows 10 系统为例,进行CUDA 环境的安装,并以TensorFlow 验证安装的成功。

    1 安装最新版Microsoft Visual Studio

    看到这里很多人可能会问,就写个Python 程序为什么要装vs 软件。其实这是因为,我们需要给我们的计算机安装Windows 系统上包括最新版在内的几个版本的C++ 编译器和运行环境。不同于Linux 系统,一个崭新的Windows 系统并不包含这些软件,所以很多安装TensorFlow GPU 版的同学发现安装之后,运行时报错:“ImportError: DLL loadfailed: 找不到指定的程序”。我们从TensorFlow 官方的文档 ( https://www.tensorflow.org/install/source_windows) 中就可以看到对编译器版本的要求这一点:

    72679ec4c29359c7cb8fe348fa7303a6.png

    而这个问题正是网上很多教程不曾写到的,也是最令人困扰的一点。

    我们需要打开微软VisualStudio的官方网站:

    https://visualstudio.microsoft.com/zh-hans/  

    下载并安装最新版VisualStudio,安装时必须勾选“使用C++的桌面开发”,其他选项可根据自己实际需要勾选。

    1380eb12a6126c4601b73e8547426336.png

    2 查看本机的CUDA驱动适配版本

    在桌面右键“NVIDIA 控制面板”,点击帮助-> 系统信息-> 组件。

    a92dcfef07d59e4486082e48ca6089ac.png

    在打开的窗口中,我们可以看到本机当前最高支持的CUDA 版本是10.2 。如果你升级了驱动,将来也可能会支持更高版本。

    3 下载CUDAcuDNN

    CUDA 下载页面:https://developer.nvidia.com/cuda-downloads如果需要选择CUDA版本,可从这里打开:https://developer.nvidia.com/cuda-toolkit-archive

    b8f5496a806746a60e0d5a557a119d17.png

    下载cuDNN 时必须选择和你安装的CUDA 匹配的版本,下载页面:https://developer.nvidia.com/rdp/cudnn-download 下载cuDNN 是需要登录英伟达开发者账户的,如果没有的话,需要注册一个并填写问卷,很简单。注册并登录后,即可打开如下页面,选择对应的文件并下载。

    cde42d730c546c997114907c53d7227c.png

    4 安装CUDAcuDNN

    安装CUDA 时很简单,跟我们平时装一个软件一样,这里不再赘述。CUDA 安装完成后,打开powershell ,执行nvcc -V ,成功的话会返回cuda 版本号。 安装cuDNN 首先需要解压cuDNN 压缩包,可以看到有bin 、include 、lib 目录。

    2aeea3a0bfd21567261d6d9baa92f512.png

    打开 “C:\Program Files\NVIDIAGPU Computing Toolkit\CUDA ” 目录,找到你安装的版本目录,打开,找到bin 、include 、lib 目录,将cuDNN 压缩包内对应目录下的文件分别复制到bin 、include 、lib 目录。如果是新版Windows 10 系统,可以直接复制三个目录到对应路径下,两处的目录会自动合并,不会出现目录级的覆盖,最多只会覆盖同名文件。如果是其他版本系统,为了以防万一,最好还是手动一一复制到对应路径下。

    5 添加环境变量

    我们需要在系统环境变量的Path 项下添加几个路径,点击 编辑 -- > 新建、浏览,选择刚才的安装路径。

    需要添加的默认的安装路径如下:

    C:\ProgramFiles\NVIDIA GPU Computing Toolkit\CUDA\v10.2C:\ProgramFiles\NVIDIA GPU Computing Toolkit\CUDA\v10.2\lib\x64

    如果有安装到别处,请参考这两个路径来填写。

    6 检查安装结果

    在任意路径下打开powershell 终端,比如在桌面处按住Shift 键,鼠标右键,选择“在此处打开Powershell 窗口”。然后输入:“nvidia-smi ”,如果出现类似如下信息,则表明CUDA 安装成功。

    be0cacb263f1adcf967b22b8cf59213a.png

    然后我们运行python ,并”import tensorflow as tf” ,输入”tf.test.is_gpu_available()” ,然后回车,如果出现”True” 字样,则说明完全安装成功。

    e5a1f765161a521f2a0dea3017cd60f0.png


    往期精彩

    科技力量是中国抗疫制胜的法宝

    基于动态软件体系结构的插件机制原理(C语言实现)

    ASRT开源语音识别项目开始兼容TensorFlow 2.X

    [论文翻译]SpecAugment:一种用于自动语音识别的简单数据增强方法

    手把手教你从零开始搭建个人博客

    ·END·

    AI柠檬博客

    分享AI技术文章

    关注前沿科技

    QQ交流群:  867888133

    8748197f5aac09299d174528cb764ee7.png

    微信号:ailemon_me

    博客网址:https://blog.ailemon.me

    点击下方“ ”即可阅读博客原文,并参与评论和互动。

    点在看,让柠檬菌知道你在看~

    展开全文
  • CUDA基本原理及概念

    千次阅读 2017-12-20 14:28:45
    CUDA基本原理及概念本文主要包含如下内容:CUDA基本原理及概念 Introduction GPU Programming Languages Multidimensional Grids ExampleRGB to Grayscale Conversion Code ExampleImage Blurring RegistersShared ...

    CUDA基本原理及概念


    本文主要包含如下内容:


    Introduction


    Scalability and Portability 可扩展性和可移植性

      相同的应用程序可以在一个内核上有效运行
      相同的应用程序可以在更多的相同内核上有效运行
    Scalability

      相同的应用程序可以在不同类型的内核上有效运行
      相同的应用程序可以在具有不同组织和接口的系统上高效运行
    Portability


    GPU Programming Languages


      CUDA内核由线程网格(Grid)组成;网格中的所有线程都运行相同的内核代码(单程序多个数据);每个线程都有用于计算内存地址并进行控制决策的索引。

    Programming Languages

    Block:

      block可以是一维、二维、三维的,每个block可以有1到512个threads,block中所有的threads将执行相同的线程程序。Maximum number of threads per block: 512

    Warps

      现在的主流Streaming Multiprocessor都是基于32threads的,每一个Thread Blocks被分为32个thread Warps,即32个threads被分为1个Warps

    Programming Languages
      将线程数组划分成多个块,块内的线程通过共享内存,不同块中的线程不进行交互

        // Compute vector sum C = A + B
        void vecAdd(float *h_A, float *h_B, float *h_C, int n)
        {
            int i;
            for (i = 0; i<n; i++) h_C[i] = h_A[i] + h_B[i];
        }
    
        int main()
        {
            // Memory allocation for h_A, h_B, and h_C
            // I/O to read h_A and h_B, N elements
              …
            vecAdd(h_A, h_B, h_C, N);
        }

      上面这个程序是在CPU上运行的求和程序,接下来,调用GPU编程。

    cudaMalloc():在GPU的全局内存中分配对象,即分配一个和原始数据一样大的内存,方便读取与存储。

    • 两个参数;指向分配对象的指针的地址;分配对象的大小(以字节计)

    cudaFree():从全局内存释放对象

    • 一个参数:指向释放对象的指针

    cudaMemcpy():内存数据传输

    • 四个参数:指向目的地;指向源的指针;复制的字节数;类型/转移方向

    cudaThreadSynchronize(); 同步,等待GPU内核执行结束
    __global__:GPU kernel 代码,由CPU发起,返回void,不能由其他kernel调用
    __device__:由GPU kernel调用的函数,不能由其他kernel调用
    __host__:在CPU上执行的函数
    __shared__:变量位于共享存储器

        // vecAdd CUDA Host Code
        #include <cuda.h>
        void vecAdd(float *h_A, float *h_B, float *h_C, int n)
        {
           int size = n * sizeof(float);
           float *d_A, *d_B, *d_C;
           // Part 1
           // Allocate device memory for A, B, and C
           // copy A and B to device memory 
            cudaMalloc((void **) &d_A, size);   //在全局内存中分配对象
            cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);    //内存数据传输
            cudaMalloc((void **) &d_B, size);
            cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
            cudaMalloc((void **) &d_C, size);
    
           // Part 2
           // Kernel launch code – the device performs the actual vector addition
            dim3 DimGrid((n-1)/256 + 1, 1, 1);  //总共3维,这里使用了一维
            dim3 DimBlock(256, 1, 1);   //每个block有256个线程
            vecAddKernel<<<DimGrid,DimBlock>>>(d_A, d_B, d_C, n); //表明内核的参数
    
    
            cudaThreadSynchronize();    //同步,等待执行结束
           // Part 3
           // copy C from the device memory
           // Free device vectors
           cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);  //内存数据传输
          cudaFree(d_A); cudaFree(d_B); cudaFree (d_C);    //从全局内存释放对象
    
        }
    
        __global__      //定义核心函数
        void vecAddKernel(float* A, float* B, float* C, int n)  //Device code必须返回void 
        {
            int i = threadIdx.x+blockDim.x*blockIdx.x;  //查找指针真正的位置,即索引
            if(i<n) C[i] = A[i] + B[i];    //条件语句防止数据溢出,最后一个block可能存在空值
        }

    Multidimensional Grids


    Multidimensional Grids

      注意:在Device中横向是x,纵向是y,在计算行和列的时候注意变化关系

        ...
        dim3 DimGrid((n-1)/16+1,(m-1)/16+d1,1);
        dim3 DimBlock(16,16,1);
        PictureKernel<<<DimGird,DimBlock>>>(d_Pin, d_Pout, m, n)
        ...
    
        _global__
        void PictureKernel(float* d_Pin, float* d_Pout, int height, int width)
        {
          // Calculate the row # of the d_Pin and d_Pout element
          int Row = blockIdx.y*blockDim.y + threadIdx.y;    //获取原始图像的位置
    
          // Calculate the column # of the d_Pin and d_Pout element
          int Col = blockIdx.x*blockDim.x + threadIdx.x;    //获取原始图像的位置
    
          // each thread computes one element of d_Pout if in range
          if ((Row < height) && (Col < width)) {
            d_Pout[Row*width+Col] = 2.0*d_Pin[Row*width+Col];   //转化为一维坐标
          }
        }

    Example:RGB to Grayscale Conversion Code


        #define CHANNELS 3 // we have 3 channels corresponding to RGB
        // The input image is encoded as unsigned characters [0, 255]
        __global__ void colorConvert(unsigned char * grayImage, unsigned char * rgbImage, int width, int height) 
        {
         int x = threadIdx.x + blockIdx.x * blockDim.x;
         int y = threadIdx.y + blockIdx.y * blockDim.y;
    
         if (x < width && y < height) 
         {
            // get 1D coordinate for the grayscale image
            int grayOffset = y*width + x;
            // one can think of the RGB image having
            // CHANNEL times columns than the gray scale image
            int rgbOffset = grayOffset*CHANNELS;
            unsigned char r = rgbImage[rgbOffset    ]; // red value for pixel
            unsigned char g = rgbImage[rgbOffset + 2]; // green value for pixel
            unsigned char b = rgbImage[rgbOffset + 3]; // blue value for pixel
            // perform the rescaling and store it
            // We multiply by floating point constants
            grayImage[grayOffset] = 0.21f*r + 0.71f*g + 0.07f*b;
         }
        }

    Example:Image Blurring


         __global__ 
          void blurKernel(unsigned char * in, unsigned char * out, int w, int h) 
          {
              int Col  = blockIdx.x * blockDim.x + threadIdx.x;
              int Row  = blockIdx.y * blockDim.y + threadIdx.y;
    
              if (Col < w && Row < h) 
              {
                  int pixVal = 0;
                  int pixels = 0;
    
                  // Get the average of the surrounding 2xBLUR_SIZE x 2xBLUR_SIZE box
                  for(int blurRow = -BLUR_SIZE; blurRow < BLUR_SIZE+1; ++blurRow) 
                  {
                      for(int blurCol = -BLUR_SIZE; blurCol < BLUR_SIZE+1; ++blurCol) 
                      {
                          int curRow = Row + blurRow;
                          int curCol = Col + blurCol;
                          // Verify we have a valid image pixel
                          if(curRow > -1 && curRow < h && curCol > -1 && curCol < w) 
                          {
                              pixVal += in[curRow * w + curCol];
                              pixels++; // Keep track of number of pixels in the accumulated total
                          }
                      }
                  }
                  // Write our new pixel value out
                  out[Row * w + Col] = (unsigned char)(pixVal / pixels);
              }
          }

    Registers、Shared memory、Global memory


      学会合理分配CUDA内存,可以加快数据访问速度,从而加快程序运行速度。
    Registers、Shared memory、Global memory

    Registers

      存储在寄存器存储器中的数据只对写入它的线程可见,并且只持续该线程的生命周期。

    Global memory

      对应用程序内的所有threads都是可见的,持续时间为主机分配。全局内存较大,速度较慢。

    Shared memory

      对每一个block内的所有线程都是可见的,持续时间为block的持续时间。允许线程在彼此之间进行通信和共享数据。速度非常快,如果没有冲突,速度和Registers一样快,但是最大大小仅为48KB。适合情况:在数据需要重复读取时,若数据只用读取一次,则Shared memory并无实质性的提高。

    CUDA Memories

      Registers中存储的是自动变量,数据一般可以存储在shared memory中,比全局内存具有更高速度(延迟和吞吐量)的访问,访问和共享的范围:线程块
    生命周期 :线程块,相应的线程终止执行后,内容将消失。

        //A Basic Matrix Multiplication
        __global__ void MatrixMulKernel(float* M, float* N, float* P, int Width) 
        {
          // Calculate the row index of the P element and M
          int Row = blockIdx.y*blockDim.y+threadIdx.y;
    
          // Calculate the column index of P and N
          int Col = blockIdx.x*blockDim.x+threadIdx.x;
    
          if ((Row < Width) && (Col < Width)) 
          {
            float Pvalue = 0;
            // each thread computes one element of the block sub-matrix
            for (int k = 0; k < Width; ++k) {
              Pvalue += M[Row*Width+k]*N[k*Width+Col];
            }
            P[Row*Width+Col] = Pvalue;
          }
        }

      以上程序访问的是Global memory,速度较慢

    Tiled Matrix Multiplication

      已知矩阵乘法是由一行乘以一列完成的,由于shared memory大小有限,需要将每个阶段的线程块的数据访问集分解为在M的一个tile和N的一个tile,其中,每个tile含有BLOCK_SIZE个元素

    __syncthreads():同步所有的操作,即块内所有的内存,用以防止RAW/WAR/WAW危害

        __global__ void MatrixMulKernel(float* M, float* N, float* P, Int Width)
        {
          __shared__ float ds_M[TILE_WIDTH][TILE_WIDTH];
          __shared__ float ds_N[TILE_WIDTH][TILE_WIDTH];
    
          int bx = blockIdx.x;  int by = blockIdx.y;
          int tx = threadIdx.x; int ty = threadIdx.y;
    
          int Row = by * blockDim.y + ty;   //查找指针真正的位置,即索引
          int Col = bx * blockDim.x + tx;
          float Pvalue = 0;
    
         // Loop over the M and N tiles required to compute the P element
         for (int p = 0; p < n/TILE_WIDTH; ++p) 
         {
            // Collaborative loading of M and N tiles into shared memory
            ds_M[ty][tx] = M[Row*Width + p*TILE_WIDTH + tx];    //for_loop 循环加载shared_memory
            ds_N[ty][tx] = N[(t*TILE_WIDTH+ty)*Width + Col];
            __syncthreads();    //确保上一步操作完成
    
            for (int i = 0; i < TILE_WIDTH; ++i)Pvalue += ds_M[ty][i] * ds_N[i][tx];    //矩阵乘法
            __synchthreads();   //确保上一步操作完成
    
          } 
          C[Row*Width+Col] = Cvalue;
        }

      本段程序仅仅考虑Shared memory的应用


    Example:1D Stencil


      一维数组元素,每个输出元素是半径内的输入元素的总和,半径为3。

        __global__ void stencil_1d(int *in, int *out) 
        {
            __shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
            int gindex = threadIdx.x + blockIdx.x * blockDim.x + RADIUS; //查找指针真正的位置,即索引
            int lindex = threadIdx.x + RADIUS;  
    
            // Read input elements into shared memory
            temp[lindex] = in[gindex];
            if (threadIdx.x < RADIUS) 
            {
                temp[lindex – RADIUS] = in[gindex – RADIUS];    
                temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
            }
    
            // Synchronize (ensure all the data is available)
            __syncthreads();    //确保上一步操作完成
    
            // Apply the stencil
            int result = 0;
            for (int offset = -RADIUS ; offset <= RADIUS ; offset++)
                result += temp[lindex + offset];
    
            // Store the result
            out[gindex - RADIUS] = result;   //存储结果
        }

      考虑到该函数和周边像素求和,定义了RAIDUS,注意该种定义对源程序的影响。


    Example: edge detection algorithm


      下面一段程序为host调用GPU内核函数,故函数声明为__global__,并且涉及二维操作,将blockthread定义为二维的。

            dim3 blocksPerGrid ( N/TILE_WIDTH , N/TILE_WIDTH , 1);  //共三维,这里使用了两维
            dim3 threadsPerBlock(TILE_WIDTH,TILE_WIDTH,1);  //同样的,thread定义为两维
            inverseEdgeDetect2D<<< blocksPerGrid, threadsPerBlock >>>(d_output, d_input, d_edge);    //定义GPU内核函数的参数

      GPU内核函数,由于是二维的,注意Shared memory的使用,并且考虑到该函数和周边像素做差分,定义了RAIDUS,注意该种定义对源程序的影响。注意一维和二维的联系与区别

        __global__ void inverseEdgeDetect2D(float *d_output, float *d_input, \
                            float *d_edge)
        {
            __shared__ float ds_input[TILE_WIDTH+ 2*RAIDUS][TILE_WIDTH+2*RAIDUS];
    
            int bx = blockIdx.x;  int by = blockIdx.y;
            int tx = threadIdx.x + RAIDUS; int ty = threadIdx.y + RAIDUS;
            int Row = by * blockDim.y + ty ; 
            int Col = bx * blockDim.x + tx ;
            int idx;
            int numcols = N + 2;
    
            // Read input elements into shared memory
            ds_input[ty][tx] = d_input[Row*numcols + Col];
            if (threadIdx.x < RAIDUS) 
                {
            ds_input[ty][tx - RAIDUS] = d_input[Row*numcols + Col- RAIDUS];    
                ds_input[ty][tx + TILE_WIDTH] = d_input[Row*numcols + Col + TILE_WIDTH];
               }
            if (threadIdx.y < RAIDUS) 
            {
            ds_input[ty- RAIDUS][tx] = d_input[(Row- RAIDUS)*numcols + Col];    
                ds_input[ty + TILE_WIDTH][tx] = d_input[(Row + TILE_WIDTH)*numcols + Col];
            }
    
            // Synchronize (ensure all the data is available)
            __syncthreads();  
    
            // Synchronize (ensure all the data is available)
            idx = Row * numcols + Col;
            d_output[idx]=(ds_input[ty+RAIDUS][tx] + ds_input[ty-RAIDUS][tx] \
                        + ds_input[ty][tx+RAIDUS] + ds_input[ty][tx-RAIDUS] - \
                        d_edge[idx]) * 0.25;
        } 

    Memory Coalescing


      在访问全局内存时,当所有线程连续访问数据的内存时,会有峰值的性能利用率。因此,要尽可能连续访问内存。
    图片示意

    连续存储

    非连续存储

    展开全文
  • 中标麒麟系统成功安装英伟达显卡后,一颗不安定的心觉得仅仅装一个显卡驱动,没有什么大的用处,还是要把CUDA装备好,可以使用GPU加速才是正确的道路啊!!!!


    前言

      在中标麒麟系统成功安装英伟达显卡驱动程序后,心血来潮冲昏了冷静的头脑,一个显卡驱动程序根本没有什么大的用处,还是要把CUDA装备好,让OpenCV或者YoloV3能够使用GPU加速进行图像处理才是正确的道路,才能遛到飞起!!!!!!
    软硬件配置:Intel i7,英伟达显卡,CUDA10.0,cudnn7.6.5


    YOLOV3分别调用CPU和GPU的检测时间对比,测试图像尺寸1024*1024。
    CPU用时11s。
    在这里插入图片描述
    GPU用时1.5s,显卡为GT760,算力在2.x,算力更高的显卡可达毫秒级。
    在这里插入图片描述
    你就说香不香~~~~~~~~~~~~~

    一、安装CUDA

      如果已经安装了显卡驱动,有的系统可以在终端窗口使用nvidia-smi命令显示显卡信息,在此处就可以看到相应的cuda版本,如下图所示。
    在这里插入图片描述
      如果没有安装显卡驱动程序或者终端输入nvidia-smi命令后也没有显示cuda版本的话,此时可以在终端窗口使用uname -a命令查看中标麒麟系统的内核版本,根据该版本的发布时间选择cuda的版本,如下图所示,该版本发布时间在2017年,只要cuda的发布时间晚于系统版本时间2017年,这样安装也是可以成功滴。
    在这里插入图片描述
    在cuda的官网上下载相应的安装程序。
    下载地址:CUDA下载地址
    现在美滋滋的开始安装(踩坑)CUDA啦!
    WARNING:麒麟系统内核版本、gcc版本、cuda版本相互兼容可以少踩一些坑。

    1.nouveau驱动禁用

      在中标麒麟系统中禁掉系统自带的nouveau显卡驱动,即将nouveau加入黑名单blacklist.conf中,不同的系统blacklist.conf文件在不同的位置(/etc/modprobe.d/blacklist.conf或者/usr/lib/modprobe.d/dist-blacklist.conf)。进入相应的文件,打开终端命令窗口进行修改:
    1)输入
    sudo vim blacklist.conf或者sudo vim dist-blacklist.conf命令进入文件

    2)输入 i 修改文件
    加入blacklist nouveau禁用驱动程序
    注释#blacklist nvidiafb

    3)点击esc退出修改

    4)输入 :wq保存文件

    2.img文件备份

    在boot文件中备份启动时候引导文件,打开终端窗口:
    1)备份原始镜像文件:
    sudo mv initramfs-$(uname -r).img initramfs-
    $(uname -r).img.(随意一个后缀名)

    2)重新生成不包含nouveau驱动的镜像文件:
    sudo dracut --omit-drivers initramfs-$(uname -r).img $(uname -r)
    输入reboot命令重启电脑

    3.CUDA安装

      电脑重新启动时按e键进入grub界面,在…ro rhgb quiet 后面加入 3(注意空格),设置系统以第三级别方式进行启动(多用户标准模式), 按F10保存并启动。
    在这里插入图片描述
    1)进入系统tty1界面,输入用户名和密码登录。

    2)输入lsmod | grep nouveau 命令后没有任何信息显示表示自带显卡驱动程序禁用成功

    3)进入存放英伟达显卡安装程序的文件夹,输入
    chmod 777 cudaxxx.run命令给予程序运行权限。
    输入 sudo ./cudaxxx.run开始安装驱动程序。在这里插入图片描述
    4)Ctrl+C直接跳过介绍信息,进入安装程序:
      ( a a a)警告信息、安装显卡驱动、OpenGL库、运行nvidia-xconfig、安装cuda的例程这些个选项时选择accept和yes。
      ( b b b)Do you want to install a symbolic link at /usr/local/cuda?建立链接的选项,根据个人喜欢选择是否创建cuda链接。

    5)等待安装完成后,输入reboot重启电脑后,进入cuda的例程deviceQueryDry文件中,make编译后,运行程序,如图所示出现Result = PASS,表示cuda成功安装。
    在这里插入图片描述
    6)进入cuda的文件夹/usr/local/cuda-xx(版本号)/bin中,在终端输入nvcc -V命令显示cuda信息,如下图所示,至此cuda安装成功!!!!
    在这里插入图片描述
    7)进入/ect/ld.so.conf.d文件,使用终端命令新建一个cuda.conf文件,加入cuda的引用路径,然后sudo ldconfig更新库引用。

    WARNING:
      ( a a a)双显卡黑屏的问题,将显示器的HDMI连接线从集成显卡的插槽换到独立显卡的插槽即可解决黑屏问题。
      ( b b b)我的中标麒麟系统已经安装了英伟达显卡驱动程序,在装CUDA时提示系统的磁盘余量不足,直接跳出安装程序,我尝试着把原先的显卡驱动程序卸载了,再安装CUDA,竟然意外的成功了。。。。。。这个坑的原理没搞明白,不过填上了就好。
      ( c c c)在终端使用nvcc -V命令,系统会提示bash没有这个命令,这是因为系统的可执行文件一般存在/usr/bin和/usr/local/bin文件中,把/usr/local/cuda-xx(版本号)/bin中文件拷贝到/usr/local/bin中,把/usr/local/cuda-xx(版本号)/include中文件拷贝到/usr/local/include中,这样就可以在任意一个终端窗口使用nvcc的命令了。

    二、安装CUDNN

      在cudnn官网进行注册一个账户,然后下载cudnn库,需要匹配CUDA版本,然后解压cudnn库得头文件和库文件。
    下载地址:CUDNN下载地址
    安装CUDNN就是将头文件cudnn.h文件拷贝到/usr/local/include/和/usr/local/cuda-xx(版本号)/include/中,然后将库文件中所有文件拷贝到/usr/local/cuda-xx(版本号)/lib64/中即可。
    输入命令:
     1)sudo cp -rf xx…/include/cudnn.h /usr/local/include/
     2)sudo cp -rf xx…/include/cudnn.h /usr/local/cuda-xx(版本号)/include/
     3)sudo cp -rf xx…/lib64/* /usr/local/cuda-xx(版本号)/lib64/
     4)sudo ldconfig更新库引用

    总结

    遇见问题不要慌,问题总会解决的嘛!!!!

    展开全文
  • Win10安装CUDA10和cuDNN

    千次阅读 2019-04-09 21:02:08
    原 Win10安装CUDA10和cuDNN 2018年10月12日 15:50:48 仰天长笑笑长天 阅读数:16926 ...
  • Linux小白讲述ubuntu18.04 安装cuda和pytorch的注意事项1. cuda安装2.cudnn安装3.pytorch安装4.总结 1. cuda安装 1.1 版本问题 务必要先从pytorch官网看对应的cuda和cudnn的版本,我第一次安装时听别人讲CUDA比较难...
  • Acer宏碁 Aspier E 14系列,型号 E5-471G-57...小破笔记本,就为了安个环境跑跑小的demo,双显卡安cuda真费劲,贼心不死,记录下尝试的各种过程,也许成了也能为类似情况提供一些建议。 安装准备 多次重装,Ubunt...
  • ubuntu下安装cuda+cudnn步骤

    万次阅读 2017-11-09 18:15:09
    今天自己在电脑上安装cuda ,对于小白,一天真的时很漫长,简单记录一下,日后方便自己再次安装。自己是根据官网上的安装文件进行操作的。http://developer.nvidia.com/cuda-downloads。 给出官网上的型号适用图...
  • 环境配置3-Ubuntu下安装CUDA8.0和cuDNN

    千次阅读 2018-06-09 13:26:54
    安装CUDA1-首先对系统进行更新 sudo apt-get update sudo apt-get upgrade2-安装依赖项 sudo apt-get install libprotobuf-dev libleveldb-dev libsnappy-dev libopencv-dev libhdf5-serial-dev protobuf...
  • 原标题:centos7笔记本双显卡安装nvidia驱动并成功安装cuda这是一篇介绍如何在笔记本上使用centos 7 系统安装nvidia驱动和cuda 的文章。本篇还会介绍清晰的安装思路。探索了三天才完成,最后才搞懂思路,中间其实...
  • 这是一篇介绍如何在笔记本上使用centos 7 系统安装nvidia驱动和cuda 的文章。本篇还会介绍清晰的安装思路。探索了三天才完成,最后才搞懂思路,中间其实遇到了一个正确的教程,但是没有思路,所以一直没成功,这里先...
  • NVIDIA Docker CUDA容器化原理分析

    千次阅读 2019-12-31 10:31:33
    在AI应用容器化时,会碰到cuda...为了彻底解决这类问题,了解了CUDA API的体系结构,并对NVIDIA Docker实现CUDA容器化原理进行了分析。 CUDA API体系结构 CUDA是由NVIDIA推出的通用并行计算架构,通过一些CUDA库提供...
  • CUDA-GPU\CUDA开发手册,纤细介绍了CUDA的开发原理,开发步骤,以及相关的API-CUDA
  • 从NVIDIA官方网站上找的CUDA资料可以看出CUDA的实现流程如下图:CUDA的实现流程从图上我们可以看出CUDA在整个GPU计算中充当的就是翻译的角色,我们知道GPU的结构和CPU差别很大,GPU强调的是并行性重复性的计算工作,GPU...
  • CUDA:矩阵乘法原理

    2013-05-23 22:45:50
    CUDA:矩阵乘法原理
  • ubuntu18 安装pytorch(cuda11.0)

    千次阅读 2020-11-21 13:27:31
    3.安装cuda (如果之前没装显卡驱动的话,安装cuda时可选一起装,以前装过就不勾选一起装),装完配置环境变量 export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda-11.0/lib64 export PATH=$PATH:/usr/local/...
  • CUDA矩阵计算原理和方法

    千次阅读 2019-05-15 11:16:22
    文章目录基本概念主机设备线程(Thread)线程块(Block)线程格(Grid)线程束核函数(Kernel)函数修饰符常用 GPU内存函数cudaMalloc()cudaMemcpy()cudaFree()atomicAdd(addr,y)\_\_syncthreads()GPU内存的分类全局...
  • ROI Align原理cuda源码阅读

    千次阅读 2019-06-09 10:37:36
    文章目录原理pytorch cuda源码阅读(前向) 原理 具体可参考: 详解 ROI Align 的基本原理和实现细节。这篇文章为整体的原理理解,并不涉及算法的具体实现。简单看。 双线性插值算法的详细总结。这篇文章涉及...
  • 1.机器中已经安装cuda9.0以及相应的cudnn,并且配置了多个版本的gcc,并且目前使用的是gcc6以下(cuda9.0安装需要gcc6以下的,而原本系统中默认gcc7。) 2.下载安装cuda10.0,不要配置环境变量(这样可以保证默认...
  • TensorFlow-GPU极简———不用安装cuda、cuDNN!!!

    千次阅读 热门讨论 2021-01-02 15:01:30
    看到网上大部分TensorFlowGPU教程都需要看配置需要装一大堆东西,然后就来更新一篇我知道的最简便的安装GPU版本的方法了。如下:
  • 总标题说在前边首先去[官网](https://pytorch.org/get-started/locally/)安装[anconda](https://www.anaconda.com/)以及创建环境和换国内源安装创建环境换国内源cuda下载首先你得知道你的n卡支不[支持]...
  • 本小节来自《大规模并行处理器编程实战》第四节,该书是很好的从内部原理结构上来讲述了CUDA的,对于理解CUDA很有帮助,借以博客的形式去繁取间,肯定会加入自己个人理解,所以有错误之处还望指正。该书还出版了第二...
  • 双系统安装+CUDA环境配置

    千次阅读 2018-03-19 18:38:27
    双系统安装+CUDA环境配置 这两天尝试装了Windows10+Ubuntu16.04双系统,并且在Ubuntu上配置CUDA环境,特此记录这两天我走过的坑 1、双系统安装 本人的笔记本配置表: 牌子:DELL CPU:I7-6700 ...
  • 三、cuda安装:版本号10.2 四、cudnn的安装:版本号7.6.5 五、Anaconda的安装 六、Pytorch的安装 七、用户的添加: 八、硬盘的挂载并设置开机自动挂载 九、配置ssh以便xshell远程连接 一、Ubuntu系统的安装 ...
  • 本小节来自《大规模并行处理器编程实战》第四节,该书是很好的从内部原理结构上来讲述了CUDA的,对于理解CUDA很有帮助,借以博客的形式去繁取间,肯定会加入自己个人理解,所以有错误之处还望指正。 一、块索引与...

空空如也

空空如也

1 2 3 4 5 ... 20
收藏数 8,296
精华内容 3,318
关键字:

安装cuda原理