初学如何选择thread数量

本文以求两个矢量的平方和矢量的程序来简单分析如何选择thread的尺寸，使用的环境为win7+cuda+GTX1050显卡

代码如下：

#include< stdio.h>
//#include <thrust/device_vector.h>
#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#define N 10000
#define THREAD_NUM 1023

bool InitCUDA()
{
	int count;
	cudaGetDeviceCount(&count);
	if (count == 0)
	{
		fprintf(stderr, "There is no device.\n");
		return false;
	}
	int i;
	for (i = 0; i < count; i++)
	{
		cudaDeviceProp prop;
		if (cudaGetDeviceProperties(&prop, i) == cudaSuccess)
		{
			if (prop.major >= 1)
			{
				break;
			}
		}
	}
	if (i == count)
	{
		fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
		return false;
	}
	cudaSetDevice(i);
	return true;
}


__global__ void add(int *a, int *b, int *c)
{
	int tid = threadIdx.x + blockIdx.x*blockDim.x;
	int size = N / THREAD_NUM;
	//for(int i=0; i<size; i++)
	{
		if (tid<N)
		{
			//c[tid+i*THREAD_NUM]=a[tid+i*THREAD_NUM]*a[tid+i*THREAD_NUM]+b[tid+i*THREAD_NUM]*b[tid+i*THREAD_NUM];
			c[tid] = a[tid] * a[tid] + b[tid] * b[tid];
		}
	}
}

int main()
{
	if (!InitCUDA())
	{
		return 0;
	}
	printf("HelloWorld, CUDA has been initialized.\n");

	int a[N], b[N], c[N];
	/*int *a = new int[N];
	int *b = new int[N];
	int *c = new int[N];*/
	int *dev_a, *dev_b, *dev_c;
	for (int i = 0; i<N; i++)
	{
		a[i] = i % 5;
		b[i] = i % 5;
	}
	cudaMalloc((void**)&dev_a, N * sizeof(int));
	cudaMalloc((void**)&dev_b, N * sizeof(int));
	cudaMalloc((void**)&dev_c, N * sizeof(int));

	cudaMemcpy(dev_a, &a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, &b, N * sizeof(int), cudaMemcpyHostToDevice);

	add << <(N + THREAD_NUM - 1) / THREAD_NUM, THREAD_NUM >> >(dev_a, dev_b, dev_c);

	cudaMemcpy(&c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
	int count = 0;
	for (int i = 0; i < N; i++)
	{
		
		if (c[i] != a[i] * a[i] + b[i] * b[i])
		{
			printf("%d + %d = %d \n", a[i], b[i], c[i]);
			count++;
		}
	}
	printf("count=%d\n", count);

	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	//delete[] a;
	//delete[] b;
	//delete[] c;

	system("pause");
	return 0;
}

总的数据量为N    单个块中线程数THREAD_NUM，如下

#define N 10000

#define THREAD_NUM 1

核函数为

__global__ void add(int *a, int *b, int *c)
{
	int tid = threadIdx.x + blockIdx.x*blockDim.x;
	int size = N / THREAD_NUM;
	//for(int i=0; i<size; i++)
	{
		if (tid<N)
		{
			//c[tid+i*THREAD_NUM]=a[tid+i*THREAD_NUM]*a[tid+i*THREAD_NUM]+b[tid+i*THREAD_NUM]*b[tid+i*THREAD_NUM];
			c[tid] = a[tid] * a[tid] + b[tid] * b[tid];
		}
	}
}

现设置THREAD_NUM 为 1 10 100 640 1024时所耗时为

由图可见，在CUDA允许的每个块中最大的线程数时计算速度最快，当超过该数时结果会出错。当每块线程数比较小时增加线程数对计算速度有较大提升，线程数较多时提升不明显。

Occupancy实际为85.67%，最高100%，活跃warps54.83，共有64个，可见warps活跃率比较高，性能较好，但离性能高峰还有15%左右的距离。

可以看到左图实现过程中有60000个加法 ，181440个乘法，20000个移位运算，可以使用中间变量或合理调整算法，减少运算量或耗时较长的运算。可以看出实际的运算量不太好使用数kernel代码中的各种运算，原因是有些运算被分成了好多步骤，如32位乘法运算可能被分为了多次24位乘法运算和求和运算（这里还不清楚）。

右图是运算的速度，可以此结果进行算法效率的评估和优化。

对于CUDA程序的性能分析还是比较重要的，目前还在学习分析报表的各种项，有机会再专门了解这方面内容吧。