tensorflow的CPU版本

如果要在笔记本等没有GPU的机器上使用tensorflow,需要支持新的指令集(FMA, AVX, AVX2, SSE4.1, SSE4.2),可以去这个老兄的github上下载。

https://github.com/lakshayg/tensorflow-build

2020.08.02 ~ Latest Builds

TFHWOSGCCPythonSupports
2.2.0CPUUbuntu 18.047.5.03.7.7FMA, AVX, AVX2, SSE4.1, SSE4.2Download
2.1.0CPUUbuntu 16.045.4.03.5.2FMA, AVX, AVX2, SSE4.1, SSE4.2Download
1.14.1CPUmacOS Mojaveclang-1001.0.46.43.7.4FMA, AVX, AVX2, SSE4.1, SSE4.2Download

cuda资料整理(5)

CUDA并行算法系列之FFT快速卷积

卷积定义

在维基百科上,卷积定义为:

连续卷积公式

离散卷积定义为:

离散卷积公式

[ 0, 1, 2, 3][0, 1, 2]的卷积例子如下图所示:

卷积示例

Python实现(直接卷积)

根据离散卷积的定义,用Python实现:

def conv(a, b):
    N = len(a)
    M = len(b)
    YN = N + M - 1
    y = [0.0 for i in range(YN)]
    for n in range(YN):
        for m in range(M):
            if 0 <= n - m and n - m < N:
                y[n] += a[n - m] * b[m]
    return y

把数组b逆序,则可以不交叉计算卷积(使用numpyarray[::-1]即可实现逆序):

import numpy as np
def conv2(a, b):
    N = len(a)
    M = len(b)
    YN = N + M - 1
    y = [0.0 for i in range(YN)]
    b = np.array(b)[::-1]       # 逆序
    for n in range(YN):
        for m in range(M):
            k = n - M + m + 1;
            if 0 <= k and k < N:
                y[n] += a[k] * b[m]
    return y

测试

可以利用numpy.convolve来检验计算结果的正确性:

if __name__ == '__main__':
    a = [ 0, 1, 2, 3 ]
    b = [ 0, 1, 2 ]
    print(conv2(a, b))
    print(np.convolve(a, b))

完整代码可以在Github上找到。

利用FFT快速卷积

时域的卷积和频域的乘法是等价的,同时时域的乘法和频域的卷积也是等价的。基于这个这个前提,可以把待卷积的数组进行FFT变换,在频域做乘法,然后再进行IFFT变换即可得到卷积结果。算法流程描述如下:

  1. N=len(a)M = len(b), 其中ab为待卷积的数组,将长度增加到L>=N+M−1,L=2n,n∈ZL>=N+M−1,L=2n,n∈Z,即 L=2logN+M−12+1L=2log2N+M−1+1。
  2. 增加ab的长度到L,后面补零。
  3. 分别计算afft=fft(a)afft=fft(a),bfft=fft(b)bfft=fft(b)。
  4. abfft=afft×bfftabfft=afft×bfft。
  5. 用IFFT计算abaft的FFT逆变换,取前(N + M – 1)个值即为卷积结果。

FFT快速卷积Python代码如下:

def convfft(a, b):
    N = len(a)
    M = len(b)
    YN = N + M - 1
    FFT_N = 2 ** (int(np.log2(YN)) + 1)
    afft = np.fft.fft(a, FFT_N)
    bfft = np.fft.fft(b, FFT_N)
    abfft = afft * bfft
    y = np.fft.ifft(abfft).real[:YN]
    return y

测试

对比直接卷积、FFT卷积、numpy的卷积结果:

if __name__ == '__main__':
    a = [ 0, 1, 2, 3 ]
    b = [ 0, 1, 2 ]
    print(conv2(a, b))
    print(convfft(a, b))
    print(np.convolve(a, b))

可以看到,3个版本的计算结果是一致的。完整代码可以在Github上找到。

性能分析

复杂度分析

直接卷积的时间复杂度为o(MN),即o(n2)o(n2)。
FFT的时间复杂度为o(nlogn),FFT卷积复杂度为3次FFT+L次乘法,3o(nlogn)+o(n)=o(nlogn)3o(nlogn)+o(n)=o(nlogn),及o(nlogn)o(nlogn)。
在实际应用中,卷积核(b)被提前计算,则只需2次FFT变换。

运行测试

分别测试3个版本在数组长度为n * 1000 + 10, n=0,1,…,9的运行时间,并绘制运行时间曲线,编写如下测试代码:

def time_test():
    import time
    import matplotlib.pyplot as plt

    def run(func, a, b):
        n = 1
        start = time.clock()
        for j in range(n):
            func(a, b)
        end = time.clock()
        run_time = end - start
        return run_time / n

    n_list = []
    t1_list = []
    t2_list = []
    t3_list = []
    for i in range(10):
        count = i * 1000 + 10
        print(count)
        a = np.ones(count)
        b = np.ones(count)
        t1 = run(conv, a, b)    # 直接卷积
        t2 = run(conv2, a, b)
        t3 = run(convfft, a, b) # FFT卷积
        n_list.append(count)
        t1_list.append(t1)
        t2_list.append(t2)
        t3_list.append(t3)

    # plot
    plt.plot(n_list, t1_list, label='conv')
    plt.plot(n_list, t2_list, label='conv2')
    plt.plot(n_list, t3_list, label='convfft')
    plt.legend()
    plt.title(u"convolve times")
    plt.ylabel(u"run times(ms/point)")
    plt.xlabel(u"length")
    plt.show()

运行得到的曲线图如下:

卷积运行时间

从图中可知,FFT卷积比直接卷积速度要快很多。完整代码可以在Github上找到

CUDA实现

直接卷积

只需要把外层循环并行化就可以在CUDA上实现卷积,代码如下:

// 直接计算卷积
__global__ void conv_kernel(const float *ina, const float *inb, float *out, size_t len_a, size_t len_b, size_t len_out)
{
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= len_out)
    {
        return;
    }

    float sum = 0.0f;
    for (int m = 0; m < len_b; ++m)
    {
        int k = tid - m;
        if (0 <= k && k < len_a)
        {
            sum += ina[k] * inb[m];
        }
    }
    out[tid] = sum;
}

当然,可以使用共享内存和常量内存(卷积核存入常量内存)进行优化,优化的代码请查看Github

cuFFT卷积

使用CUDA的cuFFT可以方便的进行快速傅里叶变换,cuFFT的详细说明可以查看NVIDIA的官方文档。本文主要使用到一下两个函数:

  • cufftExecR2C:实数到复数的快速傅里叶变换(FFT)
  • cufftExecC2R:复数到实数的快速傅里叶逆变换(IFFT)

基于cuFFT的实数到复数的快速傅里叶变换代码如下:

void fft(float *in, Complex *out, size_t size)
{
    cufftHandle plan;
    cufftPlan1d(&plan, size, CUFFT_R2C, 1);
    cufftExecR2C(plan, in, out);
    cufftDestroy(plan);
}

基于cuFFT的复数到实数的快速傅里叶逆变换代码如下:

void ifft(Complex *in, float *out, size_t size)
{
    cufftHandle plan;
    cufftPlan1d(&plan, size, CUFFT_C2R, 1);
    cufftExecC2R(plan, in, out);
    cufftDestroy(plan);
}

其中Complex被定义为float2typedef float2 Complex;

有了FFT,那么基于CUDA的卷积代码可如下编写:

void convfft( float *ina, float *inb, float *out, size_t len_out, size_t L, size_t numThreads)
{
    thrust::device_vector<Complex> d_a_fft(L);
    thrust::device_vector<Complex> d_b_fft(L);
    thrust::device_vector<Complex> d_c_fft(L);

    Complex *raw_point_a_fft = thrust::raw_pointer_cast(&d_a_fft[0]);
    Complex *raw_point_b_fft = thrust::raw_pointer_cast(&d_b_fft[0]);
    Complex *raw_point_c_fft = thrust::raw_pointer_cast(&d_c_fft[0]);

    fft(ina, raw_point_a_fft, L);
    fft(inb, raw_point_b_fft, L);

    // 计算 d_c_fft = d_a_fft * d_b_fft;

    ifft(raw_point_c_fft, out, L);
}

最后只剩下乘法运算了,可以自己编写一个复数乘法的内核,也可以使用thrushtransform。使用thrush实现复数乘法,首先定义一个复数乘法操作函数(可以参考Transformations):

struct complex_multiplies_functor
{
    const int N;

    complex_multiplies_functor(int _n) : N(_n) {}

    __host__ __device__ Complex operator()(const Complex &a, const Complex &b) const
    {
        Complex c;
        c.x = (a.x * b.x - a.y * b.y) / N;
        c.y = (a.x * b.y + a.y * b.x) / N;
        return c;
    }
};

然后使用thrush::transform即可完成计算:

// 计算 d_c_fft = d_a_fft * d_b_fft;
thrust::transform(d_a_fft.begin(), d_a_fft.end(), d_b_fft.begin(), d_c_fft.begin(), complex_multiplies_functor(L));

结语

本文首先简要介绍了卷积运算,然后使用Python实现了卷积运行的代码,接着讨论了基于FFT的快速卷积算法,并使用Python实现了FFT卷积,接着对直接卷积和基于FFT的快速卷积算法的性能进行了分析,从实验结果可以看出,FFT卷积相比直接卷积具有更快的运行速度。最后,基于CUDA实现了直接卷积算法,并且使用cuFFT和thrush在CUDA平台实现了基于FFT的快速卷积算法。

本文完整代码可在Github上下载

参考文献

  1. 维基百科.卷积.https://zh.wikipedia.org/zh/%E5%8D%B7%E7%A7%AF
  2. 百度文库.利用FFT计算卷积.http://wenku.baidu.com/view/5606967101f69e3143329407.html
  3. 用Python做科学计算.FFT卷积的速度比较.http://old.sebug.net/paper/books/scipydoc/example_spectrum_fft_convolve_timeit.html
  4. NVIDIA.cuFFT.https://developer.nvidia.com/cufft
  5. thrust. https://github.com/thrust/thrust/tree/master/thrust

cuda资料整理(4)

傅立叶变换—FFT(cuda实现)

背景:

无意间看到cuda解决FFT有一个cufft函数库,大体查看了有关cufft有关知识,写了一个解决一维情况的cuda代码,据调查知道cufft在解决1D,2D,3D的情况时间复杂度都为O(nlogn),附上解决一维情况的代码,准备后面找一些详细的资料去学习一下cuda的函数库。

复制代码
#include "stdio.h"
#include "cuda_runtime.h"
#include "cufft.h"
#include "device_launch_parameters.h"


#define LENGTH 4
int main()
{

float Data[LENGTH] = {1,2,3,4}; cufftComplex *CompData=(cufftComplex*)malloc(LENGTH*sizeof(cufftComplex)); int i; for(i=0;i<LENGTH;i++) { CompData[i].x=Data[i]; CompData[i].y=0; } cufftComplex *d_fftData; cudaMalloc((void**)&d_fftData,LENGTH*sizeof(cufftComplex)); cudaMemcpy(d_fftData,CompData,LENGTH*sizeof(cufftComplex),cudaMemcpyHostToDevice); cufftHandle plan; cufftPlan1d(&plan,LENGTH,CUFFT_C2C,1); cufftExecC2C(plan,(cufftComplex*)d_fftData,(cufftComplex*)d_fftData,CUFFT_FORWARD); cudaDeviceSynchronize(); cudaMemcpy(CompData,d_fftData,LENGTH*sizeof(cufftComplex),cudaMemcpyDeviceToHost); for(i=0;i<LENGTH;i++) { if(CompData[i].x != 0) { printf("%3.1f",CompData[i].x); } if(CompData[i].y != 0 ) { printf("+%3.1fi",CompData[i].y); } printf("\n"); } cufftDestroy(plan); free(CompData); cudaFree(d_fftData); }
复制代码

在Linux下运行的这段代码:

编译命令:nvcc -o fftcu FFT.cu -I /usr/local/cuda/include  -L /usr/local/cuda/lib64 -lcufft

运行命令:./fftcu

注:/usr/local/cuda/include中有cufft.h头文件,/usr/local/cuda/lib64中有libcufft.so库文件

cuda资料整理(3)

CUDA基础介绍 https://blog.csdn.net/fengbingchun/article/details/54691225

一、GPU简介

1985年8月20日ATi公司成立,同年10月ATi使用ASIC技术开发出了第一款图形芯片和图形卡,1992年4月ATi发布了Mach32图形卡集成了图形加速功能,1998年4月ATi被IDC评选为图形芯片工业的市场领导者,但那时候这种芯片还没有GPU的称号,很长的一段时间ATi都是把图形处理器称为VPU,直到AMD收购ATi之后其图形芯片才正式采用GPU的名字。

NVIDIA公司在1999年发布GeForce 256图形处理芯片时首先提出GPU的概念。GPU使显卡削减了对CPU的依赖,并实现部分原本CPU的工作,尤其是在3D图形处理时。GPU所采用的核心技术有硬体T&L(Transform and Lighting,多边形转换和光源处理)、立方环境材质贴图与顶点混合、纹理压缩及凹凸映射贴图、双重纹理四像素256位渲染引擎等,而硬体T&L技术能够说是GPU的标志。

GPU(Graphics Processing Unit)即图形处理器,又称显示核心、视觉处理器、显示芯片,是一种专门在个人电脑、工作站、游戏机和一些移动设备(如平板电脑、智能手机等)上作图像运算工作的微处理器。

显卡作为电脑主机里的一个重要组成部分,承担输出显示图形的任务。显卡的处理器称为图形处理器(GPU),它是显卡的”心脏”,与CPU类似,只不过GPU是专为执行复杂的数学和几何计算而设计的,这些计算是图形渲染所必需的。

时下的GPU多数拥有2D或3D图形加速功能。有了GPU,CPU就从图形处理的任务中解放出来,可以执行其他更多的系统任务,这样可以大大提高计算机的整体性能。

GPU会产生大量热量,所以它的上方通常安装有散热器或风扇。

GPU是显示卡的”大脑”,GPU决定了该显卡的档次和大部分性能,同时GPU也是2D显示卡和3D显示卡的区别依据。2D显示芯片在处理3D图像与特效时主要依赖CPU的处理能力,称为软加速。3D显示芯片是把三维图像和特效处理功能集中在显示芯片内,也就是所谓的”硬件加速”功能。显示芯片一般是显示卡上最大的芯片(也是引脚最多的)。时下市场上的显卡大多采用NVIDIA和 AMD-ATI 两家公司的图形处理芯片。

GPU已经不再局限于3D图形处理了,GPU通用计算技术发展已经引起业界不少的关注,在浮点运算、并行计算等部分计算方面,GPU可以提供数十倍乃至于上百倍于CPU的性能。

GPU通用计算方面的标准目前有OpenCL、CUDA、AMD APP、DirectCompute。

二、GPU通用计算编程

对GPU通用计算进行深入研究从2003年开始,并提出了GPGPU概念,前一个GP则表示通用目的(General Purpose),所以GPGPU一般也被称为通用图形处理器或通用GPU。

GPU通用计算通常采用CPU+GPU异构模式,由CPU负责执行复杂逻辑处理和事务处理等不适合数据并行的计算,由GPU负责计算密集型的大规模数据并行计算。

OpenCL(Open Computing Language,开放运算语言)是第一个面向异构系统通用目的并行编程的开放式、免费标准,也是一个统一的编程环境,便于软件开发人员为高性能计算服务器、桌面计算系统、手持设备编写高效轻便的代码,而且广泛适用于多核心处理器(CPU)、图形处理器(GPU)、Cell类型架构以及数字信号处理器(DSP)等其他并行处理器,AMD-ATI、NVIDIA时下的产品都支持OpenCL。目前,OpenCL最新版本为2.2.

CUDA(Compute Unified Device Architecture)是一种将GPU作为数据并行计算设备的软硬件体系,硬件上NVIDIA GeForce 8系列以后的GPU(包括GeForce、ION、Quadro、Tesla系列)已经采用支持CUDA的架构,软件开发包上CUDA也已经发展到CUDA Toolkit 8.0,并且支持Windows、Linux、MacOS三种主流操作系统。CUDA采用比较容易掌握的类C语言进行开发。

         AMD APP(AMD Accelerated Parallel Processing)是AMD加速并行处理技术。是AMD针对旗下图形处理器(GPU)所推出的通用并行计算技术。利用这种技术可以充分发挥AMD GPU的并行运算能力,用于对软件进行加速运算或进行大型的科学运算。AMD APP技术的前身称作ATI Stream。2010年10月,随着AMD Radeon HD6800系列显卡的发布,ATI品牌正式被AMD取代。ATI Stream技术也随着技术升级并更名为AMD APP技术。目前,AMD APP SDK最新版本为3.0.

         DirectCompute是一种用于GPU通用计算的应用程序接口,由Microsoft(微软)开发和推广,集成在Microsoft DirectX内。目前,最新的DirectX版本为DirectX 12,安装在windows 10上。DirectX 11内集成DirectCompute 5.0,那DirectX 12内应该是集成DirectCompute 6.0吧。

         其中OpenCL、DirectCompute、AMD APP(基于开放型标准OpenCL开发)是开放标准,CUDA是私有标准。

三、NVIDIA 显卡系列

NVIDIA(英伟达)创立于1993年1月,是一家以设计智核芯片组为主的无晶圆(Fabless)IC半导体公司。

NVIDIA已经开发出了五大产品系列,以满足特定细分市场需求,包括:GeForce、Tegra、ION、Quadro、Tesla。

Geforce系列主要面向家庭和企业的娱乐应用,该系列又可以分为面向性能的GTX系列,面向主流市场的GTS和GT系列,已经具有高性价比的GS系列。

Quadro系列主要应用于图形工作站中,对专业领域应用进行了专门优化。

Tesla系列是专门用于高性能通用计算的产品线。

Tegra系列是NVIDIA为便携式和移动领域推出的全新解决方案,在极为有限的面积上集成了通用处理器、GPU、视频解码、网络、音频输入输出等功能,并维持了极低的功耗。

针对Geforce显卡系列,NVIDIA各代显卡都遵循了由高至低命名规则:GTX>GTS>GT>GS

从GTX 500系开始,为避免命名复杂带来的产品线识别困扰,NVIDIA显卡将取消GTS级别的显卡,中高端全部使用GTX命名,而低端使用GT命名,带Ti后缀为更高一级显卡,如GTX 560 Ti > GTX 560.

         NVIDIA显卡末尾数字解读,以GeForce GTX 980M:GTX代表是高端显卡的意思;980M:第一位数字9,代表第几代的意思(9是高端显卡第九代的意思,如果末尾数字有四位,则前两位表示是第多少代的意思,如GeForce GTX 1080)。第二位至关重要,因为显卡分高端显卡,中端显卡,入门级显卡就是取决于第二位数字的。第二位数字是1-2代表是入门级显卡;第二位数字是3-5代表是中端显卡;第二位数字是6-9代表是高端显卡。第三位数字是一个特殊的标志,几乎能在市场上买到的显卡都是0结尾的,如果第三位数字为5的显卡一般都是OEM显卡,即只给大厂子做品牌机的特供。数字越大,性能越好。显卡数字后缀Ti,代表加强。

如果用显卡来进行各种运算,衡量显卡性能的参数可包括:(1)、核心数目;(2)、显存带宽(GPU计算能力太强,很多时候瓶颈都在数据传输上);(3)、峰值单精度浮点计算能力;(4)、峰值双精度浮点计算能力;(5)、时钟频率;(6)、架构版本。

四、CUDA基础

1.      简介

CUDA(Compute Unified Device Architecture,统一计算设备架构),是显卡厂商NVIDIA在2007年推出的并行计算平台和编程模型。它利用图形处理器(GPU)能力,实现计算性能的显著提高。CUDA是一种由NVIDIA推出的通用并行计算架构,该架构使GPU能够解决复杂的计算问题,从而能通过程序控制底层的硬件进行计算。它包含了CUDA指令集架构(ISA)以及GPU内部的并行计算引擎。开发人员可以使用C/C++/C++11语言来为CUDA架构编写程序。CUDA提供host-device的编程模式以及非常多的接口函数和科学计算库,通过同时执行大量的线程而达到并行的目的。

3.0以下版本仅支持C编程,从3.0版本开始支持C++编程,从7.0版本开始支持C++11编程。

CUDA仅能在有NVIDIA显卡的设备上才能执行,并不是所有的NVIDIA显卡都支持CUDA,目前NVIDIA的GeForce、ION、Quadro以及Tesla显卡系列上均可支持。根据显卡本身的性能不同,支持CUDA的版本也不同。

2.      安装

(1)、在windows上的安装可以参考:http://blog.csdn.net/fengbingchun/article/details/53892997

(2)、在ubuntu上的安装可以参考:http://blog.csdn.net/fengbingchun/article/details/53840684

3.      使用CUDA C编写代码的前提条件

(1)、支持CUDA的图形处理器:从2007年开始,NVIDIA新推出的并且显存超过256MB的GPU都可以用于开发和运行基于CUDAC编写的代码。

(2)、NVIDIA设备驱动程序:NVIDIA提供了一些系统软件来实现应用程序与支持CUDA的硬件之间的通信,即显卡驱动程序。要确保安装匹配的驱动程序,选择与开发环境相符的图形卡和操作系统。

(3)、CUDA开发工具箱:CUDA Toolkit,注意选择与操作系统相匹配的CUDA Toolkit。

(4)、标准C编译器:由于CUDA C应用程序将在两个不同的处理器上执行计算,因此需要两个编译器。其中一个编译器为GPU编译代码,而另一个为CPU编译代码。下载并安装CUDA Toolkit后,就会获得一个编译GPU代码的编译器。对于CPU编译器,Windows推荐使用Visual Studio,Linux使用GNU C编译器(gcc),Mac使用Xcode。

4.      设备计算能力

设备计算能力的版本描述了一种GPU对CUDA功能的支持程度。计算能力版本中小数点前的第一位用于表示设备核心架构,小数点后的第一位则表示更加细微的进步,包括对核心架构的改进以及功能的完善等。例如,计算能力1.0的设备能够支持CUDA,而计算能力1.1设备加入了对全局存储器原子操作的支持,计算能力1.2的设备则可以支持warp vote函数等更多功能,而计算能力1.3的设备又加入了对双精度浮点运算功能。

GeForce GTX 970型号计算能力为5.2,GeForce GT 640M型号计算能力为3.0,目前GeForce系列最高的计算能为6.1,可在https://developer.nvidia.com/cuda-gpus中查找各种系列型号的计算能力以及查找指定的显卡型号是否支持CUDA。

5.      软件体系

CUDA的软件堆栈由三层构成,如下图,CUDA Library、CUDA runtimeAPI、CUDA driver API. CUDA的核心是CUDA C语言,它包含对C语言的最小扩展集和一个运行时库,使用这些扩展和运行时库的源文件必须通过nvcc编译器进行编译。

CUDA C语言编译得到的只是GPU端代码,而要管理GPU资源,在GPU上分配显存并启动内核函数,就必须借助CUDA运行时API(runtime API)或者CUDA驱动API(driver API)来实现。在一个程序中只能使用CUDA运行时API与CUDA驱动API中的一种,不能混合使用

6.      CUDA C语言

CUDA C语言为程序员提供了一种用C语言编写设备端代码的编程方式,包括对C的一些必要扩展和一个运行时库,CUDA对C的扩展主要包括以下几个方面:

(1)、引入了函数类型限定符,用来规定函数是在host还是在device上执行,以及这个函数是从host调用还是从device调用。这些限定符有:__device__、__host__、__global__。

(2)、引入了变量类型限定符,用来规定变量被存储在哪一类存储器上。传统的在CPU上运行的程序,编译器能自动决定将变量存储在CPU的寄存器还是内存中。在CUDA编程模型中,一共抽象出来8种不同的存储器。为了区分各种存储器,引入了一些限定符,包括:__device__、__shared__、__constant__。

(3)、引入了内置矢量类型,如char4、ushort3、double2、dim3等,它们是由基本的整形或浮点型构成的矢量类型,通过x、y、z、w访问每一个分量,在设备端代码中各矢量类型有不同的对齐要求。

(4)、引入了4个内置变量:blockIdx和threadIdx用于索引线程块和线程,gridDim和blockDim用于描述线程网格和线程块的维度。warpSize用于查询warp中的线程数量。

(5)、引入了<<<>>>运算符,用于指定线程网格和线程块维度,传递执行参数。

对__global__函数的任何调用都必须指定该调用的执行配置(execution configuration)。执行配置用于定义在设备上执行函数时的grid和block的维度,以及相关的流。

使用驱动API时,需要通过一系列驱动函数设置执行配置参数。

使用运行时API时,需要在调用的内核函数名与参数列表直接以<<<Dg,Db,Ns,S>>>的形式设置执行配置,其中:

Dg是一个dim3型变量,用于设置grid的维度和各个维度上的尺寸。设置好Dg后,grid中将有Dg.x*Dg.y个block,Dg.z必须为1.

Db是一个dim3型变量,用于设置block的维度和各个维度上的尺寸。设置好Db后,每个block中将有Db.x*Db.y*Db.z个thread。

Ns是一个size_t型变量,指定各块为此调用动态分配的共享存储器大小,这些动态分配的存储器可供声明为外部数组(extern __shared__)的其他任何变量使用;Ns是一个可选参数,默认值为0.

S为cudaStream_t类型,用于设置与内核函数关联的流。S是一个可选参数,默认值为0.

(6)、引入了一些函数:memory fence函数、同步函数、数学函数、纹理函数、测时函数、原子函数、warp vote函数。

以上扩展均有一些限制,如果违背了这些限制,nvcc将给出错误或警告信息,但有时也不会报错,程序无法运行。

7.      常用术语

(1)、主机(host):将CPU及系统的内存称为主机。

(2)、设备(device):将GPU及GPU本身的显示内存称为设备,在一个系统中可以存在一个主机和若干个设备。

CUDA编程模型中,CPU与GPU协同工作,CPU负责进行逻辑性强的事务处理和串行计算,GPU则专注于执行高度线程化的并行处理任务。CPU、GPU各自拥有相互独立的存储器地址空间:主机端的内存和设备端的显存。

(3)、线程(Thread):一般通过GPU的一个核进行处理,可以表示成一维、二维、三维。一个block中的所有thread在一个时刻执行指令并不一定相同。

(4)、线程块(Block):由多个线程组成,可以表示成一维、二维、三维;各block是并行执行的,block间无法通信,也没有执行顺序;注意线程块的数量有限制(硬件限制)。

Block内,可以通过__syncthreads()进行线程同步;thread间通过shared memory进行通信。

在实际运行中,block会被分割成更小的线程束(warp)。线程束的大小由硬件的计算能力版本决定。Warp中的线程只与thread ID有关,而与block的维度和每一维的尺度没有关系。

(5)、线程格(Grid):由多个线程块组成,可以表示成一维、二维、三维。

(6)、线程束:在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被”编织在一起”并且”步调一致”的形式执行,在程序中的每一行,线程束中的每个线程都将在不同数据上执行相同的命令

(7)、核函数(Kernel):运行在GPU上的CUDA并行计算函数称为kernel(内核函数)。内核函数必须通过__global__函数类型限定符定义,并且只能在主机端代码中调用。在调用时,必须声明内核函数的执行参数即”<<< >>>”,用于说明内涵函数中的线程数量,以及线程是如何组织的。不同计算能力的设备对线程的总数和组织方式有不同的约束。必须先为Kernel中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误,例如越界或报错,甚至导致蓝屏和死机。

在设备端运行的线程之间是并行执行的,其中的每个线程则按照指令的顺序串行执行一次kernel函数每一个线程有自己的block ID和thread ID用于与其它线程相区分。blockID和thread ID只能在kernel中通过内置变量访问。内置变量不需要由程序员自己定义,是由设备中的专用寄存器提供的。因此,内置变量是只读的,并且只能在GPU端的kernel函数中使用。

Kernel是以block为单位执行的,CUDA引入grid只是用来表示一系列可以被并行执行的block的集合。各block是并行执行的,block间无法通信,也没有执行顺序,在同一个block中的线程,可以进行数据通信,在同一个block中的线程通过共享存储器(shared memory)交换数据,并通过栅栏同步(可以在kernel函数中需要同步的位置调用__syncthreads()函数)保证线程间能够正确地共享数据。这样,无论是只能同时处理一个线程块的GPU上,还是在能同时处理数十乃至上百个线程块的GPU上,这一CUDA编程模型都能很好地适用。

一个kernel函数并不是一个完整的程序,而是整个CUDA程序中一个可以被并行执行的步骤。一个完整的CUDA程序是由一系列的设备端kernel函数并行步骤和主机端的串行处理步骤共同组成的。如下图(CUDA编程模型):

CPU串行代码完成的工作包括在kernel启动前进行数据准备和设备初始化的工作,以及在kernel之间进行一些串行计算。理想情况下,CPU串行代码的作用应该只是清理上一个内核函数,并启动下一个内核函数。在这种情况下,可以在设备上完成尽可能多的工作,减少主机与设置之间的数据传输。

8.      内置变量

内置变量用于确定grid和block的维度,以及block和thread在其中的索引。这些内置变量只能在设备端执行的函数(__global__、__device__)中使用。

(1)、dim3:基于uint3定义的矢量类型,相当于由3个unsigned int类型组成的结构体,可表示一个三维数组,在定义dim3类型变量时,凡是没有赋值的元素都会被赋予默认值1.其它常用基本数据类型可参考include/vector_types.h文件。

(2)、threadIdx:内置变量,用于说明当前thread在block中的位置;如果线程是一维的可获取threadIdx.x,如果是二维的还可获取threadIdx.y,如果是三维的还可获取threadIdx.z;为uint3类型,包含了一个thread在block中各个维度的索引信息。可参考include/device_launch_parameters.h文件。

threadIdx.x取值范围是[0,blockDim.x -1],threadIdx.y取值范围[0, blockDim.y-1],threadIdx.z取值范围[0, blockDim.z-1]。

(3)、blockIdx:内置变量,用于说明当前thread所在的block在整个grid中的位置,blockIdx.x取值范围是[0,gridDim.x-1],blockIdx.y取值范围是[0, gridDim.y-1]。为uint3类型,包含了一个block在grid中各个维度上的索引信息。

对于一维的block,线程的threadID就是threadIdx.x;

对于大小为(Dx, Dy)的二维block,线程的threadID是(threadIdx.x+ threadIdx.y * Dx);

对于大小为(Dx, Dy, Dz)的三维block,线程的threadID是(threadIdx.x+ threadIdx.y * Dx + threadIdx.z * Dx * Dy).

(4)、blockDim:内置变量,用于说明每个block的维度与尺寸。为dim3类型,包含了block在三个维度上的尺寸信息。

(5)、gridDim:内置变量,用于说明整个网格的维度与尺寸,一个grid最多只有二维。为dim3类型,包含了grid在三个维度上的尺寸信息。

        uint3 __device_builtin__ __STORAGE__threadIdx;	uint3 __device_builtin__ __STORAGE__ blockIdx;	dim3 __device_builtin__ __STORAGE__ blockDim;	dim3 __device_builtin__ __STORAGE__ gridDim;

         (6)、warpSize:内置变量,用于引用warpsize。为int类型,用于确定设备中一个warp包含多少个thread.

         以上这些内置变量只能在设备端代码中使用,这些变量是只读的,不能对它们赋值,也不能对它们取地址。

9.      变量类型限定符

变量类型限定符用于指明变量存储在设备端的哪一类存储器上。

(1)、__device__:声明的变量存在于设备上。当__device__变量限定符不与其他限定符连用时,这个变量将:位于全局存储器空间中;与应用程序具有相同的生命周期;可以通过运行时库从主机端访问,设备端的所有线程也可访问。

(2)、__constant__:使用__constant__限定符,或者与__device__限定符连用,这样声明的变量:存在于常数存储器空间;与应用程序具有相同的生命周期;可以通过运行时库从主机端访问,设备端的所有线程也可访问。

(3)、__shared__:使用__shared__限定符,或者与__device__限定符连用,此时声明的变量:位于block中的共享存储器空间中;与block具有相同的生命周期;仅可通过block内的所有线程访问。

(4)、volatile:存在于全局或者共享存储器中的变量通过volatile关键字声明为敏感变量,编译器认为其他线程可能随时会修改变量的值,因此每次对该变量的引用都会被编译成一次真实的内存读指令。

以上限定符不能用于struct与union成员、在主机端执行的函数的形参以及局部变量。

__shared__和__constant__变量默认为是静态存储。

__device__、__shared__和__constant__不能用extern关键字声明为外部变量。在__shared__前可以加extern关键字,但表示的是变量大小由执行参数确定

__device__和__constant__变量只能在文件作用域中声明,不能再函数体内声明。

__constant__变量不能从device中赋值,只能从host中通过host运行时函数赋值。

__shared__变量在声明时不能初始化。

在设备代码中(__global__或者__device__函数中),如果一个变量前没有任何限定符,这个变量将被分配到寄存器中。但如果寄存器资源不足,编译器会把这些变量存放在local memory中。Local memory中的数据被存放于显存中,而且没有任何缓存可以加速local memory的读写,因此会大大降低程序的速度。

只要编译器能够解析出设备端代码中的指针指向的地址,指向shared memory或者global memory,这样的指针即受支持。如果编译器不能正确地解析指针指向的地址,那么只能使用指向global memory的指针。

在host端代码中使用指向global或者shared memory的指针,或者在device端代码中使用指向host memory的指针都将引起不确定的行为,通常会报分区错误(segmentation fault)并导致程序终止运行。

在device端通过取址符号&获得的__device__、__constant__、__shared__的地址,这样得到的地址只能在device端使用。通过在host端调用cudaGetSymbolAddress()函数可以获得__device__、__constant__的地址,这样得到的地址只能在host端使用。

10.  函数类型限定符

(1)、__global__:表明被修饰的函数在设备上执行,可以从主机端调用;

(2)、__device__:表明被修饰的函数在设备上执行,只能从设备上调用,但只能在其它__device__函数或者__global__函数中调用;

(3)、__host__:在主机端上执行,只能从主机端调用。

没有__host__、__device__、__global__限定符修饰的函数,等同于只用__host__限定符修饰的函数,函数都将仅为主机端进行编译,即编译出只能在主机端运行的版本。__host__可以与__device__一起使用,此时函数将为主机和设备进行编译,即分别编译出在主机和设备端运行的版本。

使用限制:

(1)、__device__和__global__函数不支持递归;

(2)、__device__和__global__函数的函数体内不能声明静态变量;

(3)、__device__和__global__函数的参数数目是不可变化的;

(4)、不能对__device__取指针,但可以对__global__函数取指针;

(5)、__global__与__host__不能连用;

(6)、__global__函数的返回类型必须为void;

(7)、调用__global__函数必须指明其执行配置;

(8)、对__global__函数的调用是异步的,控制权在设备执行完成之前就会返回;

(9)、__global__函数的参数目前通过共享存储器传递,总的大小不能超过256Byte。

11.  CUDA存储器模型

每一个线程拥有自己的私有存储器寄存器和局部存储器;每一个线程块拥有一块共享存储器(shared memory);最后,grid中所有的线程都可以访问同一个全局存储器(global memory)。除此以外,还有两种可以被所有线程访问的只读存储器:常数存储器(constant memory)和纹理存储器(texture memory),它们分别为不同的应用进行了优化。全局存储器、常数存储器和纹理存储器中的值在一个内核函数执行完成后将被继续保持,可以被同一程序中的其他内核函数调用。

八种存储器比较如下图:

(1)、寄存器(register):是GPU片上高速缓存器,执行单元可以以极低的延迟访问寄存器。寄存器的基本单元是寄存器文件(register file),每个寄存器文件大小为32 bit。

(2)、局部存储器(local memory):对于每个线程,局部存储器也是私有的。如果寄存器被消耗完,数据将被存储在局部存储器中。如果每个线程使用了过多的寄存器,或声明了大型结构体或数组,或者编译器无法确定数组的大小,线程的私有数据就有可能会被分配到local memory中。一个线程的输入和中间变量将被保存在寄存器或者局部存储器中。局部存储器中的数据被保存在显存中,而不是片上的寄存器或者缓存中,因此对local memory的访问速度很慢

(3)、共享存储器(shared memory):也是GPU片内的高速存储器。它是一块可以被同一block种的所有线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器一样快,是实现线程间通信的延迟最小的方法。共享存储器可用于实现多种功能,如用于保存共用的计数器或者block的公用结果。

可以将CUDA  C的关键字__shared__添加到变量声明中,这将使这个变量驻留在共享内存中。CUDA C编译器对共享内存中的变量与普通变量将分别采取不同的处理方式。对于GPU上启动的每个线程块,CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存,但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

(4)、全局存储器(global memory):全局存储器位于显存(占据了显存的绝大部分),CPU、GPU都可以进行读写访问。整个网格中的任意线程都能读写全局存储器的任意位置由于全局存储器是可写的。在目前的架构中,全局存储器没有缓存。

全局存储器能够提供很高带宽,但同时也具有较高的访存延迟。要有效地利用全局存储器带宽,必须遵守和并访问要求,并避免分区冲突。

在运行时API中,显存中的全局存储器也称为线性内存。线性内存通常使用cudaMalloc()函数分配,cudaFree()函数释放,并由cudaMemcpy()进行主机端与设备端的数据传输。通过CUDA API分配的空间未经过初始化,初始化共享存储器需要调用cudaMemset函数。

此外,也可以使用__device__关键字定义的变量分配全局存储器。这个变量应该在所有函数外定义,必须对使用这个变量的host端和device端函数都可见才能成功编译。在定义__device__变量的同时可以对其赋值。

在驱动API中,线性内存由cuMemAlloc()或cuMemAllocPitch()来分配,cuMemFree()来释放。

(5)、主机端内存(host memory):在CUDA中,主机端内存分为两种。可分页内存(pageable memory)和页锁定(page-locked或pinned)内存。可分页内存即为通过操作系统API(malloc(), new())分配的存储器空间;而页锁定内存始终不会被分配到低速的虚拟内存中,能够保证存在于物理内存中,并且能够通过DMA加速与设备端的通信。一般的主机端内存操作方法与其他程序没有任何区别。

(6)、主机端页锁定内存(pinned memory):它有一个重要的属性,即操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存上。因此,操作系统能够安全地使某个应用程序访问该内存的物理地址,因为这块内存将不会被破坏或者重新定位。它可以提高访问速度,由于GPU知道主机内存的物理地址,因此可以通过”直接内存访问DMA(Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。

pinned memory是一把双刃剑。当使用pinned memory时,你将失去虚拟内存的所有功能。特别是,在应用程序中使用每个页锁定内存时都需要分配物理内存,因为这些内存不能交换到磁盘上。这意味着,与使用标准的malloc函数调用相比,系统将更快地耗尽内存。因此,应用程序在物理内存较少的机器上会运行失败,而且意味着应用程序将影响在系统上运行的其它应用程序的性能。建议,仅对cudaMemcpy()调用中的源内存或者目标内存,才使用页锁定内存,并且在不再需要使用它们时立即释放,而不是等到应用程序关闭时才释放。

在运行时API中,通过cudaHostAlloc()和cudaFreeHost()来分配和释放pinned memory。使用pinned memory有很多好处,比如:可以达到更高的主机端—-设备端的数据传输带宽,如果页锁定内存以write-combined方式分配,带宽还能更高一些;某些设备支持DMA功能,在执行内核函数的同时利用pinned memory进行主机端与设置端之间的通信;在某些设备上,pinned memory还可以通过zero-copy功能映射到设备地址空间,从GPU直接访问,这样就不用在主存与显存间进行数据拷贝工作了。

虽然pinned memory能带来诸多好处,但它是系统中的一种稀缺资源。如果分配过多,会导致操作系统用于分页的物理内存变小,导致系统整体性能下降。

在驱动API中,pinned memory通过cuMemHostAlloc()和一些标志分配,通过cuMemFreeHost()释放。

(7)、常数存储器(constant memory):是只读的地址空间。常数存储器中的数据位于显存,但拥有缓存加速。常数存储器的空间较小(只有64KB),在CUDA程序中用于存储需要频繁访问的只读参数。当来自同一half-warp的线程访问常数存储器中的同一数据时,如果发生缓存命中,那么只需要一个周期就可以获得数据。

常数存储器有缓存机制,用以节约带宽,加快访问速度。每个SM拥有8KB的常数存储器缓存。常数存储器是只读的,因此不存在缓存一致性问题。

constant memory用于保存在核函数执行期间不会发生变化的数据。NVIDIA硬件提供了64KB的常量内存,并且对常量内存采取了不同于标准全局内存的处理方式。在某些情况下,用常量内存来替换全局内存能有效地减少内存带宽。要使用常量内存,需在变量前面加上__constant__关键字。常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。

(8)、纹理存储器(texture memory):是一种只读存储器,由GPU用于纹理渲染的图形专用单元发展而来,具备一些特殊功能。它并不是一块专门的存储器,而是牵涉到显存、两级纹理缓存、纹理拾取单元的纹理流水线。纹理存储器中的数据以一维、二维或者三维数组的形式存储在显存中,可以通过缓存加速访问,并且可以声明大小比常数存储器要大的多。在通用计算中,纹理存储器非常适合实现图像处理和查找表,对大量数据的随机访问或非对齐访问也有良好的加速效果。

在kernel中访问纹理存储器的操作称为纹理拾取(texture fetching).纹理拾取使用的坐标与数据在显存中的位置可以不同。

与常数存储器类似,纹理存储器也有缓存机制,纹理缓存有两个作用。首先,纹理缓存中的数据可以被重复利用,当一次访问需要的数据已经存在于纹理缓存中时,就可以避免对显存的再次读取。数据重用过滤了一部分对显存的访问,节约了带宽,也不必按照显存对齐的要求读取。其次,纹理缓存一次预取拾取坐标对应位置附近的几个像元,可以实现滤波模式,也可以提高具有一定局部性的访存的效率。

纹理存储器是只读的,因此没有数据一致性可言。

与constant memory类似的是,texture memory同样缓存在芯片上,因此在某些情况中,它能够减少对内存的请求并提供更高效的内存带宽。纹理缓存是专门为那些在内存访问模式中存在大量空间局部性(Spatial Locality)的图形应用程序而设计的。纹理变量(引用)必须声明为文件作用域内的全局变量。分为一维纹理内存和二维纹理内存。

12.  CUDA通信机制

(1)、同步函数:__syncthread()实现了线程块内的线程同步,它保证线程块中的所有线程都执行到同一位置。当任意一个thread运行到BAR标记处后,就会暂停运行;直到整个block中所有的thread都运行到BAR标记处以后,才继续执行下面的语句。这样,才能保证之前语句的执行结果对块内所有线程可见。如果不做同步,一个线程块中的一些线程访问全局或者共享存储器的同一地址时,可能会发生读后写、写后读、写后写错误。而通过同步可以避免这些错误的发生。

只有当整个线程块都走向相同分支时,才能在条件语句里面使用__syncthreads(),否则可能引起错误。另外,一个warp内的线程不用同步。也就是说,如果需要同步的线程处于同一warp中,则不需要调用__syncthreads()。可以使用特别的宏函数对warp内的threads进行同步。

Memory fence函数也是用来保证线程间数据通信的可靠性的。但与同步函数不同,memory fence函数并不要求所有线程都运行到同一位置,而只保证执行memory fence函数的线程生产的数据能够安全地被其它线程消费。

kernel间通信:kernel直接的数据传递,可以通过global memory实现。

GPU与CPU线程同步:在CUDA主机端代码中使用cudaThreadSynchronize(),可以实现GPU与CPU线程的同步。Kernel启动后控制权将异步返回,利用该函数可以确定所有设备端线程均已运行结束,基本只是用来实现更加准确的计时或捕获运行错误。

(2)、原子(ATOM)操作:如果操作的执行过程不能分解为更小的部分,将满足这种条件限制的操作称为原子操作。

如函数调用,atomicAdd(addr,y)将生成一个原子的操作序列,这个操作序列包括读取地址addr处的值,将y增加到这个值,以及将结果保存回地址addr。

只有1.1或者更高版本的GPU计算功能集才能支持全局内存上的原子操作,且只能在设备端使用。此外,只有1.2或者更高版本的GPU计算功能集才能支持共享内存上的原子操作。CUDA C支持多种原子操作。可参考include/device_atomic_functions.h文件。

原子函数(atomic function)对位于全局或共享存储器的一个32位或64位字执行read-modify-write的原子操作。也就是说,当多个线程同时访问全局或共享存储器的同一位置时,保证每个线程能够实现对共享可写数据的互斥操作:在一个操作完成之前,其它任何线程都无法访问此地址。例如,atomicAdd()函数可以读入共享存储器或者全局存储器中的32bit字,与一个整数求和后,将结果写回到原位置上。之所以将这一过程称为原子操作,是因为每个线程的操作都不会影响到其它线程。换句话说,原子操作能够保证对一个地址的当前操作完成之前,其它线程都不能访问这个地址

只能对有符号或者无符号整形进行原子操作(atomicExch()函数除外,该函数的操作数可以是有符号单精度浮点型)。

各种硬件对ATOM指令的支持、以及ATOM指令支持的数据类型不尽相同。

(3)、VOTE操作:VOTE指令是CUDA 2.0的新特性,只有1.2以上版本的硬件才能支持。VOTE的作用范围不是整个block,而是一个warp。

13.  异步并行执行

为了让主机端与设备端并行执行,很多函数都是异步的:控制在设备还没完成请求任务前就被返回给主机线程,这些函数有:kernel启动、以Async为后缀的内存拷贝函数、device到device内存拷贝函数、存储器初始化函数(比如cudaMemset())。

一些CUDA设备能够在kernel执行期间,执行pinnedmemory和显存间的数据传输。

异步执行的意义在于:首先,处于同一个流内的计算与数据拷贝是依次进行的,但一个流内的计算可以和另一个流的数据传输同时进行,因此通过异步执行就能够使GPU中的执行单元与存储器控制单元同时工作,提高了资源利用率;其次,当GPU在进行计算或者数据传输时就返回给主机线程,主机线程不必等待GPU运行完毕就可以继续进行一些计算,从而使得CPU和GPU可以并行工作。

如果调用了同步版本的GPU函数,在设备完成请求任务前,都不会返回主机线程,此时主机端线程将进入让步(yield)、阻滞(block)或者自旋(spin)状态。通过设置一些特定标记并调用cudaSetDeviceFlags()或cuCtxCreate()来选择主机端在进行GPU计算时进入的状态,不过和其它设置操作一样,该操作要在主机线程执行任何CUDA操作前就进行。

14.  流

程序通过流来管理并发,每个流是按顺序执行的一系列操作,而不同的流与其它的流之间乱序则是乱序执行的,也可能是并行执行的。这样,可以使一个流的计算与另一个流的数据传输同时进行,从而提高了GPU中资源的利用率。

流的定义方法,是创建一个cudaStream_t对象,并在启动内核和进行memcpy时将该对象作为参数传入,参数相同的属于同一个流,参数不同的属于不同的流。

执行参数中没有流参数,或使用0作为流参数时,不会创建流。此时,进行任何内核启动、内存设置或内存拷贝函数时,只有在之前所有的操作(包括流的部分操作)均已完成后才会开始,是异步执行方式。

驱动API提供了类似于运行时API的函数来管理流。

15.  事件

运行时API可以通过事件管理密切监控设备进度并执行准确计时,它可以异步地记录下程序内任意点的事件,并且可以查询这些事件被记录的时间。事件使用的GPU的计时器,用于测时比使用CPU的计时器更加准确。当先于该事件的所有任务(包括特定流中的所有操作)均已完成,这个事件的时戳就会被记录下来。0号流中的事件会在设备完成对所有流的操作后记录下来。事件管理可以用于测量程序运行时间,或者管理CPU和GPU同时进行计算

驱动API提供类似于运行时API的函数来管理事件。

16.  指令与指令吞吐量

在CUDA中,吞吐量指每个多处理器在一个时钟周期下执行的操作数目。对于大小为32的warp,一条指令由32个操作构成。因此,如果记T为每个时钟下的操作数目,那么指令吞吐量就是每32/T个时钟周期一条指令。

所有的吞吐量都是针对一个多处理器而言的。所以,要计算整个设备的吞吐量需要乘以设备的多处理器个数。

17.  CUDA与图形学API互操作

(1)、通过CUDA与OpenGL的互操作可以将OpenGL缓冲对象(buffer object)映射到CUDA的地址空间,这样就可以在CUDA 中读取OpenGL写入的数据,也可以用CUDA写入数据供OpenGL使用。要实现与OpenGL的互操作,必须在调用CUDA函数之前先调用cudaGLSetGLDevice()配置设备,并且在进行映射前要将OpenGL缓冲对象注册到CUDA。

要在驱动API中实现与OpenGL的互操作,就必须使用cuGLCtxCreate()而不是cuCtxCreate()创建CUDA上下文。和在运行API中一样,在进行映射前必须将缓冲对象注册到CUDA。

(2)、通过CUDA与Direct3D的互操作可以将Direct3D资源映射到CUDA地址空间,这样就可以在CUDA中读取由Direct3D写入的数据,也可以写入数据供Direct3D使用。Direct3D 9.0/10.0才支持Direct3D互操作。只有满足一些限制的Direct3D资源才能被映射到CUDA。由于DirectX 9和DirectX 10的资源有一定的差异,因此在CUDA中分别使用了不同的API与两个版本的DirectX进行互操作。

CUDA上下文一次仅可与一个Direct3D设备互操作,并且此时CUDA上下文和Direct3D设备必须是在同一个GPU上创建的。

驱动API提供了类似于运行时API的函数管理与Direct3D的互操作。

18.  Runtime API和Driver API

Runtime API比Driver API更高级,封装的更好,在Runtime之上就是封装的更好的cuFFT等库。这两个库的函数都是能直接调用的,但Driver API相对于Runtime API对底层硬件驱动的控制会更直接更方便。Driver API向后兼容支持老版本的。大部分的功能两组API都有对应的实现,一般基于Driver API的开头会是cu,而基于RuntimeAPI的开头是cuda,但基于Driver API来写程序会比RuntimeAPI要复杂。

CUDA runtime API和CUDA driverAPI提供了实现设备管理(Device management)、上下文管理(Context management)、存储器管理(Memory management)、代码块管理(Code Module management)、执行控制(Execution Control)、纹理索引管理(Texture Reference management)、与OpenGL和Direct3D的互操作性(Interoperity with OpenGL and Direct3D)的应用程序接口。

(1)、CUDA runtimeAPI在CUDA driver API的基础上进行了封装,隐藏了一些实现细节,编程更加方便,代码更加简洁。CUDA runtime API被打包存放在CUDArt包里,其中的函数都有CUDA前缀。CUDA运行时没有专门的初始化函数,它将在第一次调用运行时函数时自动完成初始化。

(2)、CUDA driverAPI是一种基于句柄的底层接口(大多对象通过句柄被引用),可以加载二进制或汇编形式的内核函数模块,指定参数,并启动计算。CUDA driver API编程复杂,但有时能通过直接操作硬件的执行实现一些更加复杂的功能,或者获得更高的性能。由于它使用的设备端代码是二进制或者汇编代码,因此可以在各种语言中调用。CUDA driver API被存放在nvCUDA包里,所有函数前缀为cu。

在调用任何一个驱动API函数之前,必须先调用cuInit()完成初始化,创建一个CUDA上下文。

19.  多设备与设备集群

在一台计算机中可以存在多个CUDA设备,通过CUDA API提供的上下文管理和设备管理功能可以使这些设备并行工作。采取这种方式建立的多设备系统可以提高单台机器的性能,节约空间和成本。

CUDA的设备管理功能是由不同的线程管理各个GPU,每个GPU在一个时刻只能被一个线程使用。除了采用C提供的多线程库外,CUDA还支持使用OpenMP管理多个设备。

除了在单个系统中使用多个GPU外,也可以使用CPU+GPU异构系统作为节点构造集群,或者设计更大规模的CPU+GPU异构超级计算机。CUDA可以与MPI一起使用,提供成本更低,体积和功耗更小,性能更强的高性能计算解决方案。

(1)、CUDA设备控制:一个系统中可以有一个主机或多个设备。可以通过CUDA枚举这些设备,并查询它们的属性,每个主机端线程可以选取其中的一个设备执行内核程序。每个主机端线程各自管理一个设备,当主机端存在多个下线程时,就可以使多个设备能够并行工作。一个主机端线程通过CUDA运行时分配的CUDA资源不能被其它的主机端线程使用。

在默认情况下,如果没有调用设备管理函数,主机端线程将会在运行第一个运行时函数时自动使用设备0.

CUDA runtime API通过设备管理功能对多个设备进行管理。由CUDA运行时API管理多设备,需要使用多个主机端线程。每个主机端线程在第一次调用其它CUDA运行时API函数之前,必须先由设备管理函数cudaSetDevice()与一个设备关联,并且以后也不能再次调用cudaSetDevice()函数与其它设备关联。主机端线程的数量可以多于设备数量,但一个时刻一个设备上只有一个主机端线程的上下文。为了达到最高性能,最好使主机端线程数量与设备数量相同,每个线程与设备一一对应。

通过CUDA驱动API管理多设备与多个上下文要略微复杂一些。CUDA驱动API通过上下文管理功能将上下文与主机端线程关联,一个线程在一个时刻只能有一个与之关联的上下文。

(2)、CUDA与OpenMP:除了直接使用操作系统提供的API管理多线程外,CUDA也可以与OpenMP一起使用。

(3)、CUDA与集群:MPI(MessagePassing Interface, 消息传递接口)是国际上最流行的并行编程开发环境。CUDA也可以与MPI联用,实现集群或者超级计算机中的多节点多GPU并行计算。

20.  测量程序运行时间

CUDA的内核程序运行时间可以在设备端测量,也可以在主机端测量。而CUDA API的运行时间则只能从主机端测量。无论是主机端测时还是设备端测时,最好都测量内核函数多次运行的时间,然后再除以运行次数以获得更加准确的结果。使用CUDA runtime API时,会在第一次调用runtime API函数时启动CUDA环境,为了避免将这一部分时间计入,最好在正式测时开始前先进行一次包含数据输入输出的计算,这样也可以使GPU从平时的节能模式进入工作状态,使测时结果更加可靠。

(1)、设备端测时:使用GPU中的计时器的时戳计时。实现设备端测时有两种不同的方法,分别是调用clock()函数和使用CUDA API的事件管理功能。

使用clock()函数计时,在内核函数中要测量的一段代码的开始和结束的位置分别调用一次clock()函数,并将结果记录下来。由于调用__syncthreads()函数后,一个block中的所有thread需要的时间是相同的,因此只需要记录每个block执行需要的时间就行了,而不需要记录每个thread的时间。Clock()函数的返回值的单位是GPU的时钟周期,需要除以GPU的运行频率才能得到以秒为单位的时间。

在设备端执行clock()函数,将返回每一个多处理器的时间计数器中的值。该时间计数器在每一个时钟周期递增1.在内核启动和结束时对时间计数器取样,比较两个值,并由每个线程记录各自的结果,就可以知道每个线程在多处理器上运行了多长时间。但是这并不是每个线程在多处理器上实际执行的时间。实际执行的时间比按照上述测试得到的时间短,因为多处理器上的执行时间是由多个线程按照时间分片共享的。

(2)、主机端测时:与普通程序测时一样,CUDA的主机端测时也采用CPU的计时器测时。通常取得CPU中计时器的值的方法是调用汇编中的相应指令,或者操作系统提供的API。此外,一些函数库,如C标准库中的time库的clock_t()函数也可以用来测时。不过,clock_t()函数的精度很低,建议在两次调用clock_t()时,让待测程序运行至少数十次,运行时间达到数秒,再取平均求得每次运行时间。

使用CPU测时,一定要牢记CUDA API的函数都是异步的。这就是说,在一个CUDA API函数在GPU上执行完成之前,CPU线程就已经得到了它的返回值。内核函数和带有asyn后缀的存储器拷贝函数都是异步的。

要从主机端准备的测量一个或者一系列CUDA调用需要的时间,就要先调用cudaThreadSynchronize()函数,同步CPU线程与GPU之后,才能结束CPU测时。cudaThreadSynchronize()函数的功能是阻塞CPU线程,直到cudaThreadSynchronize()函数之前所有的CUDA调用都已经完成。

与cudaThreadSynchronize()函数类似的函数有cudaStreamSynchronize()和cudaEventSynchronize()。它们的作用是阻塞所有Stream/CUDA Events,直到这条函数前的所有CUDA调用都已完成。注意,同一串流中的各个流可能会交替执行,因此即使使用了cudaStreamSynchronize()函数,也很难测得准确的执行时间。不过,一串流中的第一个流(ID为0的流)的行为总是同步的,因此使用这些函数对0号流进行测试,得到的结果是可靠的。

21.  CUDA函数库

(1)、cuFFT(CUDA Fast Fourier Transform):是一个利用GPU进行傅里叶变换的函数库,提供了与广泛使用的FFTW库相似的接口。

(2)、cuSparse:稀疏矩阵运算。

(3)、cuDNN:深度学习网络库。

(4)、cuBlas(CUDA Basic Linear Algebra Subprograms):线性代数函数库,是一个基本的矩阵与向量的运算库,提供了与BLAS相似的接口,可以用于简单的矩阵计算,也可以作为基础构建更加复杂的函数包。

(5)、cuRand:随机数生成库。

(6)、cuDpp(CUDA Data Parallel Primitives):提供了很多基本的常用的并行操作函数,如排序、搜索等,可以作为基本组件快速地搭建出并行计算程序。

22.  注意事项

(1)、在GPU上进行整数的除法和求模非常慢,避免这些运算能够有效地提高程序效率。

(2)、通常,block的数量都应该至少是处理核心的数量的几倍,才能有效地发挥GPU的处理能力。

(3)、在开发CUDA程序时应尽量避免分支,并尽量做到warp内不分支,否则将会导致性能急剧下降。

23.  CUDA Toolkit

Toolkit是CUDA的核心软件包,打开toolkit的安装目录,如C:\ProgramFiles\NVIDIA GPU Computing Toolkit\CUDA\v7.5,此目录下主要目录介绍:

(1)、bin目录:包含一些工具程序如nvcc.exe(CUDAC编译器)、ptxas.exe(ptx转机器码);一些动态链接库文件,包含w32和x64,如cudart64_75.dll(CUDA运行时API动态链接库)。

(2)、doc目录:里面包含了各种文档,包括pdf和html,可以根据实际需要查看相关文档说明。

(3)、include目录:包含常用的头文件,如cuda.h(CUDA驱动API头文件)。

(4)、lib目录:包含静态链接库,包含win32和x64,如cuda.lib(CUDA驱动库)、cudart.lib(CUDA运行时库)。

24.  Samples

在C:\ProgramData\NVIDIACorporation\CUDA Samples\v7.5 目录下包含了很多CUDA例子程序,对进一步掌握CUDA很有帮助。

五、CUDA架构

NVIDIA GPU是基于CUDA架构而构建的。可以将CUDA架构视为NVIDIA构建GPU的模式,其中GPU既可以完成传统的图形渲染任务,又可以完成通用计算任务。要在CUDA GPU上编程,需要使用CUDA C语言。

CUDA架构包含了一个统一的着色器流水线,使得执行通用计算的程序能够对芯片上的每个数学逻辑单元(Arithmetic Logic Unit, ALU)进行排列。由于NVIDIA希望使新的图形处理器能适应于通用计算,因此在实现这些ALU时都确保它们满足IEEE单精度浮点数学运算的需求,并且可以使用一个裁剪后的指令集来执行通用计算,而不是仅限于执行图形计算。此外,GPU上的执行单元不仅能任意地读/写内存,同时还能访问由软件管理的缓存,也称为共享内存。CUDA架构的所有这些功能都是为了使GPU不仅能执行传统的图形计算,还能高效地执行通用计算。

NVIDIA采取工业标准的C语言,并且增加了一小部分关键字来支持CUDA架构的特殊功能。NVIDIA公布了一款编译器来编译CUDA C语言。这样,CUDA C就成为了第一款专门由GPU公司设计的编程语言,用于在GPU上编写通用计算。

除了专门设计一种语言来为GPU编写代码之外,NVIDIA还提供了专门的硬件驱动程序来发挥CUDA架构的大规模计算功能。

六、NVCC编译器

NVCC编译器根据配置编译CUDA C代码,可以生成三种不同的输出:PTX、CUDA二进制序列和标准C。nvcc是一种编译器驱动,通过命令行选项,nvcc可以在编译的不同阶段启动不同的工具完成编译工作。

nvcc工作的基本流程是:首先通过CUDAfe分离源文件中的主机端和设备端代码,然后再调用不同的编译器分别编译。设备端代码由nvcc编译成ptx代码或者二进制代码;主机端代码则将以C文件形式输出,由其他高性能编译器,如ICC、GCC或者其他合适的高性能编译器等进行编译。不过,也可以直接在编译的最后阶段,将主机端代码交给其他编译器生成.obj或者.o文件。在编译时,可以将设备端代码链接到所生成的主机端代码,将其中的cubin对象作为全局初始化数据数组包含进来。此时,内核执行配置也要被转换为CUDA运行启动代码,以加载和启动编译后的内核函数。使用CUDA驱动API时,可以单独执行ptx代码或者cubin对象,而忽略nvcc编译得到的主机端代码。

nvcc大概的编译流程如下图:

PTX(Parallel Thread eXecution)类似于汇编语言,是为动态编译器JIT(Just in time compiler, JIT包含在标准的NVIDIA驱动中)设计的输入指令序列。这样,虽然不同的显卡使用的机器语言不同,JIT却可以运行同样的PTX。这样做使PTX成为一个稳定的接口,带来了很多好处:向后兼容性、更长的寿命、更好的可扩展性和更高的性能,但在一定程度上也限制了工程上的自由发挥。这种技术保证了兼容性,但也使新一代的产品必须拥有上代产品的所有能力,这样才能让今天的PTX代码在未来的系统上仍然可以运行。

编译器前端按照C++语法规则对CUDA源文件进行处理。CUDA主机端代码可以支持完整的C++语法,而设备端代码则不能完全支持。

内核函数可以通过PTX编写,但通常还是通过CUDA C一类的高级语言进行编写。PTX或CUDA C语言编写的内核函数都必须通过nvcc编译器编译成二进制代码。一部分PTX指令只能在拥有较高计算能力的硬件上执行。nvcc通过-arch编译选项来指定要输出的PTX代码的计算能力。

在程序编译时,要使目标代码和目标硬件版本与实际使用的硬件一致,可以使用-arch、-gencode和-code编译选项。

关于nvcc编译选项的更详细信息可以参考:C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v7.5\doc\html\cuda-compiler-driver-nvcc

以上部分内容整理自:《GPU高性能运算之CUDA》、《GPU高性能编程CUDA实战》

cuda资料整理(2)

CUDA——”从入门到放弃” https://www.jianshu.com/p/34a504af8d51

1. 知识准备

1.1 中央处理器(CPU)

中央处理器(CPU,Central Processing Unit)是一块超大规模的集成电路,是一台计算机的运算核心(Core)和控制核心( Control Unit)。它的功能主要是解释计算机指令以及处理计算机软件中的数据。
中央处理器主要包括运算器(算术逻辑运算单元,ALU,Arithmetic Logic Unit)和高速缓冲存储器(Cache)及实现它们之间联系的数据(Data)、控制及状态的总线(Bus)。它与内部存储器(Memory)和输入/输出(I/O)设备合称为电子计算机三大核心部件。

CPU的结构主要包括运算器(ALU, Arithmetic and Logic Unit)、控制单元(CU, Control Unit)、寄存器(Register)、高速缓存器(Cache)和它们之间通讯的数据、控制及状态的总线

简单来说就是:计算单元、控制单元和存储单元,架构如下图所示:

CPU微架构示意图

什么?架构记不住?来,我们换种表示方法:

CPU微架构示意图(改)

嗯,大概就是这个意思。

从字面上我们也很好理解,计算单元主要执行算术运算、移位等操作以及地址运算和转换;存储单元主要用于保存运算中产生的数据以及指令等;控制单元则对指令译码,并且发出为完成每条指令所要执行的各个操作的控制信号。

所以一条指令在CPU中执行的过程是这样的:读取到指令后,通过指令总线送到控制器(黄色区域)中进行译码,并发出相应的操作控制信号;然后运算器(绿色区域)按照操作指令对数据进行计算,并通过数据总线将得到的数据存入数据缓存器(大块橙色区域)。过程如下图所示:

CPU执行指令图

是不是有点儿复杂?没关系,这张图完全不用记住,我们只需要知道,CPU遵循的是冯诺依曼架构,其核心就是:存储程序,顺序执行

讲到这里,有没有看出问题,没错——在这个结构图中,负责计算的绿色区域占的面积似乎太小了,而橙色区域的缓存Cache和黄色区域的控制单元占据了大量空间。

高中化学有句老生常谈的话叫:结构决定性质,放在这里也非常适用。

因为CPU的架构中需要大量的空间去放置存储单元(橙色部分)和控制单元(黄色部分),相比之下计算单元(绿色部分)只占据了很小的一部分,所以它在大规模并行计算能力上极受限制,而更擅长于逻辑控制。

另外,因为遵循冯诺依曼架构(存储程序,顺序执行),CPU就像是个一板一眼的管家,人们吩咐的事情它总是一步一步来做。但是随着人们对更大规模与更快处理速度的需求的增加,这位管家渐渐变得有些力不从心。

于是,大家就想,能不能把多个处理器放在同一块芯片上,让它们一起来做事,这样效率不就提高了吗?

没错,GPU便由此诞生了。

1.2 显卡

显卡(Video card,Graphics card)全称显示接口卡,又称显示适配器,是计算机最基本配置、最重要的配件之一。显卡作为电脑主机里的一个重要组成部分,是电脑进行数模信号转换的设备,承担输出显示图形的任务。显卡接在电脑主板上,它将电脑的数字信号转换成模拟信号让显示器显示出来,同时显卡还是有图像处理能力,可协助CPU工作,提高整体的运行速度。对于从事专业图形设计的人来说显卡非常重要。 民用和军用显卡图形芯片供应商主要包括AMD(超微半导体)Nvidia(英伟达)2家。现在的top500计算机,都包含显卡计算核心。在科学计算中,显卡被称为显示加速卡

为什么GPU特别擅长处理图像数据呢?这是因为图像上的每一个像素点都有被处理的需要,而且每个像素点处理的过程和方式都十分相似,也就成了GPU的天然温床。

GPU微架构示意图

从架构图我们就能很明显的看出,GPU的构成相对简单,有数量众多的计算单元和超长的流水线,特别适合处理大量的类型统一的数据。

再把CPU和GPU两者放在一张图上看下对比,就非常一目了然了。

GPU的工作大部分都计算量大,但没什么技术含量,而且要重复很多很多次。

但GPU无法单独工作,必须由CPU进行控制调用才能工作。CPU可单独作用,处理复杂的逻辑运算和不同的数据类型,但当需要大量的处理类型统一的数据时,则可调用GPU进行并行计算。

借用知乎上某大佬的说法,就像你有个工作需要计算几亿次一百以内加减乘除一样,最好的办法就是雇上几十个小学生一起算,一人算一部分,反正这些计算也没什么技术含量,纯粹体力活而已;而CPU就像老教授,积分微分都会算,就是工资高,一个老教授资顶二十个小学生,你要是富士康你雇哪个?

注:GPU中有很多的运算器ALU和很少的缓存cache,缓存的目的不是保存后面需要访问的数据的,这点和CPU不同,而是为线程thread提高服务的。如果有很多线程需要访问同一个相同的数据,缓存会合并这些访问,然后再去访问dram。

可爱的你如果对CUDA硬件有更多的兴趣,可移步NVIDIA中文官网进一步学习。

1.3 内存

内存是计算机中重要的部件之一,它是与CPU进行沟通的桥梁。计算机中所有程序的运行都是在内存中进行的,因此内存的性能对计算机的影响非常大。内存(Memory)也被称为内存储器,其作用是用于暂时存放CPU中的运算数据,以及与硬盘外部存储器交换的数据。只要计算机在运行中,CPU就会把需要运算的数据调到内存中进行运算,当运算完成后CPU再将结果传送出来,内存的运行也决定了计算机的稳定运行。 内存是由内存芯片、电路板、金手指等部分组成的。

1.4 显存

显存,也被叫做帧缓存,它的作用是用来存储显卡芯片处理过或者即将提取的渲染数据。如同计算机的内存一样,显存是用来存储要处理的图形信息的部件。

1.5 显卡、显卡驱动、CUDA之间的关系

显卡:(GPU)主流是NVIDIA的GPU,深度学习本身需要大量计算。GPU的并行计算能力,在过去几年里恰当地满足了深度学习的需求。AMD的GPU基本没有什么支持,可以不用考虑。

驱动:没有显卡驱动,就不能识别GPU硬件,不能调用其计算资源。但是呢,NVIDIA在Linux上的驱动安装特别麻烦,尤其对于新手简直就是噩梦。得屏蔽第三方显卡驱动。下面会给出教程。

CUDA:是NVIDIA推出的只能用于自家GPU的并行计算框架。只有安装这个框架才能够进行复杂的并行计算。主流的深度学习框架也都是基于CUDA进行GPU并行加速的,几乎无一例外。还有一个叫做cudnn,是针对深度卷积神经网络的加速库。

查看显卡驱动信息(以实验室服务器为例)

ssh ubuntu@192.168.1.158

输入服务器密码登陆
然后,进入cuda

cd /usr/local/cuda-8.0/samples/1_Utilities/deviceQuery

运行其中的可执行文件

./deviceQuery

得到如下信息

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 4 CUDA Capable device(s)

Device 0: "GeForce GTX 1080 Ti"
  CUDA Driver Version / Runtime Version          9.0 / 8.0
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 11171 MBytes (11713708032 bytes)
  (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
  GPU Max Clock rate:                            1620 MHz (1.62 GHz)
  Memory Clock rate:                             5505 Mhz
  Memory Bus Width:                              352-bit
  L2 Cache Size:                                 2883584 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 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 2 copy engine(s)
  Run time limit on kernels:                     No
  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 PCI Domain ID / Bus ID / location ID:   0 / 2 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

Device 1: "GeForce GTX 1080 Ti"
  CUDA Driver Version / Runtime Version          9.0 / 8.0
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 11172 MBytes (11715084288 bytes)
  (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
  GPU Max Clock rate:                            1620 MHz (1.62 GHz)
  Memory Clock rate:                             5505 Mhz
  Memory Bus Width:                              352-bit
  L2 Cache Size:                                 2883584 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 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 2 copy engine(s)
  Run time limit on kernels:                     No
  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 PCI Domain ID / Bus ID / location ID:   0 / 3 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

Device 2: "GeForce GTX 1080 Ti"
  CUDA Driver Version / Runtime Version          9.0 / 8.0
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 11172 MBytes (11715084288 bytes)
  (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
  GPU Max Clock rate:                            1620 MHz (1.62 GHz)
  Memory Clock rate:                             5505 Mhz
  Memory Bus Width:                              352-bit
  L2 Cache Size:                                 2883584 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 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 2 copy engine(s)
  Run time limit on kernels:                     No
  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 PCI Domain ID / Bus ID / location ID:   0 / 130 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

Device 3: "GeForce GTX 1080 Ti"
  CUDA Driver Version / Runtime Version          9.0 / 8.0
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 11172 MBytes (11715084288 bytes)
  (28) Multiprocessors, (128) CUDA Cores/MP:     3584 CUDA Cores
  GPU Max Clock rate:                            1620 MHz (1.62 GHz)
  Memory Clock rate:                             5505 Mhz
  Memory Bus Width:                              352-bit
  L2 Cache Size:                                 2883584 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 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 2 copy engine(s)
  Run time limit on kernels:                     No
  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 PCI Domain ID / Bus ID / location ID:   0 / 131 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
> Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU1) : Yes
> Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU2) : No
> Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU3) : No
> Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU0) : Yes
> Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU2) : No
> Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU3) : No
> Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU0) : No
> Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU1) : No
> Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU3) : Yes
> Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU0) : No
> Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU1) : No
> Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU2) : Yes

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 9.0, CUDA Runtime Version = 8.0, NumDevs = 4, Device0 = GeForce GTX 1080 Ti, Device1 = GeForce GTX 1080 Ti, Device2 = GeForce GTX 1080 Ti, Device3 = GeForce GTX 1080 Ti
Result = PASS

大家可以在自己PC或者工作机上尝试一下。

再啰嗦两句

GPU就是用很多简单的计算单元去完成大量的计算任务,纯粹的人海战术。这种策略基于一个前提,就是小学生A和小学生B的工作没有什么依赖性,是互相独立的。

但有一点需要强调,虽然GPU是为了图像处理而生的,但是我们通过前面的介绍可以发现,它在结构上并没有专门为图像服务的部件,只是对CPU的结构进行了优化与调整,所以现在GPU不仅可以在图像处理领域大显身手,它还被用来科学计算、密码破解、数值分析,海量数据处理(排序,Map-Reduce等),金融分析等需要大规模并行计算的领域。

所以GPU也可以认为是一种较通用的芯片。

2. CUDA软件构架

CUDA是一种新的操作GPU计算的硬件和软件架构,它将GPU视作一个数据并行计算设备,而且无需把这些计算映射到图形API。操作系统的多任务机制可以同时管理CUDA访问GPU和图形程序的运行库,其计算特性支持利用CUDA直观地编写GPU核心程序。目前Tesla架构具有在笔记本电脑、台式机、工作站和服务器上的广泛可用性,配以C/C++语言的编程环境和CUDA软件,使这种架构得以成为最优秀的超级计算平台。

CUDA软件层次结构

CUDA在软件方面组成有:一个CUDA库、一个应用程序编程接口(API)及其运行库(Runtime)、两个较高级别的通用数学库,即CUFFT和CUBLAS。CUDA改进了DRAM的读写灵活性,使得GPU与CPU的机制相吻合。另一方面,CUDA提供了片上(on-chip)共享内存,使得线程之间可以共享数据。应用程序可以利用共享内存来减少DRAM的数据传送,更少的依赖DRAM的内存带宽。

3. 编程模型

CUDA程序构架分为两部分:HostDevice。一般而言,Host指的是CPUDevice指的是GPU。在CUDA程序构架中,主程序还是由CPU来执行,而当遇到数据并行处理的部分,CUDA 就会将程序编译成GPU能执行的程序,并传送到GPU。而这个程序在CUDA里称做(kernel)。CUDA允许程序员定义称为核的C语言函数,从而扩展了C语言,在调用此类函数时,它将由N个不同的CUDA线程并行执行N次,这与普通的C语言函数只执行一次的方式不同。执行核的每个线程都会被分配一个独特的线程ID,可通过内置的threadIdx变量在内核中访问此ID。在 CUDA 程序中,主程序在调用任何GPU内核之前,必须对核进行执行配置,即确定线程块数和每个线程块中的线程数以及共享内存大小。

3.1 线程层次结构

在GPU中要执行的线程,根据最有效的数据共享来创建块(Block),其类型有一维、二维或三维。在同一个块内的线程可彼此协作,通过一些共享存储器来共享数据,并同步其执行来协调存储器访问。一个块中的所有线程都必须位于同一个处理器核心中。因而,一个处理器核心的有限存储器资源制约了每个块的线程数量。在早期的NVIDIA 架构中,一个线程块最多可以包含 512个线程,而在后期出现的一些设备中则最多可支持1024个线程。一般GPGPU程序线程数目是很多的,所以不能把所有的线程都塞到同一个块里。但一个内核可由多个大小相同的线程块同时执行,因而线程总数应等于每个块的线程数乘以块的数量。这些同样维度和大小的块将组织为一个一维或二维线程块网格(Grid)。具体框架如下图所示。

线程块网格

NOTICE:

线程(Thread)
一般通过GPU的一个核进行处理。(可以表示成一维,二维,三维,具体下面再细说)。
线程块(Block)

  1. 由多个线程组成(可以表示成一维,二维,三维,具体下面再细说)。
  2. 各block是并行执行的,block间无法通信,也没有执行顺序。
  3. 注意线程块的数量限制为不超过65535(硬件限制)。

线程格(Grid)
由多个线程块组成(可以表示成一维,二维,三维,具体下面再细说)。
线程束
在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行,线程束中的每个线程都将在不同数据上执行相同的命令。

从硬件上看

SP:最基本的处理单元,streaming processor,也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理。
SM:多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核,其他资源如:warp scheduler,register,shared memory等。SM可以看做GPU的心脏(对比CPU核心),register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此,这些有限的资源就使每个SM中active warps有非常严格的限制,也就限制了并行能力。

从软件上看

thread:一个CUDA的并行程序会被以许多个threads来执行。
block:数个threads会被群组成一个block,同一个block中的threads可以同步,也可以通过shared memory通信。
grid:多个blocks则会再构成grid。
warp:GPU执行程序时的调度单位,目前cuda的warp的大小为32,同在一个warp的线程,以不同数据资源执行相同的指令,这就是所谓 SIMT。

3.2 存储器层次结构

CUDA设备拥有多个独立的存储空间,其中包括:全局存储器、本地存储器、共享存储器、常量存储器、纹理存储器和寄存器,如图

CUDA设备上的存储器
NOTICE:

主机(Host)
将CPU及系统的内存(内存条)称为主机。
设备(Device)
将GPU及GPU本身的显示内存称为设备。
动态随机存取存储器(DRAM)
DRAM(Dynamic Random Access Memory),即动态随机存取存储器,最为常见的系统内存DRAM只能将数据保持很短的时间。为了保持数据,DRAM使用电容存储,所以必须隔一段时间刷新(refresh)一次,如果存储单元没有被刷新,存储的信息就会丢失。 (关机就会丢失数据)

CUDA线程可在执行过程中访问多个存储器空间的数据,如下图所示其中:

  • 每个线程都有一个私有的本地存储器。
  • 每个线程块都有一个共享存储器,该存储器对于块内的所有线程都是可见的,并且与块具有相同的生命周期。
  • 所有线程都可访问相同的全局存储器。
  • 此外还有两个只读的存储器空间,可由所有线程访问,这两个空间是常量存储器空间和纹理存储器空间。全局、固定和纹理存储器空间经过优化,适于不同的存储器用途。纹理存储器也为某些特殊的数据格式提供了不同的寻址模式以及数据过滤,方便Host对流数据的快速存取。

存储器的应用层次

3.3 主机(Host)和设备(Device)

如下图所示,CUDA假设线程可在物理上独立的设备上执行,此类设备作为运行C语言程序的主机的协处理器操作。内核在GPU上执行,而C语言程序的其他部分在CPU上执行(即串行代码在主机上执行,而并行代码在设备上执行)。此外,CUDA还假设主机和设备均维护自己的DRAM,分别称为主机存储器和设备存储器。因而,一个程序通过调用CUDA运行库来管理对内核可见的全局、固定和纹理存储器空间。这种管理包括设备存储器的分配和取消分配,还包括主机和设备存储器之间的数据传输。

4. CUDA软硬件

4.1 CUDA术语

由于CUDA中存在许多概念和术语,诸如SM、block、SP等多个概念不容易理解,将其与CPU的一些概念进行比较,如下表所示。

CPUGPU层次
算术逻辑和控制单元流处理器(SM)硬件
算术单元批量处理器(SP)硬件
进程Block软件
线程thread软件
调度单位Warp软件

4.2 硬件利用率

当为一个GPU分配一个内核函数,我们关心的是如何才能充分利用GPU的计算能力,但由于不同的硬件有不同的计算能力,SM一次最多能容纳的线程数也不尽相同,SM一次最多能容纳的线程数量主要与底层硬件的计算能力有关,如下表显示了在不同的计算能力的设备上,每个线程块上开启不同数量的线程时设备的利用率。

计算能力 每个线 程块的线程数1.01.11.21.32.02.13.0
6467505050333350
961001007575505075
1281001001001006767100
192100100949410010094
96100100100100100100100
······

查看显卡利用率 (以实验室服务器为例)
输入以下命令

nvidia-smi

Thu Aug 23 21:06:36 2018       
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 384.130                Driver Version: 384.130                   |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  GeForce GTX 108...  Off  | 00000000:02:00.0 Off |                  N/A |
| 29%   41C    P0    58W / 250W |      0MiB / 11171MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   1  GeForce GTX 108...  Off  | 00000000:03:00.0 Off |                  N/A |
| 33%   47C    P0    57W / 250W |      0MiB / 11172MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   2  GeForce GTX 108...  Off  | 00000000:82:00.0 Off |                  N/A |
| 36%   49C    P0    59W / 250W |      0MiB / 11172MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   3  GeForce GTX 108...  Off  | 00000000:83:00.0 Off |                  N/A |
| 33%   46C    P0    51W / 250W |      0MiB / 11172MiB |      1%      Default |
+-------------------------------+----------------------+----------------------+
                                                                               
+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID   Type   Process name                             Usage      |
|=============================================================================|
|  No running processes found                                                 |
+-----------------------------------------------------------------------------+

5. 并行计算

5.1 并发性

CUDA将问题分解成线程块的网格,每块包含多个线程。快可以按任意顺序执行。不过在某个时间点上,只有一部分块处于执行中。一旦被调用到GUP包含的N个“流处理器簇(SM)”中的一个上执行,一个块必须从开始到结束。网格中的块可以被分配到任意一个有空闲槽的SM上。起初,可以采用“轮询调度”策略,以确保分配到每一个SM上的块数基本相同。对绝大多数内核程序而言,分块的数量应该是GPU中物理SM数量的八倍或更多倍。

以一个军队比喻,假设有一支由士兵(线程)组成的部队(网格)。部队被分成若干个连(块),每个连队由一位连长来指挥。按照32名士兵一个班(一个线程束),连队又进一步分成若干个班,每个班由一个班长来指挥。

基于GPU的线程视图

要执行某个操作,总司令(内核程序/ 主机程序)必须提供操作名称及相应的数据。每个士兵(线程)只处理分配给他的问题中的一小块。在连长(负责一个块)或班长(负责一个束)的控制下,束与束之间的线程或者一个束内部的线程之间,要经常地交换数据。但是,连队(块)之间的协同就得由总司令(内核函数/ 主机程序)来控制。

5.2 局部性

对于GPU程序设计,程序员必须处理局部性。对于一个给定的工作,他需要事先思考需要哪些工具或零件(即存储地址或数据结构),然后一次性地把他们从硬件仓库(全局内存)可能把与这些数据相关的不同工作都执行了,避免发生“取来–存回–为了下一个工作再取”。

5.3 缓存一致性

GPU与CPU在缓存上的一个重要差别就是“缓存一致性”问题。对于“缓存一致”的系统,一个内存的写操作需要通知所有核的各个级别的缓存。因此,无论何时,所有的处理器核看到的内存视图是完全一样的。随着处理器中核数量的增多,这个“通知”的开销迅速增大,使得“缓存一致性”成为限制一个处理器中核数量不能太多的一重要因素。“缓存一致”系统中最坏的情况是,一个内存操作会强迫每个核的缓存都进行更新,进而每个核都要对相邻的内存单元写操作。

相比之下,非“缓存一致”系统不会自动地更新其他核的缓存。它需要由程序员写清楚每个处理器核输出的各自不同的目标区域。从程序的视角看,这支持一个核仅负责一个输出或者一个小的输出集。通常,CPU遵循“缓存一致性”原则,而GPU则不是。故GPU能够扩展到一个芯片内具有大数量的核心(流处理器簇)。

5.4 弗林分类法

根据弗林分类法,计算机的结构类型有:

SIMD–单指令,多数据
MIMD–多指令,多数据
SISD–单指令,单数据
MISD–多指令,单数据

5.5 分条 / 分块

CUDA提供的简单二维网格模型。对于很多问题,这样的模型就足够了。如果在一个块内,你的工作是线性分布的,那么你可以很好地将其他分解成CUDA块。由于在一个SM内,最多可以分配16个块,而在一个GPU内有16个(有些是32个)SM,所以问题分成256个甚至更多的块都可以。实际上,我们更倾向于把一个块内的元素总数限制为128、256、或者512,这样有助于在一个典型的数据集内划分出更多数量的块。

5.6 快速傅氏变换(FFT)

FFT: FFT(Fast Fourier Transformation)是离散傅氏变换(DFT)的快速算法。即为快速傅氏变换。它是根据离散傅氏变换的奇、偶、虚、实等特性,对离散傅立叶变换的算法进行改进获得的。

由于不是刚需,这里不展开讲。好奇的你可以点击楼下时光机,通过下面的教程进行学习。
FFT(最详细最通俗的入门手册)

5.7 CUDA计算能力的含义

体现GPU计算能力的两个重要特征:
1)CUDA核的个数;
2)存储器大小。
描述GPU性能的两个重要指标: :
1)计算性能峰值;
2)存储器带宽。

参考
1.CUDA计算能力的含义
2.CUDA GPUs

6. 实践

6.1 Ubuntu 系统下环境搭建

6.1.1 系统要求

要搭建 CUDA 环境,我们需要自己的计算机满足以下这三个条件:
1. 有至少一颗支持 CUDA 的 GPU(我的是GeForece GT 650M)
2. 有满足版本要求的 gcc 编译器和链接工具
3. 有 NVIDIA 提供的 CUDA 工具包(点击神奇的小链接下载)

6.1.2 准备工作

下面,我们一步一步来验证自己的系统是否满足安装要求。
Step 1: 验证计算机是否拥有至少一颗支持 CUDA 的 GPU
打开终端(Ctrl + Alt + T),键入以下命令:

lspci | grep -i nvidia

可以看到以下内容(结果因人而异,与具体的GPU有关)

看到这个就说明至少有一颗支持 CUDA 的 GPU,可以进入下一步了。

Step 2: 验证一下自己操作系统的版本
键入命令:

lsb_release -a

No LSB modules are available.
Distributor ID: Ubuntu
Description:    Ubuntu 16.04.4 LTS
Release:    16.04
Codename:   xenial

更多信息请移步Ubuntu查看版本信息

Step 3: 验证 gcc 编译器的版本
键入命令:

gcc --version

或者

gcc -v

得到如下信息

gcc (Ubuntu 5.4.0-6ubuntu1~16.04.10) 5.4.0 20160609
Copyright (C) 2015 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

Step 4: 验证系统内核版本
键入命令:

uname -r

得到如下信息

对照官方提供的对各种 Linux 发行版的安装要求进行安装

6.1.3 搭建 CUDA 环境

Step 1: 安装 CUDA 工具包
在前面几项验证都顺利通过以后就来到最关键的一步。首先下载对应自己系统版本的 CUDA 工具包(以CUDA Toolkit 9.2 为例),然后进入到安装包所在目录:

sudo dpkg -i cuda-repo-ubuntu1604-9-2-local_9.2.148-1_amd64.deb

sudo apt-key add /var/cuda-repo-<version>/7fa2af80.pub

sudo apt-get update

sudo apt-get install cuda

NOTICE:

Other installation options are available in the form of meta-packages. For example, to install all the library packages, replace “cuda” with the “cuda-libraries-9-2” meta package. For more information on all the available meta packages click here.

此时静静地等待安装完成。不出意外,一段时间后安装完成了。
Step 2: 设置环境变量
首先在 PATH 变量中加入 /usr/local/cuda-9.2/bin,在Terminal中执行:

export PATH=/usr/local/cuda-9.2/bin:$PATH

然后在 LD_LIBRARY_PATH 变量中添加 /usr/local/cuda-9.2/lib64,执行:

export  LD_LIBRARY_PATH=/usr/local/cuda-9.2/lib64:$LD_LIBRARY_PATH

Step 3: 验证环境搭建是否成功
首先执行命令:

nvcc -V

关于测试…聪明的你一定想起来了,我们前面是讲过怎么做的。
对,没错,就在1.5小节,话不多说,自行上翻吧。

看到通过测试,到这里,64位 Ubuntu 16.04 系统下 CUDA 环境搭建就完成了。

6.2 CUDA编程

6.2.1 核函数

1. 在GPU上执行的函数通常称为核函数。
2. 一般通过标识符__global__修饰,调用通过<<<参数1,参数2>>>,用于说明内核函数中的线程数量,以及线程是如何组织的。
3. 以线程格(Grid)的形式组织,每个线程格由若干个线程块(block)组成,而每个线程块又由若干个线程(thread)组成。
4.是以block为单位执行的。
5. 叧能在主机端代码中调用。
6. 调用时必须声明内核函数的执行参数。
7. 在编程时,必须先为kernel函数中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误,例如越界或报错,甚至导致蓝屏和死机。

看完基本知识,装好CUDA以后,就可以开始写第一个CUDA程序了:

#include <cuda_runtime.h>
 
int main(){
printf("Hello world!\n");
}

慢着,这个程序和C有什么区别?用到显卡了吗?
答:没有区别,没用显卡。如果你非要用显卡干点什么事情的话,可以改成这个样子:

/*
 * @file_name HelloWorld.cu  后缀名称.cu
 */

#include <stdio.h>
#include <cuda_runtime.h>  //头文件

//核函数声明,前面的关键字__global__
__global__ void kernel( void ) {
}

int main( void ) {
    //核函数的调用,注意<<<1,1>>>,第一个1,代表线程格里只有一个线程块;第二个1,代表一个线程块里只有一个线程。
    kernel<<<1,1>>>();
    printf( "Hello, World!\n" );
    return 0;
}

6.2.2 dim3结构类型

  1. dim3是基于uint3定义的矢量类型,相当亍由3个unsigned int型组成的结构体。uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z;
  2. 可使用于一维、二维或三维的索引来标识线程,构成一维、二维或三维线程块。
  3. dim3结构类型变量用在核函数调用的<<<,>>>中。
  4. 相关的几个内置变量
    4.1. threadIdx,顾名思义获取线程thread的ID索引;如果线程是一维的那么就取threadIdx.x,二维的还可以多取到一个值threadIdx.y,以此类推到三维threadIdx.z
    4.2. blockIdx,线程块的ID索引;同样有blockIdx.xblockIdx.yblockIdx.z
    4.3. blockDim,线程块的维度,同样有blockDim.xblockDim.yblockDim.z
    4.4. gridDim,线程格的维度,同样有gridDim.xgridDim.ygridDim.z
  5. 对于一维的block,线程的threadID=threadIdx.x
  6. 对于大小为(blockDim.x, blockDim.y)的 二维block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x
    1. 对于大小为(blockDim.x, blockDim.y, blockDim.z)的 三维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y
    2. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride

6.2.3 函数修饰符

1.__global__,表明被修饰的函数在设备上执行,但在主机上调用。

  1. __device__,表明被修饰的函数在设备上执行,但只能在其他__device__函数或者__global__函数中调用。

6.2.4 常用的GPU内存函数

cudaMalloc()
1. 函数原型: cudaError_t cudaMalloc (void **devPtr, size_t size)
2. 函数用处:与C语言中的malloc函数一样,只是此函数在GPU的内存你分配内存。
3. 注意事项:
3.1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数;
3.2. 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作;
3.3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数;
3.4. 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作(即不能进行解引用)。

cudaMemcpy()
1. 函数原型:cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)
2. 函数作用:与c语言中的memcpy函数一样,只是此函数可以在主机内存和GPU内存之间互相拷贝数据。
3. 函数参数:cudaMemcpyKind kind表示数据拷贝方向,如果kind赋值为cudaMemcpyDeviceToHost表示数据从设备内存拷贝到主机内存。
4. 与C中的memcpy()一样,以同步方式执行,即当函数返回时,复制操作就已经完成了,并且在输出缓冲区中包含了复制进去的内容。
5. 相应的有个异步方式执行的函数cudaMemcpyAsync(),这个函数详解请看下面的流一节有关内容。

cudaFree()
1. 函数原型:cudaError_t cudaFree ( void* devPtr )
2. 函数作用:与c语言中的free()函数一样,只是此函数释放的是cudaMalloc()分配的内存。
下面实例用于解释上面三个函数

#include <stdio.h>
#include <cuda_runtime.h>
__global__ void add( int a, int b, int *c ) {
    *c = a + b;
}
int main( void ) {
    int c;
    int *dev_c;
    //cudaMalloc()
    cudaMalloc( (void**)&dev_c, sizeof(int) );
    //核函数执行
    add<<<1,1>>>( 2, 7, dev_c );   
    //cudaMemcpy()
    cudaMemcpy( &c, dev_c, sizeof(int),cudaMemcpyDeviceToHost ) ;
    printf( "2 + 7 = %d\n", c );
    //cudaFree()
    cudaFree( dev_c );
 
    return 0;
}

6.2.5 GPU内存分类

全局内存
通俗意义上的设备内存。

共享内存
1. 位置:设备内存。
2. 形式:关键字__shared__添加到变量声明中。如__shared__ float cache[10]
3. 目的:对于GPU上启动的每个线程块,CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存,但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

常量内存
1. 位置:设备内存
2. 形式:关键字__constant__添加到变量声明中。如__constant__ float s[10];。
3. 目的:为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下,用常量内存替换全局内存能有效地减少内存带宽。
4. 特点:常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时,静态地分配空间。
5. 要求:当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol(),而cudaMemcpy()会复制到全局内存。
6. 性能提升的原因:
6.1. 对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。(为什么是15,因为“邻近”指半个线程束,一个线程束包含32个线程的集合。)
6.2. 常量内存的数据将缓存起来,因此对相同地址的连续读操作将不会产生额外的内存通信量。

纹理内存
1. 位置:设备内存
2. 目的:能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计,意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图:

3. 纹理变量(引用)必须声明为文件作用域内的全局变量。
4. 形式:分为一维纹理内存 和 二维纹理内存。
4.1. 一维纹理内存
4.1.1. 用texture<类型>类型声明,如texture<float> texIn
4.1.2. 通过cudaBindTexture()绑定到纹理内存中。
4.1.3. 通过tex1Dfetch()来读取纹理内存中的数据。
4.1.4. 通过cudaUnbindTexture()取消绑定纹理内存。
4.2. 二维纹理内存
4.2.1. 用texture<类型,数字>类型声明,如texture<float,2> texIn
4.2.2. 通过cudaBindTexture2D()绑定到纹理内存中。
4.2.3. 通过tex2D()来读取纹理内存中的数据。
4.2.4. 通过cudaUnbindTexture()取消绑定纹理内存。

固定内存
1. 位置:主机内存。
2. 概念:也称为页锁定内存或者不可分页内存,操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址,因为这块内存将不会破坏或者重新定位。
3. 目的:提高访问速度。由于GPU知道主机内存的物理地址,因此可以通过“直接内存访问DMA(Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。
4. 缺点:使用固定内存,将失去虚拟内存的所有功能;系统将更快的耗尽内存。
5. 建议:对cudaMemcpy()函数调用中的源内存或者目标内存,才使用固定内存,并且在不再需要使用它们时立即释放。
6. 形式:通过cudaHostAlloc()函数来分配;通过cudaFreeHost()释放。
7. 只能以异步方式对固定内存进行复制操作。

原子性
1. 概念:如果操作的执行过程不能分解为更小的部分,我们将满足这种条件限制的操作称为原子操作。
2. 形式:函数调用,如atomicAdd(addr,y)将生成一个原子的操作序列,这个操作序列包括读取地址addr处的值,将y增加到这个值,以及将结果保存回地址addr

6.2.6 常用线程操作函数

同步方法__syncthreads(),这个函数的调用,将确保线程块中的每个线程都执行完__syscthreads()前面的语句后,才会执行下一条语句。

使用事件来测量性能
1. 用途:为了测量GPU在某个任务上花费的时间。CUDA中的事件本质上是一个GPU时间戳。由于事件是直接在GPU上实现的。因此不适用于对同时包含设备代码和主机代码的混合代码设计。
2. 形式:首先创建一个事件,然后记录事件,再计算两个事件之差,最后销毁事件。如:

cudaEvent_t start, stop;
cudaEventCreate( &start );
cudaEventCreate( &stop );
cudaEventRecord( start, 0 );
//do something
cudaEventRecord( stop, 0 );
float   elapsedTime;
cudaEventElapsedTime( &elapsedTime,start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

6.2.7 流

  1. 扯一扯:并发重点在于一个极短时间段内运行多个不同的任务;并行重点在于同时运行一个任务。
  2. 任务并行性:是指并行执行两个或多个不同的任务,而不是在大量数据上执行同一个任务。
  3. 概念:CUDA流表示一个GPU操作队列,并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作,如核函数启动,内存复制以及事件的启动和结束等。这些操作的添加到流的顺序也是它们的执行顺序。可以将每个流视为GPU上的一个任务,并且这些任务可以并行执行。
  4. 硬件前提:必须是支持设备重叠功能的GPU。支持设备重叠功能,即在执行一个核函数的同时,还能在设备与主机之间执行复制操作。
  5. 声明与创建:声明cudaStream_t stream;,创建cudaSteamCreate(&stream);。
  6. cudaMemcpyAsync():前面在cudaMemcpy()中提到过,这是一个以异步方式执行的函数。在调用cudaMemcpyAsync()时,只是放置一个请求,表示在流中执行一次内存复制操作,这个流是通过参数stream来指定的。当函数返回时,我们无法确保复制操作是否已经启动,更无法保证它是否已经结束。我们能够得到的保证是,复制操作肯定会当下一个被放入流中的操作之前执行。传递给此函数的主机内存指针必须是通过cudaHostAlloc()分配好的内存。(流中要求固定内存)
  7. 流同步:通过cudaStreamSynchronize()来协调。
  8. 流销毁:在退出应用程序之前,需要销毁对GPU操作进行排队的流,调用cudaStreamDestroy()
  9. 针对多个流:
    9.1. 记得对流进行同步操作。
    9.2. 将操作放入流的队列时,应采用宽度优先方式,而非深度优先的方式,换句话说,不是首先添加第0个流的所有操作,再依次添加后面的第1,2,…个流。而是交替进行添加,比如将a的复制操作添加到第0个流中,接着把a的复制操作添加到第1个流中,再继续其他的类似交替添加的行为。
    9.3. 要牢牢记住操作放入流中的队列中的顺序影响到CUDA驱动程序调度这些操作和流以及执行的方式。

TIPS:

  1. 当线程块的数量为GPU中处理数量的2倍时,将达到最优性能。
  2. 核函数执行的第一个计算就是计算输入数据的偏移。每个线程的起始偏移都是0到线程数量减1之间的某个值。然后,对偏移的增量为已启动线程的总数。

6.2.8 这是一个栗子

我们尝试用一个程序来比较cuda/c在GPU/CPU的运行效率,来不及了,快上车。
这是一个CUDA程序,请保存文件名为“文件名.cu”,在你的PC或者服务器上运行。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <stdio.h>
#include <time.h>
 
#define N (1024*1024)
#define M (10000)
#define THREADS_PER_BLOCK 1024
 
void serial_add(double *a, double *b, double *c, int n, int m)
{
    for(int index=0;index<n;index++)
    {
        for(int j=0;j<m;j++)
        {
            c[index] = a[index]*a[index] + b[index]*b[index];
        }
    }
}
 
__global__ void vector_add(double *a, double *b, double *c)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
        for(int j=0;j<M;j++)
        {
            c[index] = a[index]*a[index] + b[index]*b[index];
        }
}
 
int main()
{
    clock_t start,end;
 
    double *a, *b, *c;
    int size = N * sizeof( double );
 
    a = (double *)malloc( size );
    b = (double *)malloc( size );
    c = (double *)malloc( size );
 
    for( int i = 0; i < N; i++ )
    {
        a[i] = b[i] = i;
        c[i] = 0;
    }
 
    start = clock();
    serial_add(a, b, c, N, M);
 
    printf( "c[%d] = %f\n",0,c[0] );
    printf( "c[%d] = %f\n",N-1, c[N-1] );
 
    end = clock();
 
    float time1 = ((float)(end-start))/CLOCKS_PER_SEC;
    printf("CPU: %f seconds\n",time1);
 
    start = clock();
    double *d_a, *d_b, *d_c;
 
 
    cudaMalloc( (void **) &d_a, size );
    cudaMalloc( (void **) &d_b, size );
    cudaMalloc( (void **) &d_c, size );
 
 
    cudaMemcpy( d_a, a, size, cudaMemcpyHostToDevice );
    cudaMemcpy( d_b, b, size, cudaMemcpyHostToDevice );
 
    vector_add<<< (N + (THREADS_PER_BLOCK-1)) / THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( d_a, d_b, d_c );
 
    cudaMemcpy( c, d_c, size, cudaMemcpyDeviceToHost );
 
 
    printf( "c[%d] = %f\n",0,c[0] );
    printf( "c[%d] = %f\n",N-1, c[N-1] );
 
 
    free(a);
    free(b);
    free(c);
    cudaFree( d_a );
    cudaFree( d_b );
    cudaFree( d_c );
 
    end = clock();
    float time2 = ((float)(end-start))/CLOCKS_PER_SEC;
    printf("CUDA: %f seconds, Speedup: %f\n",time2, time1/time2);
 
    return 0;
}

效率对比
我们通过修改count的值并且加大循环次数来观察变量的效率的差别。

运行结果:

可见在数据量大的情况下效率还是相当不错的。

7. GPU or FPGA

GPU优势
1.从峰值性能来说,GPU(10Tflops)远远高于FPGA(<1TFlops);

2.GPU相对于FPGA还有一个优势就是内存接口, GPU的内存接口(传统的GDDR5,最近更是用上了HBM和HBM2)的带宽远好于FPGA的传统DDR接口(大约带宽高4-5倍);

3.功耗方面,虽然GPU的功耗远大于FPGA的功耗,但是如果要比较功耗应该比较在执行效率相同时需要的功耗。如果FPGA的架构优化能做到很好以致于一块FPGA的平均性能能够接近一块GPU,那么FPGA方案的总功耗远小于GPU,散热问题可以大大减轻。反之,如果需要二十块FPGA才能实现一块GPU的平均性能,那么FPGA在功耗方面并没有优势。

4.FPGA缺点有三点:
第一,基本单元的计算能力有限。为了实现可重构特性,FPGA 内部有大量极细粒度的基本单元,但是每个单元的计算能力(主要依靠LUT 查找表)都远远低于CPU 和GPU 中的ALU模块。
第二,速度和功耗相对专用定制芯片(ASIC)仍然存在不小差距。
第三,FPGA 价格较为昂贵,在规模放量的情况下单块FPGA 的成本要远高于专用定制芯片。最后谁能胜出, 完全取决于FPGA架构优化能否弥补峰值性能的劣势。

5.个人更推荐: CPU+FPGA的组合模式; 其中FPGA用于整形计算,cpu进行浮点计算和调度,此组合的拥有更高的单位功耗性能和更低的时延。最后更想GPU稳定开放,发挥其长处, 达到真正的物美价廉!

FPGA优势
人工智能目前仍处于早期阶段,未来人工智能的主战场是在推理环节,远没有爆发。未来胜负尚未可知,各家技术路线都有机会胜出。目前英伟达的GPU在训练场景中占据着绝对领导地位,但是在未来,专注于推理环节的FPGA必将会发挥巨大的价值。

FPGA和GPU内都有大量的计算单元,因此它们的计算能力都很强。在进行神经网络运算的时候,两者的速度会比CPU快很多。但是GPU由于架构固定,硬件原生支持的指令也就固定了,而FPGA则是可编程的。其可编程性是关键,因为它让软件与终端应用公司能够提供与其竞争对手不同的解决方案,并且能够灵活地针对自己所用的算法修改电路。

在平均性能方面,GPU逊于FPGA,FPGA可以根据特定的应用去编程硬件,例如如果应用里面的加法运算非常多就可以把大量的逻辑资源去实现加法器,而GPU一旦设计完就不能改动了,所以不能根据应用去调整硬件资源。
目前机器学习大多使用SIMD架构,即只需一条指令可以平行处理大量数据,因此用GPU很适合。但是有些应用是MISD,即单一数据需要用许多条指令平行处理,这种情况下用FPGA做一个MISD的架构就会比GPU有优势。 所以,对于平均性能,看的就是FPGA加速器架构上的优势是否能弥补运行速度上的劣势。如果FPGA上的架构优化可以带来相比GPU架构两到三个数量级的优势,那么FPGA在平均性能上会好于GPU。

在功耗能效比方面,同样由于FPGA的灵活性,在架构优化到很好时,一块FPGA的平均性能能够接近一块GPU,那么FPGA方案的总功耗远小于GPU,散热问题可以大大减轻。 能效比的比较也是类似,能效指的是完成程序执行消耗的能量,而能量消耗等于功耗乘以程序的执行时间。虽然GPU的功耗远大于FPGA的功耗,但是如果FPGA执行相同程序需要的时间比GPU长几十倍,那FPGA在能效比上就没有优势了;反之如果FPGA上实现的硬件架构优化得很适合特定的机器学习应用,执行算法所需的时间仅仅是GPU的几倍或甚至于接近GPU,那么FPGA的能效比就会比GPU强。

在峰值性能比方面,虽然GPU的峰值性能(10Tflops)远大于FPGA的峰值性能(<1Tflops),但针对特定的场景来讲吞吐量并不比GPU差。

8. 深度学习的三种硬件方案:ASIC,FPGA,GPU

8.1 对深度学习硬件平台的要求

要想明白“深度学习”需要怎样的硬件,必须了解深度学习的工作原理。首先在表层上,我们有一个巨大的数据集,并选定了一种深度学习模型。每个模型都有一些内部参数需要调整,以便学习数据。而这种参数调整实际上可以归结为优化问题,在调整这些参数时,就相当于在优化特定的约束条件

  • 矩阵相乘(Matrix Multiplication)——几乎所有的深度学习模型都包含这一运算,它的计算十分密集。
  • 卷积(Convolution)——这是另一个常用的运算,占用了模型中大部分的每秒浮点运算(浮点/秒)。
  • 循环层(Recurrent Layers )——模型中的反馈层,并且基本上是前两个运算的组合。
  • All Reduce——这是一个在优化前对学习到的参数进行传递或解析的运算序列。在跨硬件分布的深度学习网络上执行同步优化时(如AlphaGo的例子),这一操作尤其有效。

除此之外,深度学习的硬件加速器需要具备数据级别和流程化的并行性、多线程和高内存带宽等特性。 另外,由于数据的训练时间很长,所以硬件架构必须低功耗。 因此,效能功耗比(Performance per Watt)是硬件架构的评估标准之一。

CNN在应用中,一般采用GPU加速,请解释为什么GPU可以有加速效果,主要加速算法的哪一个部分?

这里默认gpu加速是指NVIDIA的CUDA加速。CPU是中央处理单元,gpu是图形处理单元,gpu由上千个流处理器(core)作为运算器。执行采用单指令多线程(SIMT)模式。相比于单核CPU(向量机)流水线式的串行操作,虽然gpu单个core计算能力很弱,但是通过大量线程进行同时计算,在数据量很大是会活动较为可观的加速效果。

具体到cnn,利用gpu加速主要是在conv(卷积)过程上。conv过程同理可以像以上的向量加法一样通过cuda实现并行化。具体的方法很多,不过最好的还是利用fft(快速傅里叶变换)进行快速卷积。NVIDIA提供了cufft库实现fft,复数乘法则可以使用cublas库里的对应的level3的cublasCgemm函数。

GPU加速的基本准则就是“人多力量大”。CNN说到底主要问题就是计算量大,但是却可以比较有效的拆分成并行问题。随便拿一个层的filter来举例子,假设某一层有n个filter,每一个需要对上一层输入过来的map进行卷积操作。那么,这个卷积操作并不需要按照线性的流程去做,每个滤波器互相之间并不影响,可以大家同时做,然后大家生成了n张新的谱之后再继续接下来的操作。既然可以并行,那么同一时间处理单元越多,理论上速度优势就会越大。所以,处理问题就变得很简单粗暴,就像NV那样,暴力增加显卡单元数(当然,显卡的架构、内部数据的传输速率、算法的优化等等也都很重要)。

GPU主要是针对图形显示及渲染等技术的出众,而其中的根本是因为处理矩阵算法能力的强大,刚好CNN中涉及大量的卷积,也就是矩阵乘法等,所以在这方面具有优势。

机器学习的算法一定得经过gpu加速吗?

不一定。只有需要大量浮点数计算,例如矩阵乘法,才需要GPU加速。 用CNN对图像进行分类就是一个需要大量浮点数计算的典型案例,通常需要GPU加速

对于ASICFPGA分布式计算,这里不再展开讲,有兴趣的小伙伴可以,自行学习。不过….说不定某天博主心情好,就会梳理一下这几种硬件方案在端到端上应用的区别了。

菜鸟入门教程就到这里了,聪明的你一定不满足这个入门教程,如有兴趣进一步学习CUDA编程,可移步NVIDIA官方的课程平台CUDA ZONE(PS:中文网站,英文课程)

作者:He_Yu
链接:https://www.jianshu.com/p/34a504af8d51
来源:简书
著作权归作者所有。商业转载请联系作者获得授权,非商业转载请注明出处。

cuda资料整理(1)

from: https://bbs.csdn.net/topics/390798229?list=lz

CUDA是什么
        CUDA,Compute Unified Device Architecture的简称,是由NVIDIA公司创立的基于他们公司生产的图形处理器GPUs(Graphics Processing Units,可以通俗的理解为显卡)的一个并行计算平台和编程模型。
        通过CUDA,GPUs可以很方便地被用来进行通用计算(有点像在CPU中进行的数值计算等等)。在没有CUDA之前,GPUs一般只用来进行图形渲染(如通过OpenGL,DirectX)。
        开发人员可以通过调用CUDA的API,来进行并行编程,达到高性能计算目的。NVIDIA公司为了吸引更多的开发人员,对CUDA进行了编程语言扩展,如CUDA C/C++,CUDA Fortran语言。注意CUDA C/C++可以看作一个新的编程语言,因为NVIDIA配置了相应的编译器nvcc,CUDA Fortran一样。更多信息可以参考文献。

        如果粗暴的认为C语言工作的对象是CPU和内存条(接下来,称为主机内存),那么CUDA C工作的的对象就是GPU及GPU上的内存(接下来,称为设备内存),且充分利用了GPU多核的优势及降低了并行编程的难度。一般通过C语言把数据从外界读入,再分配数据,给CUDA C,以便在GPU上计算,然后再把计算结果返回给C语言,以便进一步工作,如进一步处理及显示,或重复此过程。
 主要概念与名称
主机
        将CPU及系统的内存(内存条)称为主机。
设备
        将GPU及GPU本身的显示内存称为设备。
线程(Thread)
        一般通过GPU的一个核进行处理。(可以表示成一维,二维,三维,具体下面再细说)。
线程块(Block)
        1. 由多个线程组成(可以表示成一维,二维,三维,具体下面再细说)。
        2. 各block是并行执行的,block间无法通信,也没有执行顺序。
        3. 注意线程块的数量限制为不超过65535(硬件限制)。
线程格(Grid)
        由多个线程块组成(可以表示成一维,二维,三维,具体下面再细说)。

线程束
        在CUDA架构中,线程束是指一个包含32个线程的集合,这个线程集合被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行,线程束中的每个线程都将在不同数据上执行相同的命令。
核函数(Kernel)
        1. 在GPU上执行的函数通常称为核函数。
        2. 一般通过标识符__global__修饰,调用通过<<<参数1,参数2>>>,用于说明内核函数中的线程数量,以及线程是如何组织的。
        3. 以线程格(Grid)的形式组织,每个线程格由若干个线程块(block)组成,而每个线程块又由若干个线程(thread)组成。
        4. 是以block为单位执行的。
        5. 叧能在主机端代码中调用。
        6. 调用时必须声明内核函数的执行参数。
        7. 在编程时,必须先为kernel函数中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在GPU计算时会发生错误,例如越界或报错,甚至导致蓝屏和死机。
C/C++ code?

dim3结构类型
        1. dim3是基亍uint3定义的矢量类型,相当亍由3个unsigned int型组成的结构体。uint3类型有三个数据成员unsigned int x; unsigned int y; unsigned int z;
        2. 可使用亍一维、二维或三维的索引来标识线程,构成一维、二维或三维线程块。
        3. dim3结构类型变量用在核函数调用的<<<,>>>中。
        4. 相关的几个内置变量
        4.1. threadIdx,顾名思义获取线程thread的ID索引;如果线程是一维的那么就取threadIdx.x,二维的还可以多取到一个值threadIdx.y,以此类推到三维threadIdx.z。
        4.2. blockIdx,线程块的ID索引;同样有blockIdx.x,blockIdx.y,blockIdx.z。
        4.3. blockDim,线程块的维度,同样有blockDim.x,blockDim.y,blockDim.z。
        4.4. gridDim,线程格的维度,同样有gridDim.x,gridDim.y,gridDim.z。
        5. 对于一维的block,线程的threadID=threadIdx.x。
        6. 对于大小为(blockDim.x, blockDim.y)的 二维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。
        7. 对于大小为(blockDim.x, blockDim.y, blockDim.z)的 三维 block,线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。
        8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride。
函数修饰符
        1. __global__,表明被修饰的函数在设备上执行,但在主机上调用。
        2. __device__,表明被修饰的函数在设备上执行,但只能在其他__device__函数或者__global__函数中调用。
常用的GPU内存函数
cudaMalloc()
        1. 函数原型: cudaError_t cudaMalloc (void **devPtr, size_t size)。
        2. 函数用处:与C语言中的malloc函数一样,只是此函数在GPU的内存你分配内存。
        3. 注意事项:
        3.1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数;
        3.2. 可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作;
        3.3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数;
        3.4. 不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存读写操作(即不能进行解引用)。
cudaMemcpy()
        1. 函数原型:cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。
        2. 函数作用:与c语言中的memcpy函数一样,只是此函数可以在主机内存和GPU内存之间互相拷贝数据。
        3. 函数参数:cudaMemcpyKind kind表示数据拷贝方向,如果kind赋值为cudaMemcpyDeviceToHost表示数据从设备内存拷贝到主机内存。
        4. 与C中的memcpy()一样,以同步方式执行,即当函数返回时,复制操作就已经完成了,并且在输出缓冲区中包含了复制进去的内容。
        5. 相应的有个异步方式执行的函数cudaMemcpyAsync(),这个函数详解请看下面的流一节有关内容。
cudaFree()
        1. 函数原型:cudaError_t cudaFree ( void* devPtr )。
        2. 函数作用:与c语言中的free()函数一样,只是此函数释放的是cudaMalloc()分配的内存。
        下面实例用于解释上面三个函数
C/C++ code?

GPU内存分类
全局内存
        通俗意义上的设备内存。
共享内存
        1. 位置:设备内存。
        2. 形式:关键字__shared__添加到变量声明中。如__shared__ float cache[10]。
        3. 目的:对于GPU上启动的每个线程块,CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存,但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。
常量内存
        1. 位置:设备内存
        2. 形式:关键字__constant__添加到变量声明中。如__constant__ float s[10];。
        3. 目的:为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下,用常量内存替换全局内存能有效地减少内存带宽。
        4. 特点:常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时,静态地分配空间。
        5. 要求:当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol(),而cudaMemcpy()会复制到全局内存。
        6. 性能提升的原因:
        6.1. 对常量内存的单次读操作可以广播到其他的“邻近”线程。这将节约15次读取操作。(为什么是15,因为“邻近”指半个线程束,一个线程束包含32个线程的集合。)
        6.2. 常量内存的数据将缓存起来,因此对相同地址的连续读操作将不会产生额外的内存通信量。
纹理内存
        1. 位置:设备内存
        2. 目的:能够减少对内存的请求并提供高效的内存带宽。是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序设计,意味着一个线程读取的位置可能与邻近线程读取的位置“非常接近”。如下图:

        3. 纹理变量(引用)必须声明为文件作用域内的全局变量。
        4. 形式:分为一维纹理内存 和 二维纹理内存。
        4.1. 一维纹理内存
        4.1.1. 用texture<类型>类型声明,如texture<float> texIn。
        4.1.2. 通过cudaBindTexture()绑定到纹理内存中。
        4.1.3. 通过tex1Dfetch()来读取纹理内存中的数据。
        4.1.4. 通过cudaUnbindTexture()取消绑定纹理内存。
        4.2. 二维纹理内存
        4.2.1. 用texture<类型,数字>类型声明,如texture<float,2> texIn。
        4.2.2. 通过cudaBindTexture2D()绑定到纹理内存中。
        4.2.3. 通过tex2D()来读取纹理内存中的数据。
        4.2.4. 通过cudaUnbindTexture()取消绑定纹理内存。
固定内存
        1. 位置:主机内存。
        2. 概念:也称为页锁定内存或者不可分页内存,操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存中。因此操作系统能够安全地使某个应用程序访问该内存的物理地址,因为这块内存将不会破坏或者重新定位。
        3. 目的:提高访问速度。由于GPU知道主机内存的物理地址,因此可以通过“直接内存访问DMA(Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因此DMA复制过程中使用固定内存是非常重要的。
        4. 缺点:使用固定内存,将失去虚拟内存的所有功能;系统将更快的耗尽内存。
        5. 建议:对cudaMemcpy()函数调用中的源内存或者目标内存,才使用固定内存,并且在不再需要使用它们时立即释放。
        6. 形式:通过cudaHostAlloc()函数来分配;通过cudaFreeHost()释放。
        7. 只能以异步方式对固定内存进行复制操作。
原子性
        1. 概念:如果操作的执行过程不能分解为更小的部分,我们将满足这种条件限制的操作称为原子操作。
        2. 形式:函数调用,如atomicAdd(addr,y)将生成一个原子的操作序列,这个操作序列包括读取地址addr处的值,将y增加到这个值,以及将结果保存回地址addr。
常用线程操作函数
        1. 同步方法__syncthreads(),这个函数的调用,将确保线程块中的每个线程都执行完__syscthreads()前面的语句后,才会执行下一条语句。
使用事件来测量性能
        1. 用途:为了测量GPU在某个任务上花费的时间。CUDA中的事件本质上是一个GPU时间戳。由于事件是直接在GPU上实现的。因此不适用于对同时包含设备代码和主机代码的混合代码设计。
        2. 形式:首先创建一个事件,然后记录事件,再计算两个事件之差,最后销毁事件。如:
C/C++ code?


        1. 扯一扯:并发重点在于一个极短时间段内运行多个不同的任务;并行重点在于同时运行一个任务。
        2. 任务并行性:是指并行执行两个或多个不同的任务,而不是在大量数据上执行同一个任务。
        3. 概念:CUDA流表示一个GPU操作队列,并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作,如核函数启动,内存复制以及事件的启动和结束等。这些操作的添加到流的顺序也是它们的执行顺序。可以将每个流视为GPU上的一个任务,并且这些任务可以并行执行。
        4. 硬件前提:必须是支持设备重叠功能的GPU。支持设备重叠功能,即在执行一个核函数的同时,还能在设备与主机之间执行复制操作。
        5. 声明与创建:声明cudaStream_t stream;,创建cudaSteamCreate(&stream);。
        6. cudaMemcpyAsync():前面在cudaMemcpy()中提到过,这是一个以异步方式执行的函数。在调用cudaMemcpyAsync()时,只是放置一个请求,表示在流中执行一次内存复制操作,这个流是通过参数stream来指定的。当函数返回时,我们无法确保复制操作是否已经启动,更无法保证它是否已经结束。我们能够得到的保证是,复制操作肯定会当下一个被放入流中的操作之前执行。传递给此函数的主机内存指针必须是通过cudaHostAlloc()分配好的内存。(流中要求固定内存)
        7. 流同步:通过cudaStreamSynchronize()来协调。
        8. 流销毁:在退出应用程序之前,需要销毁对GPU操作进行排队的流,调用cudaStreamDestroy()。
        9. 针对多个流:
        9.1. 记得对流进行同步操作。
        9.2. 将操作放入流的队列时,应采用宽度优先方式,而非深度优先的方式,换句话说,不是首先添加第0个流的所有操作,再依次添加后面的第1,2,…个流。而是交替进行添加,比如将a的复制操作添加到第0个流中,接着把a的复制操作添加到第1个流中,再继续其他的类似交替添加的行为。
        9.3. 要牢牢记住操作放入流中的队列中的顺序影响到CUDA驱动程序调度这些操作和流以及执行的方式。
技巧
        1. 当线程块的数量为GPU中处理数量的2倍时,将达到最优性能。
        2. 核函数执行的第一个计算就是计算输入数据的偏移。每个线程的起始偏移都是0到线程数量减1之间的某个值。然后,对偏移的增量为已启动线程的总数。
实例程序
感兴趣的读者可以下载本书附带的示例代码点击此处下载https://developer.nvidia.com/sites/default/files/akamai/cuda/files/cuda_by_example.zip

中芯A股科创板今天上市

今天科创板上市暴涨到300%,然后回落到200%,还是很让人吃惊,下周一20号寒武纪上市应该会有更火爆的炒作,拭目以待。

关键是港股的表现,今天00981暴跌25%,利润大幅缩水非常可惜,追求最大收益是普遍心态,所以要学习控制,特别是关键时间点的心理,离远一点可以看的更清楚。

昨天讲美股全部抛出,算是一个节点,休息一下集中精力放到新的FR2项目。

HK00981(2)

下周二申购,昨天周五H股价格已经到33港币,记录7.4日数据,看周日的发行定价有多少,港股40还是能看到的,再高就太疯了。

(6月29日)未平仓空仓比减少最大的6只股票其中中芯国际和中兴通讯都减少空仓,所以上周都有大幅上涨。

股票名称上一次空仓比最新空仓比减少值↓
联邦制药(03933)4.03%1.89%-2.14%
佐丹奴国际(00709)2.51%0.80%-1.72%
中芯国际(00981)6.75%5.82%-0.93%
VTECH HOLDINGS(00303)2.00%1.13%-0.87%
山东黄金(01787)0.82%0.24%-0.58%
中兴通讯(00763)7.73%7.23%-0.50%

HK00981

00981太火,吃瓜群众一定要积极围观并参与😄啊,端午在家查查资料,做个研究。

先是看了下面小短文,简单的分析了股本结构和后市,有点低估流通股份比例,而且我不认为它在科创板上市后会出现一地鸡毛的中石油悲剧,我们是有进步的。

https://3g.163.com/news/article/FFTN54240539JC86.html?from=history-back-list

我觉得作者把大股东作为长期投资者,不会卖出股票,这个前提不一定成立,大股东持股有50%,清华紫光和赵国伟应该已经减持了,所以流动股票应该高于20%。

其实有趣的是观察港股的空单持仓比例,从下面链接可以看到6月19日的数据,每周发布。

https://www.zhitongcaijing.com/content/search.html?keywords=%E6%B8%AF%E8%82%A1%E7%A9%BA%E4%BB%93%E6%8C%81%E5%8D%95%E7%BB%9F%E8%AE%A1

看到在6月19日有6.88%的空单,就是约3.89亿股票卖空,趋势是在增加的。

股票名称6月12日空仓比最新6月19日空仓比增加值↓
中芯国际(00981)5.42%6.88%1.45%

这段时间继续跟踪00981,希望能有好收获,另外港股打新没意思,还是大A股爽呀🤭。