深度学习细粒度分类有什么改进方法细粒度并行算法

转载

mob64ca140c75c7 2024-02-21 19:46:18

文章标签 深度学习细粒度分类有什么改进方法 cuda cg: #include 线程组 文章分类 深度学习人工智能

问题

CUDA的线程执行单元是以warp来划分的，一个warp内部包含32个线程，这32个线程存在一个隐式的线程同步。而不同warp之间是不存在隐式同步的。在一个block中，往往存在多个warp，倘若在程序中使用同步机制"__syncthreads()"令线程同步，此时可能会发生条件竞争的问题。导致指令延迟过高，性能变低，如下所示。

__global__ void test()
{
	const int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < 64)
	{
		printf("idx = %d,顺利到达。\n", idx);
		__syncthreads();
	}
}

需求

如果想只让一个块中的某些线程进行同步的话，就需要更细粒度的并行等待操作。

方法

使用CUDA的Cooperative Groups操作。头文件如下所示：

#define __CUDACC__ //该定义在开发阶段使用，方便代码自动补全
#include <cooperative_groups.h>

使用说明

对于#define __CUDACC__ 这条语句特殊说明下：由于主机端代码和设备端代码在编译的起始，是分开编译的。在设备端编译时，编译器会自定预定义__CUDACC__。但是开发阶段是在主机端，因此，可以在开发时先自定义一个，在开发完毕后再注释掉。不过不注释掉也没关系，只会警告你重复定义。

下面说明如何使用cooperative_groups

namespace cg = cooperative_groups;  //工作区名字太长了 ，这样方便。
__global__ void binomial()
{
	cg::thread_block cta = cg::this_thread_block(); //这个块中的所有线程。特指块。这个块根据包含了线程组织分布的所有信息。
	cg::thread_group tg = cg::this_thread();  //这个线程(组)，特指当前线程(组)。这个后面说
	
	if (threadIdx.x ==0)
	{
		printf("cta.size() = %d\n", cta.size());
	
		printf("tg.size() = %d\n", tg.size());
	}
}

对于cta和tg其内部有成员函数，如下所示

cta.size();  //返回值unsigned int.块的大小，即块内线程的个数
cta.group_dim();//返回值dim3.所有块在网格中的分布。，分别是xyz三个维度的大小。
cta.group_index();//返回值dim3.当前块在网格中的分布。，分别是xyz三个维度的索引。
cta.thread_index();//返回值dim3.当前线程在块中的xyz三个维度的索引
cta.thread_rank();//返回值unsigned int.当前线程是在当前块中是第几个线程。
cta.sync(); //作用等用于 __syncthreads();令当前线程

tg.size();//返回值unsigned int.块的大小，即线程(组)的个数
tg.sync();//同步线程(组)；
tg.thread_rank();//在线程(组)中的序号；

下述明明也是该空间内的，试过多种方式，但是仍然无法使用。具体原因未调查详细。使用下述会发生错误。暂时不用。

cg::grid_group gg =  cg::this_grid();
cg::multi_grid_group mgg = cg::this_multi_grid();

深度学习细粒度分类有什么改进方法细粒度并行算法_cg:

函数介绍

注意到，在前面陈述中说道 tg是线程(组)，括号里面有一个组字。这个线程组也就是更细粒度的线程划分，该组大小小于等于32(warp的大小)，且该组必须是2的幂次。

//倘若直接使用,那么一个线程就是一个线程组。
cg::thread_group tg = cg::this_thread();
//划分。由线程块划分为线程组。注意动词前后名词。
cg::thread_group tg = cg::tiled_partition(cta, 2); //将cta划分为以2个线程为一组的线程组。
//此时
tg.size();//返回值unsigned int.块的大小，即线程(组)的个数     2
tg.sync();//同步线程(组)；								对2个线程同步
tg.thread_rank();//在线程(组)中的序号；					  0、1
//只能划分为下述。下面语句可以直接抄
thread_group tile = cg::partition(this_thread_block(), 1);  //其实就是不划分。写出来是为了好看一点点
thread_group tile2 = cg::partition(this_thread_block(), 2);
thread_group tile4 = cg::partition(this_thread_block(), 4);
thread_group tile8 = cg::partition(this_thread_block(), 8);
thread_group tile16 = cg::partition(this_thread_block(), 16);
thread_group tile32 = cg::partition(this_thread_block(), 32);  //作用等用于warp，warp存在隐形同步。

//线程组也可以继续划分为更细粒度
cg::thread_group tile4 = cg::tiled_partition(tile32, 4);

到此说明结束。详细的案例如下所示，参考官方链接：Cooperative Groups: Flexible CUDA Thread Programming | NVIDIA Developer Blog

代码示例

//#define __CUDACC__

#include <iostream>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <helper_cuda.h>
#include <cooperative_groups.h>
#include <cooperative_groups_helpers.h>

namespace cg = cooperative_groups;

__global__ void binomial()
{
	cg::thread_block cta = cg::this_thread_block();
	
	cg::thread_group tile32 = cg::tiled_partition(cg::this_thread_block(), 32);
	cg::thread_group tile4 = cg::tiled_partition(tile32, 4);
	//cg::thread_group tile32 = cg::tiled_partition(cta, 32);
	
	//cg::grid_group gg =  cg::this_grid();
	//cg::multi_grid_group mgg = cg::this_multi_grid();

	cta.size();
	cta.group_dim();
	cta.group_index();
	cta.thread_index();
	cta.thread_rank();
	cta.sync();
	tile4.size();
	tile4.sync();
	tile4.thread_rank();
	if (threadIdx.x ==0)
	{
		printf("cta.size() = %d\n", cta.size());
		//printf("cta.group_dim(),%d,%d,%d\n", cta.group_dim().x, cta.group_dim().y, cta.group_dim().z);
		printf("tg.size() = %d\n", tile4.size());
		//printf("gg.size() = %d\n", gg.size());
	}
	//printf("cta.thread_index(),%d,%d,%d\n", cta.thread_index().x, cta.thread_index().y, cta.thread_index().z);
	//printf("cta.group_index(),%d,%d,%d\n", cta.group_index().x, cta.group_index().y, cta.group_index().z);
	printf("cta.thread_rank() = %d\n", cta.thread_rank());
	printf("tg.thread_rank() = %d\n", tile4.thread_rank());

	//unsigned int a = cg::group_size(cta);
	//printf("group_size = %d", a);
	//cg::thread_group tile2 = cg::tiled_partition(cta, 32);  //对线程进行划分，大小应该是2的幂次且不能大于32，否则报错。无法执行后续
	//dim3 temp = cta.thread_index();
	//printf("Dim3 from thread_index %d %d %d\n", temp.x, temp.y, temp.z);
	tile2.
	//printf("thread_rank() %d this_thread_block().thread_rank() %d\n", tile2.thread_rank(), cg::this_thread_block().thread_rank());
}

__device__ int thread_sum(int* input, int n)
{
	int sum = 0;

	for (int i = blockIdx.x * blockDim.x + threadIdx.x;
		i < n / 4;
		i += blockDim.x * gridDim.x)
	{
		int4 in = ((int4*)input)[i];
		sum += in.x + in.y + in.z + in.w;
	}
	return sum;
}

template <typename group_t>
__device__ int reduce_sum(group_t g, int* temp, int val)
{
	int lane = g.thread_rank();

	// Each iteration halves the number of active threads
	// Each thread adds its partial sum[i] to sum[lane+i]
#pragma unroll
	for (int i = g.size() / 2; i > 0; i /= 2)
	{
		temp[lane] = val;
		g.sync(); // wait for all threads to store
		if (lane < i) val += temp[lane + i];         //寄存器使用比共享内存使用快很多。
		g.sync(); // wait for all threads to load
	}

	return val; // note: only thread 0 will return full sum
}

__global__ void sum_kernel_32(int* sum, int* input, int n)
{
	int my_sum = thread_sum(input, n);  

	extern __shared__ int temp[];

	auto g = cg::this_thread_block();
	auto tileIdx = g.thread_rank() / 32;
	int* t = &temp[32 * tileIdx];

	auto tile32 = cg::tiled_partition(g, 32);
	int tile_sum = reduce_sum(tile32, t, my_sum);

	if (tile32.thread_rank() == 0) atomicAdd(sum, tile_sum);
}

int main()
{
	//测试 cg
	/*dim3 block(4, 8);
	dim3 grid(2, 2);
	binomial << <grid, block >> > ();
	cudaDeviceSynchronize();*/

	int n = 1 << 24;
	int blockSize = 256;
	int nBlocks = (n + blockSize - 1) / blockSize;
	int sharedBytes = blockSize * sizeof(int);

	int* sum, *data;
	cudaMallocManaged(&sum, sizeof(int));
	cudaMallocManaged(&data, n * sizeof(int));
	std::fill_n(data, n, 1); // initialize data
	cudaMemset(sum, 0, sizeof(int));

	sum_kernel_32 << <nBlocks/4, blockSize, sharedBytes >> > (sum, data, n);
	cudaDeviceSynchronize();
	printf("sum = %d\n", *sum);
	
	return 0;
}

本文章为转载内容，我们尊重原作者对文章享有的著作权。如有内容错误或侵权问题，欢迎原作者联系我们进行内容更正或删除文章。