CUDA线程块的分配

        为了确保能够真正地了解线程块的分配，接下来我们写一个简短的内核程序来输出线程块、线程、线程束和线程全局标号到屏幕上。现在，除非你使用的是 3.2 版本以上的 SDK否则内核中是不支持 printf的。因此，我们可以将数据传送回 CPU 端然后输出到控制台窗口，内核的代码如下:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>
#include <conio.h>

__global__ void what_is_my_id(unsigned int* const block,
	unsigned int* const thread,
	unsigned int* const warp,
	unsigned int* const calc_thread) {
	/* Thread id is block index * block size + thread offset into the block */
	const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
	block[thread_idx] = blockIdx.x; thread[thread_idx] = threadIdx.x;
	/* Calculate warp using buit in variable warpSize */
	warp[thread_idx] = threadIdx.x / warpSize;
	calc_thread[thread_idx] = thread_idx;
}

#define ARRAY_SIZE 128
#define ARRAY_SIZE_IN_BYTES (sizeof(unsigned int)* (ARRAY_SIZE))

/* Declare statically four arrays of ARRAY_SIZE each */
unsigned int cpu_block[ARRAY_SIZE];
unsigned int cpu_thread[ARRAY_SIZE];
unsigned int cpu_warp[ARRAY_SIZE];
unsigned int cpu_calc_thread[ARRAY_SIZE];

int main(void) {
	/* Total thread count =2*64=128 */
	const unsigned int num_blocks = 2;
	const unsigned int num_threads = 64;
	char ch;

	/* Declare pointers for GPU based params */
	unsigned int* gpu_block;
	unsigned int* gpu_thread;
	unsigned int* gpu_warp;
	unsigned int* gpu_calc_thread;

	/* Declare loop counter for use later */
	unsigned int i;

	/* Allocate four arrays on the GPU */
	cudaMalloc((void**)&gpu_block, ARRAY_SIZE_IN_BYTES);
	cudaMalloc((void**)&gpu_thread, ARRAY_SIZE_IN_BYTES);
	cudaMalloc((void**)&gpu_warp, ARRAY_SIZE_IN_BYTES);
	cudaMalloc((void**)&gpu_calc_thread, ARRAY_SIZE_IN_BYTES);

	/* Execute our kerne] */
	what_is_my_id <<<num_blocks, num_threads>>>(gpu_block, gpu_thread, gpu_warp, gpu_calc_thread);

	/* Copy back the gpu results to the CPU */
	cudaMemcpy(cpu_block, gpu_block, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
	cudaMemcpy(cpu_thread, gpu_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
	cudaMemcpy(cpu_warp, gpu_warp, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
	cudaMemcpy(cpu_calc_thread, gpu_calc_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);

	/* Free the arrays on the GPU as now we're done with them */
	cudaFree(gpu_block);
	cudaFree(gpu_thread);
	cudaFree(gpu_warp);
	cudaFree(gpu_calc_thread);

	/* Iterate through the arrays and print */
	for (i = 0; i < ARRAY_SIZE; i++) {
		printf("Calculated Thread: %3u - Block:%2u - Warp %2u - Thread %3u\n", cpu_calc_thread[i], cpu_block[i], cpu_warp[i], cpu_thread[i]);
	}
	ch = getch();
}

        在这个例子中，我们可以看到线程块按照线程块的编号紧密相连。由于处理的是一维数组，所以我们对线程块采用相同的布局便可简单解决问题。以下是此程序的输出结果:

                  

        正如我们计算的那样，线程索引是0~ 127。一共有两个线块，每个线程块包含 64个线程，每个线程块内部线程的索引为0~63。一个线程块包含两个线束。