星星博客 »  > 

CUDA自学-Wrap

Wrap

为什么需要wrap

虽然GPU的Grid和Block的大小很大,可以拥有上万级别的线程,但因为硬件的限制,不是所有线程都是可以平行运行的。在运行thread的时候, thread会被捆绑到一起形成一个wrap。32个thread一个wrap。 同一个wrap里的introduction 是一样的,也就是他们运行的东西是一摸一样的,数据也相同。一个wrap里的线程只允许在同一个block里面运行。为了让程序的运行更加有效,需要让同一个wrap里的线程运行同样的代码。

If语句的效率影响

如果我们有if else的语句怎么办?

__global__ void code_with_divergence()
{
	int idx = threadIdx.x;
	if(idx%2==0)
	{
		 // do A
	}
	else
	{
		// do B
	}
}

上面这个这个代码会让效率减少一半。因为当运行A的时候,wrap scheudler会让满足条件的那一半thread运行,而另一半的thread会被休眠。

需要注意的是,不是写了if语句就一定会让运行效率降低。只要能保证用一个wrap里的线程运行同样的introdution就可以,比如如下代码:

__global__ void code_without_divergence()
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	int wrap_id = idx/32;
	if(wrap_id%2==0)
	{
		 // do A
	}
	else
	{
		// do B
	}
}

上面的代码是以一个wrap为整体去运行的,所以不会影响运行效率。

测试上面程序的主程序:

int main()
{
	int size = 1<<22;

	dim3 block_size(128);
	dim3 grid_size((size + block_size.x -1) / block_size.x);

	code_without_divergence <<<grid_size, block_size>>> ();
	cudaDeviceSynchronize();

	code_with_divergence <<<grid_size, block_size>>> ();
	cudaDeviceSynchronize();
	
	cudeDeviceReset();
	return 0;
}

Branch Efficiency

这是用来描述线程效率的参数,具体计算方式可以看附件,学习材料里的 wrap+divergence的ppt。

但是这个参数是可以通过nvprof去计算的。首先你需要编译你的cu文档,

cmd>> nvcc -o myFile.out myFile.cu

然后:

cmd>> nvprof --metrics branch_efficieny myFile.out

–metrics 是option, branch_efficieny 是查看内容的指令。

你会发现,两个程序的Branch Efficiency 是一样的因为compiler自动帮识别到了 divergence然后帮我们优化了。😜

我们可以在编译的时候可以通过选项 -G取消所有编译优化,然后再对比:

cmd>> nvcc -G -o myFile.out myFile.cu

code_without_divergence的效率会是100%,但对于code_with_divergence我们最终还是会得到一个比50%大的数值。因为计算Branch Efficiency的时候不只是查看if里面的语句,还包括基本的比如idx的赋值语句等。此外,compiler还是会做一些最基本的优化,所以最终还是会大于50%。

相关文章