SP（streaming Process），SM（streaming multiprocessor）是硬件（GPU hardware）概念。而thread，block，grid，warp是软件上的（CUDA）概念。

  • 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有非常严格的限制，也就限制了并行能力。

CUDA架构专门为GPU计算设计了一种全新的模块，目的是减轻早期GPU计算中存在的一些限制，而正是这些限制使得之前的GPU在通用计算中没有得到广泛的应用。
使用CUDA C来编写代码的前提条件包括：
(1)、支持CUDA的图形处理器，即由NVIDIA推出的GPU显卡，要求显存超过256MB；
(2)、NVIDIA设备驱动程序，用于实现应用程序与支持CUDA的硬件之间的通信，确保安装最新的驱动程序，注意选择与开发环境相符的图形卡和操作系统；
(3)、CUDA开发工具箱即CUDA Toolkit，此工具箱中包括一个编译GPU代码的编译器；
(4)、标准C编译器，即CPU编译器。CUDA C应用程序将在两个不同的处理器上执行计算，因此需要两个编译器。其中一个编译器为GPU编译代码，而另一个为CPU编译代码。

第三章 第一段CUDA C代码+Host/Device

3.1 第一个程序

3.1.1 Hello World！

一般，将CPU以及系统的内存称为主机(Host)，而将GPU及其内存称为设备(Device)。在GPU设备上执行的函数通常称为核函数(Kernel)。

#include<stdio.h>
#include <iostream>

__global__ void kernel(void) {
}

int main() {
  kernel <<<1, 1>>>();
  printf("hello world!");
  getchar();
  return 0;
}

这个程序与最初的“ Hello， World！”相比，多了两个值得注意的地方：

  • 一个空的函数kernel()，并且带有修饰符__global__ 。这个修饰符将告诉编译器，函数应该编译为在设备而不是主机上运行。在这个简单的示例中，函数kernel()将被交给编译设备代码的编译器，而main()函数将被交给主机编译器。CUDA编译器和运行时将负责实现从主机代码中调用设备代码。
  • 对这个空函数的调用，并且带有修饰字符 <<<1,1>>> 。

3.1.2 传递参数给核函数

#include <stdio.h>
#include <iostream>
#include <device_launch_parameters.h>

__global__ void add(int a, int b, int *c) {
  *c = a + b;
}

int main()
{
  int a = 2, b = 3, c = 0;
  int* dev_c;
  cudaMalloc((void**)&dev_c, sizeof(int));

  //尖括号表示要将一些参数传递给CUDA编译器和运行时系统
  //尖括号中这些参数并不是传递给设备代码的参数，而是告诉运行时如何启动设备代码,
  //传递给设备代码本身的参数是放在圆括号中传递的，就像标准的函数调用一样
  add << <1, 1 >> > (a, b, dev_c);
  cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);

  printf("%d + %d = %d\n", a, b, c);
  cudaFree(dev_c);
  getchar();
  return 0;
}

注意这里增加了多行代码，在这些代码中包含两个概念：
• 可以像调用C函数那样将参数传递给核函数。
• 当设备执行任何有用的操作时，都需要分配内存，例如将计算值返回给主机。

通过cudaMalloc() 来分配内存。
这个函数调用的行为非常类似于标准的C函数malloc()，但该函数的作用是告诉CUDA运行时在设备上分配内存。

  • 第一个参数是一个指针，指向用于保存新分配内存地址的变量，
  • 第二个参数是分配内存的大小。除了分配内存的指针不是作为函数的返回值外，这个函数的行为与malloc()是相同的，并且返回类型为void*。

cudaMalloc函数使用限制总结：
1.可以将cudaMalloc()分配的指针传递给在设备上执行的函数。
2.可以在设备代码中使用cudaMalloc()分配的指针进行内存读/写操作。
3.可以将cudaMalloc()分配的指针传递给在主机上执行的函数。
4.不能在主机代码中使用cudaMalloc()分配的指针进行内存读/写操作。
要释放cudaMalloc()分配的内存，需要调用cudaFree()
设备指针的使用方式与标准C中指针的使用方式完全一样。 主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存。

3.1.3 查询设备

有可能在单块卡上包含了两个或多个GPU。
在主机代码中可以通过调用cudaMemcpy() 来访问设备上的内存，在设备和主机之间复制数据。
获得CUDA设备的数量，可以调用cudaGetDeviceCount()

int count;
HANDLE_ERROR( cudaGetDeviceCount( &count ) );

在集成的GPU上运行代码，可以与CPU共享内存。
计算功能集的版本为1.3或者更高的显卡才能支持双精度浮点数的计算。

3.1.4 设备属性的使用

cudaGetDeviceProperties()
查找主版本号大于1，或者主版本号为1且次版本号大于等于3的设备：
首先，找出我们希望设备拥有的属性并将这些属性填充到一个cudaDeviceProp结构。

cudaDeviceProp prop;
memset( &prop, 0, sizeof( cudaDeviceProp ) );
prop.major = 1;
prop.minor = 3;

在填充完 cudaDeviceProp 结构后，将其传递给 cudaChooseDevice()，这样CUDA运行时将查找是否存在某个设备满足这些条件。 cudaChooseDevice()函数将返回一个设备ID，然后我们可以将这个ID传递给 cudaSetDevice()。随后，所有的设备操作都将在这个设备上执行。

#include "../common/book.h"
int main( void )
{
	cudaDeviceProp prop;
	int dev;
	HANDLE_ERROR( cudaGetDevice( &dev ) );
	printf( "ID of current CUDA device: %d\n", dev );
	memset( &prop, 0, sizeof( cudaDeviceProp ) );
	prop.major = 1;
	prop.minor = 3;
	HANDLE_ERROR( cudaChooseDevice( &dev, &prop ) );
	printf( "ID of CUDA device closest to revision 1.3: %d\n", dev );
	HANDLE_ERROR( cudaSetDevice( dev ) );
}