本文将通过一个程序帮助了解线程块的分配,以及线程束,线程全局标号等
#include<cuda_runtime.h>
#include<conio.h>
#include<stdio.h>
#include<stdlib.h>
#include<device_launch_parameters.h>
#define ARRAY_SIZE 128
#define ARRAY_SIZE_IN_BYTES (sizeof(unsigned int)*(ARRAY_SIZE))
__global__ void what_is_my_id(unsigned int *const block,
unsigned int *const thread,
unsigned int *const warp,
unsigned int *const calc_thread)
{
const unsigned int thread_idx = blockIdx.x*blockDim.x + threadIdx.x;
block[thread_idx] = blockIdx.x;
thread[thread_idx] = threadIdx.x;//内部线程的索引
warp[thread_idx] = threadIdx.x / warpSize;
calc_thread[thread_idx] = thread_idx;
}
int main()
{
/* 本地开辟4个数组存放我们要计算的内容 */
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];
//设计线程数为2*64=128个线程
const unsigned int num_blocks = 2;
const unsigned int num_threads = 64;
/* 在GPU上分配同样大小的4个数组 */
unsigned int * gpu_block;
unsigned int * gpu_thread;
unsigned int * gpu_warp;
unsigned int * gpu_calc_thread;
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);
//执行内核函数
what_is_my_id << <num_blocks, num_threads >> >(gpu_block, gpu_thread, gpu_warp, gpu_calc_thread);
//将GPU运算完的结果复制回本地
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);
cudaFree(gpu_block);
cudaFree(gpu_thread);
cudaFree(gpu_warp);
cudaFree(gpu_calc_thread);
//输出
for (unsigned int i = 0; i < ARRAY_SIZE; i++)
{
printf("总线程数%3u-Blocks:%2u-Warp%2u-内部线程数%3u\n",
cpu_calc_thread[i], cpu_block[i], cpu_warp[i], cpu_thread[i]);
}
return 0;
}
可以看到总线程数为0~127,共有2个线程块,每个线程块包含64个 线程,
每个线程块内部线程的索引为0~63.一个线程块包含2个线程束(warp)
(1个warp包括32个线程)
#include<cuda_runtime.h>
#include<stdio.h>
#include<device_launch_parameters.h>
#define ARRAY_SIZE_X 32
#define ARRAY_SIZE_Y 16
#define ARRAY_SIZE_IN_BYTES (sizeof(unsigned int)*(ARRAY_SIZE_X)*(ARRAY_SIZE_Y))
__global__ void what_is_my_id_2d_A(unsigned int *const block_x,unsigned int *const block_y,
unsigned int *const thread,unsigned int *const calc_thread,
unsigned int *const x_thread, unsigned int *const y_thread,
unsigned int *const gird_dimx, unsigned int *const block_dimx,
unsigned int *const gird_dimy, unsigned int *const block_dimy)
{
const unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int idy = blockIdx.y*blockDim.y + threadIdx.y;
const unsigned int thread_idx =((gridDim.x*blockDim.x)*idy) + idx;
block_x[thread_idx] = blockIdx.x;//X维度线程块索引
block_y[thread_idx] = blockIdx.y;//Y维度线程块索引
thread[thread_idx] = threadIdx.x;//1个线程块内部X维度线程索引
calc_thread[thread_idx] = thread_idx;//总索引
x_thread[thread_idx] = idx;
y_thread[thread_idx] = idy;
gird_dimx[thread_idx] = gridDim.x;//线程网格X维度上线程块数量
block_dimx[thread_idx] = blockDim.x;//线程网格Y维度上线程数量
gird_dimy[thread_idx] = gridDim.y;
block_dimy[thread_idx] = blockDim.y;
}
int main()
{
/* 本地开辟4个数组存放我们要计算的内容 */
unsigned int cpu_block_x[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_block_y[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_calc_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_xthread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_ythread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_grid_dimx[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_block_dimx[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_grid_dimy[ARRAY_SIZE_Y][ARRAY_SIZE_X];
unsigned int cpu_block_dimy[ARRAY_SIZE_Y][ARRAY_SIZE_X];
//设计线程数为2*64=128个线程
const dim3 threads_rect(32, 4);
const dim3 blocks_rect(1, 4);
const dim3 threads_square(16, 8);
const dim3 blocks_square(2,2);
/* 在GPU上分配同样大小的4个数组 */
unsigned int * gpu_block_x;
unsigned int * gpu_block_y;
unsigned int * gpu_thread;
unsigned int * gpu_warp;
unsigned int * gpu_calc_thread;
unsigned int * gpu_xthread;
unsigned int * gpu_ythread;
unsigned int * gpu_grid_dimx;
unsigned int * gpu_block_dimx;
unsigned int * gpu_grid_dimy;
unsigned int * gpu_block_dimy;
cudaMalloc((void**)&gpu_block_x, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_block_y, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_thread, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_calc_thread, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_xthread, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_ythread, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_grid_dimx, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_block_dimx, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_grid_dimy, ARRAY_SIZE_IN_BYTES);
cudaMalloc((void**)&gpu_block_dimy, ARRAY_SIZE_IN_BYTES);
//执行条纹式布局左边的图
// what_is_my_id_2d_A << <blocks_rect, threads_rect >> > ( gpu_block_x, gpu_block_y, gpu_thread,gpu_calc_thread,gpu_xthread,gpu_ythread,gpu_grid_dimx,gpu_block_dimx,gpu_grid_dimy,gpu_block_dimy);
//执行方块式布局右边的图
what_is_my_id_2d_A << <blocks_square, threads_square >> >(gpu_block_x, gpu_block_y, gpu_thread,gpu_calc_thread, gpu_xthread,gpu_ythread,gpu_grid_dimx,gpu_block_dimx,gpu_grid_dimy, gpu_block_dimy);
//将GPU运算完的结果复制回本地
cudaMemcpy(cpu_block_x, gpu_block_x, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_block_y, gpu_block_y, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_thread, gpu_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_calc_thread, gpu_calc_thread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_xthread, gpu_xthread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_ythread, gpu_ythread, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_grid_dimx, gpu_grid_dimx, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_block_dimx, gpu_block_dimx, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_grid_dimy, gpu_grid_dimy, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
cudaMemcpy(cpu_block_dimy, gpu_block_dimy, ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
// printf("\nKernel &d\n", kernel);
for (int y = 0; y < ARRAY_SIZE_Y; y++)
{
for (int x = 0; x < ARRAY_SIZE_X; x++)
{
printf("总%3u X维度block索引:%1u Y维度block索引:%1u TID:%2u YTID:%2u XTID:%2uGridX维度上block数量%1u BDX:%1u GridY维度上block数量%1u blockY维度线程数量%1u\n",
cpu_calc_thread[y][x], cpu_block_x[y][x], cpu_block_y[y][x], cpu_thread[y][x], cpu_ythread[y][x], cpu_xthread[y][x],
cpu_grid_dimx[y][x], cpu_block_dimx[y][x],cpu_grid_dimy[y][x], cpu_block_dimy[y][x]);
}
}
// }
cudaFree(gpu_block_x);
cudaFree(gpu_block_y);
cudaFree(gpu_thread);
cudaFree(gpu_calc_thread);
cudaFree(gpu_xthread);
cudaFree(gpu_ythread);
cudaFree(gpu_grid_dimx);
cudaFree(gpu_block_dimx);
cudaFree(gpu_grid_dimy);
cudaFree(gpu_block_dimy);
return 0;
}