CUDA——线程束分化

本文是阅读《CUDA C 编程权威指南》所做的笔记

1. 线程束分化

    线程束是SM中基本的执行单元。 当一个线程块的网格被启动后， 网格中的线程块分布在SM中。 一旦线程块被调度到一个SM上， 线程块中的线程会被进一步划分为线程束。一个线程束由32个连续的线程组成， 在一个线程束中， 所有的线程按照单指令多线程（SIMT） 方式执行； 也就是说， 所有线程都执行相同的指令， 每个线程在私有数据上进行操作。 下图展示了线程块的逻辑视图和硬件视图之间的关系。

    GPU是相对简单的设备， 它没有复杂的分支预测机制。 一个线程束中的所有线程在同一周期中必须执行相同的指令， 如果一个线程执行一条指令， 那么线程束中的所有线程都必须执行该指令。 如果在同一线程束中的线程使用不同的路径通过同一个应用程序， 这可能会产生问题。思考下面语句：

if(cond)
{
    ...
}
else
{
    ...
}

    一半的线程束需要执行if语句块中的指令， 而另一半需要执行else语句块中的指令。 在同一线程束中的线程执行不同的指令， 被称为线程束分化。
    如果一个线程束中的线程产生分化， 线程束将连续执行每一个分支路径， 而禁用不执行这一路径的线程。 线程束分化会导致性能明显地下降。 在前面的例子中可以看到， 线程束中并行线程的数量减少了一半： 只有16个线程同时活跃地执行， 而其他16个被禁用了。条件分支越多， 并行性削弱越严重。
    注意， 线程束分化只发生在同一个线程束中。 在不同的线程束中， 不同的条件值不会引起线程束分化。

2. 简单的线程束分化代码：

#include <stdio.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <stdlib.h>
#include <time.h>

// kernel1
__global__ void mathKernel1(float *c)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	float ia, ib;
	ia = ib = 0.0f;

	if (tid % 2 == 0) // 偶数线程
	{

		ia = 100.0f;
	}
	else // 奇数线程
	{

		ib = 200.0f;
	}
	c[tid] = ia + ib;
}

// kernel2
__global__ void mathKernel2(float *c)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	float ia, ib;
	ia = ib = 0.0f;
	// 分支粒度是线程束(warpsize = 32)大小的倍数

	if (tid / 32 % 2 == 0) // 偶数
	{

		ia = 100.0f;

	}
	else // 奇数
	{

		ib = 200.0f;
	}
	c[tid] = ia + ib;
}

// Kernel3
__global__ void mathKernel3(float *c)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	float ia, ib;
	ia = ib = 0.0f;
	bool ipred = (tid % 2 == 0);
	if (ipred)
	{

		ia = 100.0f;

	}
	if (!ipred)
	{

		ib = 200.0f;
	}
	c[tid] = ia + ib;
}

// Kernel4
__global__ void mathKernel4(float *c)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	float ia, ib;
	ia = ib = 0.0f;
	int itid = tid >> 5;
	if (itid & 0x01 == 0) // 偶数
	{
		ia = 100.0f;
	}
	else // 奇数
	{

		ib = 200.0f;
	}
	c[tid] = ia + ib;
}

// Kernelwarm
__global__ void warmingup(float *c)
{
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	float ia, ib;
	ia = ib = 0.0f;
	if ((tid / warpSize) % 2 == 0)
	{
		ia = 100.0f;
	}
	else
	{
		ib = 200.0f;
	}
	c[tid] = ia + ib;
}

// 主函数
int main(int argc, char **argv)
{
	// set up device
	int dev = 0;
	cudaDeviceProp deviceProp;
	cudaGetDeviceProperties(&deviceProp, dev);
	printf("%s using Device %d: %s\n", argv[0], dev, deviceProp.name);
	int a, bc, d;
	// set up data size
	int size = 1 << 25;
	int blocksize = 256;

	printf("Data size %d ", size);

	// set up execution configuration
	dim3 block(blocksize, 1);
	dim3 grid((size + block.x - 1) / block.x, 1);
	printf("Execution Configure (block %d grid %d)\n", block.x, grid.x);

	// allocate gpu memory
	float *d_C;
	size_t nBytes = size * sizeof(float);
	cudaMalloc((float**)&d_C, nBytes);
	
	clock_t iStart, iElaps;

	// run a warmup kernel to remove overhead
	cudaDeviceSynchronize();
	iStart = clock();
	warmingup << <grid, block >> > (d_C);
	cudaDeviceSynchronize();
	iElaps = clock();
	float time = (float)(iElaps - iStart) / CLOCKS_PER_SEC;
	printf("warmup      <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time);
	cudaGetLastError();

	// run kernel 1
	iStart = clock();
	mathKernel1 << <grid, block >> > (d_C);
	cudaDeviceSynchronize();
	iElaps = clock();
	float time1 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;
	printf("mathKernel1 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time1);
	cudaGetLastError();

	// run kernel 2
	iStart = clock();
	mathKernel2 << <grid, block >> > (d_C);
	cudaDeviceSynchronize();
	iElaps = clock();
	float time2 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;
	printf("mathKernel2 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time2);
	cudaGetLastError();

	// run kernel 3
	iStart = clock();
	mathKernel3 << <grid, block >> > (d_C);
	cudaDeviceSynchronize();
	iElaps = clock();
	float time3 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;
	printf("mathKernel3 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time3);
	cudaGetLastError();

	// run kernel 4
	iStart = clock();
	mathKernel4 << <grid, block >> > (d_C);
	cudaDeviceSynchronize();
	iElaps = clock();
	float time4 = (float)(iElaps - iStart) / CLOCKS_PER_SEC;
	printf("mathKernel4 <<< %4d %4d >>> elapsed %f sec \n", grid.x, block.x, time4);
	cudaGetLastError();

	// free gpu memory and reset divece
	cudaFree(d_C);
	cudaDeviceReset();
	return EXIT_SUCCESS;
}

3. 运行结果