cuda开始使用cuda


备注

CUDA是用于其GPU的专有NVIDIA并行计算技术和编程语言。

GPU是高度并行的机器,能够并行运行数千个轻量级线程。每个GPU线程的执行速度通常较慢,而且其上下文较小。另一方面,GPU能够并行运行数千个线程,甚至更多并发运行(精确数量取决于实际的GPU模型)。 CUDA是专为NVIDIA GPU架构设计的C ++方言。但是,由于体系结构的不同,大多数算法都不能简单地从普通的C ++中复制粘贴 - 它们会运行,但速度很慢。

术语

  • host - 指在该环境中运行的普通基于CPU的硬件和普通程序
  • device - 指CUDA程序运行的特定GPU。单个主机可以支持多个设备。
  • kernel - 驻留在可以从主机代码调用的设备上的函数。

物理处理器结构

支持CUDA的GPU处理器具有以下物理结构:

  • 芯片 - GPU的整个处理器。有些GPU有两个。
  • 流式多处理器 (SM) - 每个芯片最多包含~100个SM,具体取决于型号。每个SM几乎独立于另一个SM运行,仅使用全局存储器相互通信。
  • CUDA核心 - SM的单个标量计算单元。它们的确切数量取决于架构。每个核心可以快速连续地处理几个并发执行的线程(类似于CPU中的超线程)。

此外,每个SM都具有一个或多个warp调度程序 。每个调度程序将一条指令分派给多个CUDA核心。这有效地使SM在32宽SIMD模式下操作。

CUDA执行模型

GPU的物理结构直接影响内核在设备上的执行方式,以及如何在CUDA中对它们进行编程。使用调用配置调用内核,该调用配置指定生成多少并行线程。

  • grid - 表示在内核调用时生成的所有线程。它被指定为一个或两个维数的
  • - 是一组半独立的线程 。每个块都分配给一个SM。因此,块只能通过全局存储器进行通信。块不以任何方式同步。如果块太多,则一些块可能在其他块之后顺序执行。另一方面,如果资源允许,可以在同一个SM上运行多个块,但程序员无法从中发生这种情况(除了明显的性能提升)。
  • 线程 - 由单个CUDA核心执行的标量序列指令。线程“轻量级”,具有最小的上下文,允许硬件快速交换它们。由于它们的数量,CUDA线程运行时分配了几个寄存器,并且堆栈非常短(最好是没有!)。因此,CUDA编译器倾向于内联所有函数调用以展平内核,使其仅包含静态跳转和循环。函数ponter调用和虚拟方法调用虽然在大多数较新的设备中受支持,但通常会产生重大的性能问题。

每个线程由块索引blockIdx和块threadIdx内的线程索引threadIdx 。任何正在运行的线程都可以随时检查这些数字,这是区分一个线程与另一个线程的唯一方法。

此外,线程被组织成warp ,每个warp包含32个线程。单个warp中的线程在SIMD fahsion中执行完美同步。来自不同warp但在同一块中的线程可以按任何顺序执行,但可以由程序员强制同步。来自不同块的线程无法以任何方式同步或直接交互。

记忆组织

在正常的CPU编程中,内存组织通常对程序员是隐藏的。典型的程序就好像只有RAM一样。所有内存操作(例如管理寄存器,使用L1-L2-L3缓存,交换到磁盘等)都由编译器,操作系统或硬件本身处理。

CUDA的情况并非如此。虽然较新的GPU模型部分地隐藏了负担,例如通过CUDA 6中的统一内存 ,但出于性能原因,仍然值得了解组织。基本的CUDA内存结构如下:

  • 主机内存 - 常规RAM。主要由主机代码使用,但较新的GPU型号也可以访问它。当内核访问主机内存时,GPU必须通常通过PCIe连接器与主板通信,因此它相对较慢。
  • 设备内存/全局内存 - GPU的主要片外内存,可供所有线程使用。
  • 共享内存 - 位于每个SM中,允许比全局更快的访问。共享内存对每个块都是私有的。单个块中的线程可以使用它进行通信。
  • 寄存器 - 每个线程的最快,私有,无法寻址的内存。通常,这些不能用于通信,但是一些内在函数允许在经线内混洗它们的内容。
  • 本地内存 -这寻址每个线程的私有内存。这用于寄存器溢出和具有可变索引的本地数组。在物理上,它们存在于全球记忆中。
  • 纹理内存,常量内存 - 全局内存的一部分,标记为内核不可变。这允许GPU使用专用缓存。
  • L2缓存 - 片上,可供所有线程使用。给定线程数量,每个缓存行的预期生命周期远低于CPU。它主要用于辅助未对齐和部分随机的内存访问模式。
  • L1缓存 - 与共享内存位于同一空间。同样,在给定使用它的线程数量的情况下,数量相当小,因此不要指望数据会长时间停留在那里。可以禁用L1缓存。

版本

计算能力建筑 GPU代号发布日期
1.0 特斯拉 G80 2006-11-08
1.1 特斯拉 G84,G86,G92,G94,G96,G98, 2007-04-17
1.2 特斯拉 GT218,GT216,GT215 2009-04-01
1.3 特斯拉 GT200,GT200b 2009-04-09
2.0 费米 GF100,GF110 2010-03-26
2.1 费米 GF104,GF106 GF108,GF114,GF116,GF117,GF119 2010-07-12
3.0 开普勒 GK104,GK106,GK107 2012-03-22
3.2 开普勒 GK20A 2014年4月1日
3.5 开普勒 GK110,GK208 2013年2月19日
3.7 开普勒 GK210 情节中字
5 麦克斯韦 GM107,GM108 2014年2月18日
5.2 麦克斯韦 GM200,GM204,GM206 2014-09-18
5.3 麦克斯韦 GM20B 2015-04-01
6 帕斯卡尔 GP100 2016年10月1日
6.1 帕斯卡尔 GP102,GP104,GP106 2016年5月27日

发布日期标志着支持给定计算能力的第一个GPU的发布。有些日期是近似值,例如,2014年第二季度发布了3.2张卡。

编译和运行示例程序

NVIDIA安装指南以运行示例程序结束,以验证CUDA Toolkit的安装,但没有明确说明如何安装。首先检查所有先决条件。检查示例程序的默认CUDA目录。如果不存在,可以从官方CUDA网站下载。导航到示例所在的目录。

$ cd /path/to/samples/
$ ls
 

您应该看到类似于的输出:

0_Simple     2_Graphics  4_Finance      6_Advanced       bin     EULA.txt
1_Utilities  3_Imaging   5_Simulations  7_CUDALibraries  common  Makefile
 

确保Makefile 存在于此目录中。基于UNIX的系统中的make 命令将构建所有示例程序。或者,导航到存在另一个Makefile 的子目录,并从那里运行make 命令以仅构建该样本。

运行两个建议的示例程序 - deviceQuerybandwidthTest

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

输出类似于下面显示的输出:

./deviceQuery Starting...

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

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 950M"
  CUDA Driver Version / Runtime Version          7.5 / 7.5
  CUDA Capability Major/Minor version number:    5.0
  Total amount of global memory:                 4096 MBytes (4294836224 bytes)
  ( 5) Multiprocessors, (128) CUDA Cores/MP:     640 CUDA Cores
  GPU Max Clock rate:                            1124 MHz (1.12 GHz)
  Memory Clock rate:                             900 Mhz
  Memory Bus Width:                              128-bit
  L2 Cache Size:                                 2097152 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 PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 7.5, CUDA Runtime Version = 7.5, NumDevs = 1, Device0 = GeForce GTX 950M
Result = PASS
 

结尾处的语句Result = PASS 表示一切正常。现在,以类似的方式运行其他建议的示例程序bandwidthTest 。输出类似于:

[CUDA Bandwidth Test] - Starting...
Running on...

 Device 0: GeForce GTX 950M
 Quick Mode

 Host to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10604.5

 Device to Host Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10202.0

 Device to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            23389.7

Result = PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
 

同样, Result = PASS 语句表明一切都已正确执行。所有其他示例程序都可以以类似的方式运行。

让我们推出一个单独的CUDA线程来打招呼

这个简单的CUDA程序演示了如何编写将在GPU(又称“设备”)上执行的函数。 CPU或“主机”通过调用称为“内核”的特殊函数来创建CUDA线程。 CUDA程序是具有附加语法的C ++程序。

要查看它是如何工作的,请将以下代码放在名为hello.cu 的文件中:

#include <stdio.h>

// __global__ functions, or "kernels", execute on the device
__global__ void hello_kernel(void)
{
  printf("Hello, world from the device!\n");
}

int main(void)
{
  // greet from the host
  printf("Hello, world from the host!\n");

  // launch a kernel with a single thread to greet from the device
  hello_kernel<<<1,1>>>();

  // wait for the device to finish so that we see the message
  cudaDeviceSynchronize();

  return 0;
}
 

(请注意,为了在设备上使用printf 功能,您需要一台计算能力至少为2.0的设备。有关详细信息,请参阅版本概述 。)

现在让我们使用NVIDIA编译器编译程序并运行它:

$ nvcc hello.cu -o hello
$ ./hello
Hello, world from the host!
Hello, world from the device!
 

有关上述示例的一些其他信息:

  • nvcc 代表“NVIDIA CUDA编译器”。它将源代码分为主机和设备组件。
  • __global__ 是函数声明中使用的CUDA关键字,表示该函数在GPU设备上运行并从主机调用。
  • 三角括号( <<<>>> )标记从主机代码到设备代码(也称为“内核启动”)的调用。这三个括号中的数字表示并行执行的次数和线程数。

先决条件

要开始使用CUDA进行编程,请下载并安装CUDA Toolkit和开发人员驱动程序 。该工具包包括nvcc ,NVIDIA CUDA编译器以及开发CUDA应用程序所需的其他软件。该驱动程序可确保GPU程序在支持CUDA的硬件上正确运行,您也需要这些硬件

您可以通过从命令行运行nvcc --version 来确认您的计算机上已正确安装CUDA Toolkit。例如,在Linux机器上,

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jul_12_18:28:38_CDT_2016
Cuda compilation tools, release 8.0, V8.0.32
 

输出编译器信息。如果上一个命令不成功,则可能未安装CUDA Toolkit,或者nvcc (Windows机器上的C:\CUDA\bin /usr/local/cuda/bin ,POSIX操作系统上的/usr/local/cuda/bin )的路径不属于您的PATH 环境变量。

此外,您还需要一个与nvcc 一起使用的主编译器来编译和构建CUDA程序。在Windows上,这是与Microsoft Visual Studio一起提供的Microsoft编译器cl.exe 。在POSIX OS上,可以使用其他编译器,包括gccg++ 。官方CUDA 快速入门指南可以告诉您特定平台支持哪些编译器版本。

为确保一切设置正确,让我们编译并运行一个简单的CUDA程序,以确保所有工具正确地协同工作。

__global__ void foo() {}

int main()
{
  foo<<<1,1>>>();

  cudaDeviceSynchronize();
  printf("CUDA error: %s\n", cudaGetErrorString(cudaGetLastError()));

  return 0;
}
 

要编译此程序,请将其复制到名为test.cu的文件中,然后从命令行进行编译。例如,在Linux系统上,以下应该可以工作:

$ nvcc test.cu -o test
$ ./test
CUDA error: no error
 

如果程序成功没有错误,那么让我们开始编码!

使用CUDA对两个数组求和

此示例说明如何创建一个简单的程序,该程序将使用CUDA对两个int 数组求和。

CUDA程序是异构的,由在CPU和GPU上运行的部分组成。

利用CUDA的程序的主要部分类似于CPU程序,包括

  • 将在GPU上使用的数据的内存分配
  • 数据从主机内存复制到GPU内存
  • 调用内核函数来处理数据
  • 将结果复制到CPU内存

要分配设备内存,我们使用cudaMalloc 函数。要在设备和主机之间复制数据,可以使用cudaMemcpy 函数。 cudaMemcpy 的最后一个参数指定了复制操作的方向。有5种可能的类型:

  • cudaMemcpyHostToHost - 主持人 - >主持人
  • cudaMemcpyHostToDevice - 主机 - >设备
  • cudaMemcpyDeviceToHost - 设备 - >主机
  • cudaMemcpyDeviceToDevice - 设备 - >设备
  • cudaMemcpyDefault - 基于默认的统一虚拟地址空间

接下来调用内核函数。三个V形之间的信息是执行配置,它指示有多少设备线程并行执行内核。第一个数字(示例中为2 )指定块数,第二个( (size + 1) / 2 ) - 块中的线程数。请注意,在此示例中,我们将大小添加1,以便我们请求一个额外的线程,而不是让一个线程负责两个元素。

由于内核调用是异步函数, cudaDeviceSynchronize 调用cudaDeviceSynchronize 以等待执行完成。将结果数组复制到主机内存,并使用cudaFree 释放设备上分配的所有内存。

要将函数定义为内核,请使用__global__ 声明说明符。每个线程都将调用此函数。如果我们希望每个线程处理结果数组的元素,那么我们需要一种区分和识别每个线程的方法。 CUDA定义变量blockDimblockIdxthreadIdx 。预定义变量blockDim 包含内核启动的第二个执行配置参数中指定的每个线程块的维度。预定义变量threadIdxblockIdx 包含其线程块内的线程索引和网格内的线程块。请注意,由于我们可能比数组中的元素请求多一个线程,因此我们需要传入size 以确保我们不会访问数组的末尾。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void addKernel(int* c, const int* a, const int* b, int size) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < size) {
        c[i] = a[i] + b[i];
    }
}

// Helper function for using CUDA to add vectors in parallel.
void addWithCuda(int* c, const int* a, const int* b, int size) {
    int* dev_a = nullptr;
    int* dev_b = nullptr;
    int* dev_c = nullptr;

    // Allocate GPU buffers for three vectors (two input, one output)
    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));

    // Copy input vectors from host memory to GPU buffers.
    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
    // 2 is number of computational blocks and (size + 1) / 2 is a number of threads in a block
    addKernel<<<2, (size + 1) / 2>>>(dev_c, dev_a, dev_b, size);
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaDeviceSynchronize();

    // Copy output vector from GPU buffer to host memory.
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
}

int main(int argc, char** argv) {
    const int arraySize = 5;
    const int a[arraySize] = {  1,  2,  3,  4,  5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    addWithCuda(c, a, b, arraySize);

    printf("{1, 2, 3, 4, 5} + {10, 20, 30, 40, 50} = {%d, %d, %d, %d, %d}\n", c[0], c[1], c[2], c[3], c[4]);

    cudaDeviceReset();

    return 0;
}