[CUDA编程]基础入门例程1

前言：在学习CUDA编程前有必要先了解GPU的结构，例如SM、SP、memory等
例程地址：https://github.com/TycoonL/awesome-cuda.git/

CUDA中内存组织图：

CUDA程序的基本框架：

头文件包含
常量（宏）定义
C++自定义函数和CUDA核函数的定义和声明
int main(void)
{
	分配主机和设备内存
	初始化主机中的数据
	将某些数据复制到设备中
	调用核函数怎在设备中运算
	将某些数据从设备中复制到主机
	释放内存
}

1.basic_1：用核函数查看device的参数

//打印GPU设备信息
void print_GPU_device_info()
{
	int deviceCount;
	cudaGetDeviceCount(&deviceCount);
	for (int i = 0; i < deviceCount; i++)
	{
		cudaDeviceProp devProp;
		cudaGetDeviceProperties(&devProp, i);
		std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
		std::cout << "设备全局内存总量： " << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
		std::cout << "SM的数量：" << devProp.multiProcessorCount << std::endl;
		std::cout << "每个SM的最大线程数：" << devProp.maxThreadsPerMultiProcessor << std::endl;
		std::cout << "每个SM的最大线程束数：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
		std::cout << "每个Block的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
		std::cout << "每个Block的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;
		std::cout << "每个Block中可用的32位寄存器数量： " << devProp.regsPerBlock << std::endl;
		std::cout << "======================================================" << std::endl;

	}
}

2.basic_2：使用线程索引

一个核函数可以指派多个线程，而这些线程的组织结构由执行配置**<<<grid_size,block_size>>>**决定

• 每个线程在核函数中都有一个唯一的身份标识

• grid_size→gridDim(数据类型：dim3 （x，y，z）); block_size→blockDim; 0<=blockIdx<gridDim; 0<=threadIdx<blockDim

• 在一维情况下：
  $$threadID=blockIdx.x\times blockDim.x+threadIdx.x$$

//核函数 
__global__ void hello_from_device()
{
	int bid = blockIdx.x;
	int tid = threadIdx.x;
	int threadID = blockIdx.x*blockDim.x + threadIdx.x;
	printf("hello world from block %d and thread %d and threadId %d\n", bid, tid,threadID);
}

hello_from_device<<< 2,3 >>>()的结果：
  hello world from block 1 and thread 0 and threadId 3
  hello world from block 1 and thread 1 and threadId 4
  hello world from block 1 and thread 2 and threadId 5
  hello world from block 0 and thread 0 and threadId 0
  hello world from block 0 and thread 1 and threadId 1
  hello world from block 0 and thread 2 and threadId 2

上面是一维的情况，下面介绍用结构体dim3来定义“多维”的网络和线程块：

	#dim3 _size(x,y,z)
	dim3 grid_size(2,2)  //等价于 (2,2,1)
	dim3 block_size(3,2)

对应的一维指标为：

	int tid=threadIdx.z*blockDim.x*blockDim.y+threadIdx.y*blockDim.x+threadIdx.x
	int bid=blockIdx.z*gridDim.x*gridDim.y+blockIdx.y*gridDim.x+blockIdx.x

注意：

• 网格大小在x,y,z三个方向上要分别小于$2^{31}-1$,65535,65535
• 块线程大小在x,y,z三个方向上要分别小于1024，1024，64
• 另外，块线程在三个维度上的乘积还要小于1024，即：一个块最多有1024线程