CUDA学习笔记(1)

一、环境配置与安装

• Linux下安装请参考Ubuntu12.04配置NVIDIA cuda5.5经验帖
• Windows下安装请参考CUDA在Windows的安装和使用

  我是在Windows下安装的CUDA，使用Visual Studio 2013编程。安装过程是只用点确定的傻瓜式安装，并且会自动给你添加环境变量。安装完成后，Visual Studio在新建项目的时候会多出来一个CUDA的工程选型，如下图所示：

  建立工程的时候，编译器会自动生成一个模板供我们参考。

二、GPU的硬件结构

  为了更好地编写CUDA程序，我们必须先了解我们的显卡硬件信息，不同的显卡（GPU）会有不同的硬件特性，如果编程不恰当，会直接导致CUDA程序无法在不同型号的显卡上通用！

  先集中介绍下CUDA中的常用专业名词：

英 简写 中 Stream Processor SP 流处理器 Stream MultiProcessor SM 流处理器组 Thread 线程 Block 线程块 Grid 线程网格 Warp 线程束 shared memory 共享内存 kernel function 核函数

  显卡结构方面推荐大家去看Rachel-Zhang的CUDA系列学习（一），里面做了非常详细的介绍，这里我借用CUDA_C_Programming_Guide中的两张图做简要说明。
  首先简单说下显卡的硬件结构，一块显卡的PCB上主要有 核心芯片 和 核心外的显存，前者相当于计算机中的CPU，后者相当于内存，详细参数可以通过GPU-Z或者在techpowerup上查看，下图就是显卡的结构，决定显卡计算能力的就是红框中的GPU核心芯片。

  可以看到上图中的GPU核心芯片的型号是GM107，其内部的架构如下图所示。它有3个 SM（Stream MultiProcessor），每个SM中又有128个 SP （Stream Processor）以及64KB的 共享内存 （shared memory）和 一级缓存 （L1 cache）。
  这里我们可以把每个 SM 看做一块 CPU，一个 SM 的多个 SP 可以看做 CPU 中的多个核，显然 GPU 的 核远多于 CPU 的核，这也是用它并行计算远强于 CPU 的根本原因。

  除了GPU核心，显存也十分重要，它的大小远大于GPU核心中的缓存，相当于内存对于CPU的作用。当然同一个GPU核心可以搭配不同大小的显存，例如同样是GP102核心的GTX1080和TITAN-X就有着不同的显存。

三、CUDA的软件结构

  线程（Thread）是GPU中最小的执行单元，在CPU上，同一时间一个核心上只能运行一条线程（Thread），同样的，GPU的每个SP上在同一时间也只能运行一条线程。为了方便调度GPU上成千上万条线程，CUDA为我们提供了线程块（Block）、线程网格（Grid）来帮助我们控制线程的调度。每个线程网格（Grid）执行同样的函数——核函数。
  这里先说明一个重要概念——核函数，它是同一时刻被所有使用到的SP同时执行的函数，例如：现在有两个长度都为128的 int 型数组 A[128] 和 B[128] 对应位置元素相加为 C[128]。

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

  那么只用在核函数中写以上代码，就会有128个线程（Thread）占用128个流处理器（SP）同时将对应位置的元素相加。
  但要注意，当线程（Thread）的数量大于流处理器（SP）的数量时，一个流处理器（SP）可以像CPU的核心一样在不同时间执行多个线程。

每个线程块（Block）中存放着若干线程（Thread）。
1. 一个线程块（Block）只能放在一个SM中运行，一个SM中同一时刻可以执行多个线程块（Block）。
2. 一条线程（Thread）由一个SP进行处理，并且都有自己对应的序号threadIdx。
3. 一个线程块（Block）中的线程数是有限的，一般最大为512（Fermi架构）或1024（Maxwell架构）。
4. 一个线程块（Block）中的所有线程共享同一个SM中的共享内存 （shared memory）。

每个线程网格（Grid）中存放着若干线程块（Block）。
1. 一个线程网格（Grid）中存放着多个线程块（Block），但是同时执行的线程块收到流处理器数量的限制。
2. 一个线程网格（Grid）所执行的核函数是相同的，例如：Grid0中的核函数为kernel_func0()，那么Grid0中的所有Block都执行这个核函数。

  结合硬件来看，每个SM中同一时刻只能运行一个线程块（Block），如上图所示，拥有更多的SM的GPU执行线程块（Block）的速度更快。当然，这两块GPU的SM必须是相同规格的。

四、通过代码获取GPU参数

  由于不同型号的GPU核心的SM、SP等数量还不一样，在我们执行计算程序之前，我们最好先获取该GPU的硬件参数，以保证我们的程序能够正常地执行。

  拿我的电脑举例，通过cudaDeviceProp这个类，我们可以获得对应GPU的硬件信息，可以看到显卡型号是K620，以及5.0的CUDA算力，我们需要重点注意一下这两个参数：
1. maxThreadsPerBlock，每个Block的最大线程数，我们在核函数中的线程下标不能超过这个值。
2. maxThreadsPerMultiProcessor，每个SM的最大线程数，即能够同时运行最大线程数。

// CUDA#include "cuda_runtime.h"#include "device_launch_parameters.h"// includes CUDA Runtime#include <cuda_runtime.h>// includes, project#include <helper_cuda.h>#include <helper_functions.h> // helper utility functions // C IO#include <stdio.h>// C++ IOstream#include <iostream>using namespace std;int main(int argc, char *argv[]){    cudaError_t cudaStatus;    void check_Cuda_information(int main_argc, char ** main_argv);    // 读取、检查设备信息    check_Cuda_information(argc, argv);    // 计算部分    //cudaStatus = caculate_Cuda_function();    // cudaDeviceReset must be called before exiting in order for profiling and    // tracing tools such as Nsight and Visual Profiler to show complete traces.    cudaStatus = cudaDeviceReset();    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "cudaDeviceReset failed!");        return 1;    }    return 0;}// 检查显卡硬件属性void check_Cuda_information(int main_argc, char ** main_argv){    // 设备ID    int devID;    // 设备属性    cudaDeviceProp deviceProps;    // 获取设备ID    devID = findCudaDevice(main_argc, (const char **)main_argv);    //     checkCudaErrors(cudaGetDeviceProperties(&deviceProps, devID));    cout << "devID = " << devID << endl;    // 显卡名称    cout << "CUDA device is \t\t\t" << deviceProps.name << endl;    // 每个 线程块（Block）中的最大线程数    cout << "CUDA max Thread per Block is \t" << deviceProps.maxThreadsPerBlock << endl;    // 每个 多处理器组（MultiProcessor）中的最大线程数    cout << "CUDA max Thread per SM is \t" << deviceProps.maxThreadsPerMultiProcessor << endl;    // 线程束大小    cout << "CUDA Warp size is \t\t" << deviceProps.warpSize << endl;    // 每个SM中共享内存的大小    cout << "CUDA shared memorize is \t" << deviceProps.sharedMemPerMultiprocessor << "\tbyte" << endl;}

参考：

1.Ubuntu12.04配置NVIDIA cuda5.5经验帖

2.CUDA在Windows的安装和使用

3.《CUDA并行程序设计》机械工业出版社

4.CUDA C Programming Guide

5.Fermi架构白皮书

6.CUDA系列学习（一）An Introduction to GPU and CUDA

7.GM107核心的白皮书 NVIDIA GeForce GTX 750 Ti Whitepaper