《GPU高性能编程CUDA实战》学习笔记 (2)_为什么每个线程格最多只能有65535个线程块

高性能计算——GPU篇
https://blog.csdn.net/kkk584520/column/info/hpc-gpu
CUDA入门（四）Visual Profiler
https://blog.csdn.net/qq_25819827/article/details/52107570

cudaThreadExit();
1

CUDA编程——GPU架构，由sp，sm，thread，block，grid，warp说起
https://blog.csdn.net/junparadox/article/details/50540602
https://blog.csdn.net/yu132563/article/details/50301725

• 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有非常严格的限制，也就限制了并行能力。
  每个SM包含的SP数量依据GPU架构而不同

kenerl函数执行：
一个kernel程式会有一个grid，grid底下又有数个block，每个block是一个thread群组。在同一个block中thread可以通过共享内存（shared memory） 来通信，同步。而不同block之间的thread是无法通信的。

CUDA的设备在实际执行过程中，会以block为单位。把一个个block分配给SM进行运算；而block中的thread又会以warp（线程束）为单位，对thread进行分组计算。
warp是调度和运行的基本单元。一个warp需要占用一个SM运行，多个warps需要轮流进入SM。由SM的硬件warp scheduler负责调度。目前每个warp包含32个threads（Nvidia保留修改数量的权利），同一个warp中的thread执行的指令是相同的，只是处理的数据不同。一个GPU上resident thread最多只有 SM*warp个。

第4章 CUDA C并行编程

修饰符说明：

• _global__：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
• __device__：在device上执行，单仅可以从device中调用，不可以和__global__同时用。
• __host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。

核函数参数说明

核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量。
add<<<N,1>>>( dev_a, dev_b, dev_c );
N表示设备在执行核函数时使用的并行线程块的数量。
并行线程块集合也称为一个线程格(Grid)。线程格既可以是一维的线程块集合，也可以是二维的线程块集合。
GPU有着完善的内存管理机制，它将强行结束所有违反内存访问规则的进程。
在启动线程块数组时，数组每一维的最大数量都不能超过65535.这是一种硬件限制,如果启动的线程块数量超过了这个限值,那么程序将运行失败。
注：blockIdx.x是一个内置变量，包含的值是当前执行设备代码的线程块的索引，第一个线程块的blockIdx.x为0，最后一个为N-1
CUDA的软件架构由网格（Grid）、线程块（Block）和线程（Thread） 组成，相当于把GPU上的计算单元分为若干（2~3）个网格，每个网格内包含若干（65535）个线程块，每个线程块包含若干（512）个线程，三者的关系如下图：

• thread：一个CUDA的并行程序会被以许多个threads来执行。
• block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过共享存储器shared memory通信。
• grid：多个blocks则会再构成grid。
  
  硬件将线程块的数量限制为不超过65535.同样，对于启动核函数时每个线程块中的线程数量，硬件也进行了限制。具体来说，最大的线程数量不能超过设备属性结构中maxThreadsPerBlock域的值。这个值并不固定，有的是512，有的是1024.

1：每个block 最大1024个线程（视不同的卡来定），这个是线程总数的限制。

2：每个线程块最大维度的限制为x方向1024，y方向1024，z方向64（视不同的卡来定）。

3：一个线程块的线程情况同时收到上述两条的约束，即，如果在x方向排布了1024个线程，那么y和z方向的维度只能是1，否则将超出第一条的约束。

• 内置变量blockDim，对于所有线程块来说，这个变量是一个常数，保存的是线程块中每一维的线程数量。

• 内置变量gridDim，对于所有线程块来说，这个变量是一个常数，用来保存线程格每一维的大小，即每个线程格中线程块的数量。

• 内置变量blockId.x，变量中包含的值就是当前执行设备代码的线程块的索引。

• 内置变量threadId.x，变量中包含的值就是当前执行设备代码的线程索引。

• gridDim.x：线程网格x维度上线程块的数量
  gridDim.y：线程网格y维度上线程块的数量
  blockDim.x：一个线程块x维度上的线程数量
  blockDlm.y：一个线程块y维度上的线程数量
  theadIdx.x：线程块x维度上的线程索引
  theadIdx.y：线程块y维度上的线程索引

// 基于GPU的矢量求和
#include<stdio.h>
#include <iostream>
#include<device_launch_parameters.h>

#define N   10

__global__ void add(int *a, int *b, int *c) {
	int tid = blockIdx.x;    // blockIdx.x是一个内置变量，包含的值是当前执行设备代码的线程块的索引，第一个线程块的blockIdx.x为0，最后一个为N-1
	if (tid < N)
		c[tid] = a[tid] + b[tid];
}

int main(void) {
	int a[N], b[N], c[N];
	int *dev_a, *dev_b, *dev_c;

	// allocate the memory on the GPU
	cudaMalloc((void**)&dev_a, N * sizeof(int));
	cudaMalloc((void**)&dev_b, N * sizeof(int));
	cudaMalloc((void**)&dev_c, N * sizeof(int));

	// fill the arrays 'a' and 'b' on the CPU
	for (int i = 0; i<N; i++) {
		a[i] = -i;
		b[i] = i * i;
	}

	// copy the arrays 'a' and 'b' to the GPU
	cudaMemcpy(dev_a, a, N * sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, N * sizeof(int),cudaMemcpyHostToDevice);

	add << <N, 1 >> >(dev_a, dev_b, dev_c);

	// copy the array 'c' back from the GPU to the CPU
	cudaMemcpy(c, dev_c, N * sizeof(int),cudaMemcpyDeviceToHost);

	// display the results
	for (int i = 0; i<N; i++) {
		printf("%d + %d = %d\n", a[i], b[i], c[i]);
	}

	// free the memory allocated on the GPU
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	getchar();
	return 0;
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50

// 有error处理的版本
#include "cuda_runtime.h"  
#include "device_launch_parameters.h"  
#include <stdio.h>  

cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);

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

int main()
{
	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 };

	// Add vectors in parallel.  
	cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "addWithCuda failed!");
		return 1;
	}

		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]);

	// cudaThreadExit must be called before exiting in order for profiling and  
	// tracing tools such as Nsight and Visual Profiler to show complete traces.  
	cudaStatus = cudaThreadExit();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaThreadExit failed!");
		return 1;
	}
	getchar();
	return 0;
}

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

	// Choose which GPU to run on, change this on a multi-GPU system.  
	cudaStatus = cudaSetDevice(0);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
		goto Error;
	}

	// Allocate GPU buffers for three vectors (two input, one output)    .  
	cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	// Copy input vectors from host memory to GPU buffers.  
	cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

	cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

	// Launch a kernel on the GPU with one thread for each element.  
	addKernel << <1, size >> >(dev_c, dev_a, dev_b);

	// cudaThreadSynchronize waits for the kernel to finish, and returns  
	// any errors encountered during the launch.  
	cudaStatus = cudaThreadSynchronize();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
		goto Error;
	}

	// Copy output vector from GPU buffer to host memory.  
	cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

Error:
	cudaFree(dev_c);
	cudaFree(dev_a);
	cudaFree(dev_b);


	return cudaStatus;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114