CUDA性能优化系列——Kmeans算法调优（二）

最新推荐文章于 2023-06-01 19:43:04 发布
最新推荐文章于 2023-06-01 19:43:04 发布
阅读量1.6k 收藏 3
点赞数
分类专栏： # CUDA性能调优 CUDA 文章标签： cuda gpgpu
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/yan31415/article/details/115153893
版权

本篇介绍Kmeans算法中计算新的聚类中心部分。这部分主要逻辑：根据计算出的新的分类信息，对全部数据点依次对每个类别求出所属当前类别的数据点个数与坐标和。本质上就是进行规约运算。

V1 Atomic实现全局规约

由于最终生成16个聚类中心，因此这里的规约操作需要针对算法进行一定的调整。V1在实现逻辑为：先在共享内存上分别通过原子操作，对16个类进行规约，再通过原子操作进行设备内存的全局规约操作。

/*V1 atomic规约*/
__global__ void kKmeansSumAtomic(int SrcCount,
	int Coords/*[1]*/,
	int ClusterCount/*[16]*/,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	//TODO:local reduce cluster result in share memory
	extern __shared__ float sm[];/*[Coords][blockDim.x] + [ClusterCount][Coords] + [ClusterCount]*/
	for (size_t i = threadIdx.x; i < Coords * blockDim.x + ClusterCount * Coords + ClusterCount; i += blockDim.x)
	{
		sm[i] = 0.0f;
	}
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int* pClusterCount/*[ClusterCount]*/ = (int*)(sm + blockDim.x * Coords + ClusterCount * Coords);
	float* pPointsReduce/*[ClusterCount][Coords]*/ = (float*)(sm + blockDim.x * Coords);
	int CurCluster = d_pMembership[tid];
	for (size_t iCoord = 0; iCoord < Coords; iCoord++)
	{
		sm[threadIdx.x + iCoord * blockDim.x] = d_pSrc[tid + iCoord * SrcCount];
	}
	//Reduce count info
	atomicAdd(pClusterCount + CurCluster, 1);
	__syncthreads();
	if (threadIdx.x < ClusterCount)
	{
		atomicAdd(d_pMemberCount + threadIdx.x, pClusterCount[threadIdx.x]);
	}
	//Reduce points info
	for (size_t iCoord = 0; iCoord < Coords; iCoord++)
	{
		atomicAdd(pPointsReduce + CurCluster * Coords + iCoord, sm[iCoord * blockDim.x + threadIdx.x]);
	}
	__syncthreads();
	if (threadIdx.x < ClusterCount)
	{
		for (size_t iCoord = 0; iCoord < Coords; iCoord++)
		{
			atomicAdd(d_pDst + threadIdx.x * Coords + iCoord, pPointsReduce[threadIdx.x * Coords + iCoord]);
		}
	}
}

执行情况如下
在这里插入图片描述
此时耗时1720us
在这里插入图片描述
block 1024 grid 1024
在这里插入图片描述
调整执行配置后运行时间393us
在这里插入图片描述
grid 8192 block 128
由于原子操作串行执行，因此整个规约计算全部通过原子操作并不是一个很好的选择。这里仅作为一种实现的思路。

V2 分类归约

主要思路：依次按照所属类别中心，在共享内存中对其分别进行归约，最终将每个block的归约的结果通过原子操作更新至设备内存中。

/*V2 分类归约*/
__global__ void kKmeansSum(int SrcCount,
	int Coords,
	int ClusterCount,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	//TODO:local reduce cluster result in share memory
	extern __shared__ float sm[];/*[Coords][blockDim.x] + [blockDim.x]*/
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int* pClusterCount = (int*)(sm + blockDim.x * kCoords);
	int RealCluster = d_pMembership[tid];
	//TODO:sum each cluster
	for (int itcluster = 0; itcluster < ClusterCount; itcluster++)
	{
		//TODO:load
		bool bCurCluster = (RealCluster == itcluster);
		pClusterCount[threadIdx.x] = bCurCluster ? 1 : 0;
		for (size_t idim = 0; idim < Coords; idim++)
		{
			sm[threadIdx.x + idim * blockDim.x] = bCurCluster ? d_pSrc[idim * SrcCount + tid] : 0.0f;
		}
		//__syncthreads();/*stall barrier 原因*/
		//TODO:reduce
		for (int step = (blockDim.x >> 1); step > 0; step = (step >> 1))
		{
			__syncthreads();
			if (threadIdx.x < step)
			{
				pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + step];
				for (size_t idim = 0; idim < Coords; idim++)
				{
					sm[threadIdx.x + idim * blockDim.x] += sm[threadIdx.x + step + idim * blockDim.x];

				}
			}
		}
		if (0 == threadIdx.x)
		{
			atomicAdd(d_pMemberCount + itcluster, pClusterCount[0]);
			for (size_t idim = 0; idim < Coords; idim++)
			{
				atomicAdd((d_pDst + itcluster * Coords + idim), sm[idim * blockDim.x]);
			}
		}
	}
}

执行情况如下
在这里插入图片描述
耗时850us。
在这里插入图片描述
block128 grid 8192.
在这里插入图片描述
其中性能问题是由block内同步导致的，对应代码位置如下：
在这里插入图片描述

V3 减少block内同步

不同的工作线程处理数据时，由于程逻辑上存在依赖关系，因此block内同步函数__syncthreads()不能省略。但是当归约尺寸小于一个warp的尺寸32时，情况就发生变化了，由于warp SIMT的执行机制，可以消除同步。这里是通过shf指令，通过每个线程的寄存器交换数据。

/*V3 一个warp内归约，通过shf指令交换数据+坐标维度展开*/
template<int COORDS>
__global__ void kKmeansSumCutSync(int SrcCount,
	int Coords/*[1]*/,
	int ClusterCount/*[16]*/,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	extern __shared__ float sm[];/*[Coords][blockDim.x] + [blockDim.x]*/
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int* pClusterCount = (int*)(sm + blockDim.x * COORDS);
	int RealCluster = d_pMembership[tid];
	//TODO:sum each cluster
	for (int itcluster = 0; itcluster < ClusterCount; itcluster++)
	{
		//TODO:load
		bool bCurCluster = (RealCluster == itcluster);
		pClusterCount[threadIdx.x] = bCurCluster ? 1 : 0;
#pragma unroll
		for (size_t idim = 0; idim < COORDS; idim++)
		{
			sm[threadIdx.x + idim * blockDim.x] = bCurCluster ? d_pSrc[idim * SrcCount + tid] : 0.0f;
		}
		//__syncthreads();/*stall barrier 原因*/
		//TODO:reduce
		int step;
		for (step = (blockDim.x >> 1); step > 16; step = (step >> 1))
		{
			__syncthreads();
			if (threadIdx.x < step)
			{
				pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + step];
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					sm[threadIdx.x + idim * blockDim.x] += sm[threadIdx.x + step + idim * blockDim.x];
				}
			}
		}
		//TODO:reduce warp
		if (threadIdx.x < kWarpSize)
		{
			int LocalClusterCount = pClusterCount[threadIdx.x];
			float LocalPoints[kCoords];
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				LocalPoints[i] = sm[threadIdx.x + i * blockDim.x];
			}
			for (; step > 0; step = (step >> 1))
			{
				LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, step, 32);
				LocalPoints[0] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[0], step, 32);
			}
			if (0 == threadIdx.x)
			{
				atomicAdd(d_pMemberCount + itcluster, LocalClusterCount);
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					atomicAdd((d_pDst + itcluster * Coords + idim), LocalPoints[idim]);
				}
			}
		}
	}
}

执行情况如下
在这里插入图片描述
耗时379us，其中内存和计算超过峰值的60%。
在这里插入图片描述

在这里插入图片描述
这里总的stall_barrier相对上一个版本减少，但是执行的总的指令数量也相应减少，因此相对的总体指标反而升高了。

V4 进一步展开

在block大小这里通过编译器展开。

/*V4 blockdim展开*/
template<int COORDS,int BLOCKISZE>
__global__ void kKmeansSumUnroll(int SrcCount,
	int Coords/*[1]*/,
	int ClusterCount/*[16]*/,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	extern __shared__ float sm[];/*[Coords][blockDim.x] + [blockDim.x]*/
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int* pClusterCount = (int*)(sm + BLOCKISZE * COORDS);
	int RealCluster = d_pMembership[tid];
	//TODO:sum each cluster
	for (int itcluster = 0; itcluster < ClusterCount; itcluster++)
	{
		//TODO:load
		bool bCurCluster = (RealCluster == itcluster);
		pClusterCount[threadIdx.x] = bCurCluster ? 1 : 0;
#pragma unroll
		for (size_t idim = 0; idim < COORDS; idim++)
		{
			sm[threadIdx.x + idim * blockDim.x] = bCurCluster ? d_pSrc[idim * SrcCount + tid] : 0.0f;
		}
		//__syncthreads();/*stall barrier 原因*/
		//TODO:reduce
		int step;
#pragma unroll
		for (step = (BLOCKISZE >> 1); step > 16; step = (step >> 1))
		{
			__syncthreads();
			if (threadIdx.x < step)
			{
				pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + step];
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					sm[threadIdx.x + idim * blockDim.x] += sm[threadIdx.x + step + idim * blockDim.x];
				}
			}
		}
		//TODO:reduce warp
		if (threadIdx.x < kWarpSize)
		{
			int LocalClusterCount = pClusterCount[threadIdx.x];
			float LocalPoints[kCoords];
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				LocalPoints[i] = sm[threadIdx.x + i * blockDim.x];
			}
#pragma unroll
			for (; step > 0; step = (step >> 1))
			{
				LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, step, 32);
				LocalPoints[0] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[0], step, 32);
			}
			if (0 == threadIdx.x)
			{
				atomicAdd(d_pMemberCount + itcluster, LocalClusterCount);
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					atomicAdd((d_pDst + itcluster * Coords + idim), LocalPoints[idim]);
				}
			}
		}
	}
}

虽然占用率提升不少，但执行速度并没有快多少363us
在这里插入图片描述

V5 增加负载

基于V5，加载内存部分进行16倍展开

template<int BLOCKSIZE,int UNROLL,int COORDS>
__global__ void kkKmeansUnrollAdv(int SrcCount,
	int Coords,
	int ClusterCount,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	extern __shared__ float sm[];/*[Coords][blockDim.x] + [blockDim.x]*/
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int* pClusterCount = (int*)(sm + blockDim.x * COORDS);
	//int RealCluster = d_pMembership[tid];
	//TODO:sum each cluster
	for (int itcluster = 0; itcluster < ClusterCount; itcluster++)
	{
		//TODO:load
		bool bCurCluster;
		pClusterCount[threadIdx.x] = 0;
		sm[threadIdx.x] = 0.0f;
#pragma unroll
		for (size_t i = 0; i < UNROLL; i++)
		{
			bCurCluster = (d_pMembership[tid + (i * SrcCount / UNROLL)] == itcluster);
			pClusterCount[threadIdx.x] += bCurCluster ? 1 : 0;
			sm[threadIdx.x] += bCurCluster ? d_pSrc[tid + (i * SrcCount / UNROLL)] : 0.0f;
		}
		//__syncthreads();
		TODO:reduce
		//if ((BLOCKSIZE >= 256) && threadIdx.x < 128)
		//{
		//	sm[threadIdx.x] += sm[threadIdx.x + 128];
		//	pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + 128];
		//}
		__syncthreads();
		if ((BLOCKSIZE >= 128) && threadIdx.x < 64)
		{
			sm[threadIdx.x] += sm[threadIdx.x + 64];
			pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + 64];
		}
		__syncthreads();
		if ((BLOCKSIZE >= 64) && threadIdx.x < 32)
		{
			sm[threadIdx.x] += sm[threadIdx.x + 32];
			pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + 32];
		}
		__syncthreads();
		//TODO:reduce warp
		if (threadIdx.x < kWarpSize)
		{
			int LocalClusterCount = pClusterCount[threadIdx.x];
			float LocalPoints[COORDS];
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				LocalPoints[i] = sm[threadIdx.x + i * blockDim.x];
			}
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 16, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 8, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 4, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 2, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 1, 32);
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 16, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 8, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 4, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 2, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 1, 32);
			}
			
			if (0 == threadIdx.x)
			{
				atomicAdd(d_pMemberCount + itcluster, LocalClusterCount);
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					atomicAdd((d_pDst + itcluster * Coords + idim), LocalPoints[idim]);
				}
			}
		}
	}
}

执行情况
在这里插入图片描述
在这里插入图片描述

耗时179us
在这里插入图片描述

block 128 grid 1024
当前的主要性能问题为 Stall MIO Throttle
在这里插入图片描述
代码位置
在这里插入图片描述
这里由于频繁访问共享内存导致性能瓶颈。

V6 寄存器中实现16倍的展开操作

如题，代码实现如下

template<int BLOCKSIZE, int UNROLL, int COORDS>
__global__ void kkKmeansUnrollAdvopt(int SrcCount,
	int Coords,
	int ClusterCount,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	extern __shared__ float sm[];/*[Coords][blockDim.x] + [blockDim.x]*/
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int* pClusterCount = (int*)(sm + blockDim.x * COORDS);
	int LocalCount = 0;
	float LocalClusters[COORDS];

	//int RealCluster = d_pMembership[tid];
	//TODO:sum each cluster
	for (int itcluster = 0; itcluster < ClusterCount; itcluster++)
	{
		//TODO:load
		bool bCurCluster;
#pragma unroll
		for (size_t i = 0; i < COORDS; i++)
		{
			LocalClusters[i] = 0.0f;
		}
		LocalCount = 0;
		pClusterCount[threadIdx.x] = 0;
		sm[threadIdx.x] = 0.0f;
#pragma unroll
		for (size_t i = 0; i < UNROLL; i++)
		{
			bCurCluster = (d_pMembership[tid + (i * SrcCount / UNROLL)] == itcluster);
			LocalCount += bCurCluster ? 1 : 0;
#pragma unroll
			for (size_t idim = 0; idim < COORDS; idim++)
			{
				LocalClusters[idim] += bCurCluster ? d_pSrc[tid + (i * SrcCount / UNROLL)] : 0.0f;
			}
			
		}
		pClusterCount[threadIdx.x] = LocalCount;
#pragma unroll
		for (size_t idim = 0; idim < COORDS; idim++)
		{
			sm[threadIdx.x] = LocalClusters[idim];
		}
		
		__syncthreads();
		if ((BLOCKSIZE >= 128) && threadIdx.x < 64)
		{
			sm[threadIdx.x] += sm[threadIdx.x + 64];
			pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + 64];
		}
		__syncthreads();
		if ((BLOCKSIZE >= 64) && threadIdx.x < 32)
		{
			sm[threadIdx.x] += sm[threadIdx.x + 32];
			pClusterCount[threadIdx.x] += pClusterCount[threadIdx.x + 32];
		}
		__syncthreads();
		//TODO:reduce warp
		if (threadIdx.x < kWarpSize)
		{
			int LocalClusterCount = pClusterCount[threadIdx.x];
			float LocalPoints[COORDS];
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				LocalPoints[i] = sm[threadIdx.x + i * blockDim.x];
			}
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 16, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 8, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 4, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 2, 32);
			LocalClusterCount += __shfl_down_sync(0xFFFFFFFF, LocalClusterCount, 1, 32);
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 16, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 8, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 4, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 2, 32);
				LocalPoints[i] += __shfl_down_sync(0xFFFFFFFF, LocalPoints[i], 1, 32);
			}

			if (0 == threadIdx.x)
			{
				atomicAdd(d_pMemberCount + itcluster, LocalClusterCount);
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					atomicAdd((d_pDst + itcluster * Coords + idim), LocalPoints[idim]);
				}
			}
		}
	}
}

执行情况如下

V7 寄存器版本2

进一步减少共享内存的访问。

/*V6 寄存器进行归约*/
template<int BLOCKSIZE, int UNROLL, int COORDS>
__global__ void kkKmeansReg(int SrcCount,
	int Coords,
	int ClusterCount,
	float* d_pSrc, /*[Coords][SrcCount]*/
	int* d_pMembership,/*[SrcCount]*/
	float* d_pDst,/*[ClusterCount][Coords]*/
	int* d_pMemberCount/*[ClusterCount]*/
)
{
	extern __shared__ float sm[];/*[blockDim.x/kWarpSize] + [blockDim.x/kWarpSize*COORDS]*/
	float* smLCluster = (float*)(sm +( BLOCKSIZE >> 5));
	int* smClusterCount = (int*)sm;
	int LocalClusCount = 0;
	float LocalClus[COORDS];
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	//TODO:sum each cluster
#pragma unroll
	for (int itcluster = 0; itcluster < ClusterCount; itcluster++)
	{
		//TODO:load
		bool bCurCluster;
		LocalClusCount = 0;
#pragma unroll
		for (size_t i = 0; i < COORDS; i++)
		{
			LocalClus[i] = 0.0f;
		}
#pragma unroll
		for (size_t unroll = 0; unroll < UNROLL; unroll++)
		{
			bCurCluster = (d_pMembership[tid + (unroll * SrcCount / UNROLL)] == itcluster);
			LocalClusCount += bCurCluster ? 1 : 0;
#pragma unroll
			for (size_t idim = 0; idim < COORDS; idim++)
			{
				LocalClus[idim] += bCurCluster ? d_pSrc[tid * COORDS + idim + (unroll * SrcCount / UNROLL)] : 0.0f;
			}
		}
		__syncthreads();
		LocalClusCount += __shfl_down_sync(0xFFFFFFFF, LocalClusCount, 16, 32);
		LocalClusCount += __shfl_down_sync(0xFFFFFFFF, LocalClusCount, 8, 32);
		LocalClusCount += __shfl_down_sync(0xFFFFFFFF, LocalClusCount, 4, 32);
		LocalClusCount += __shfl_down_sync(0xFFFFFFFF, LocalClusCount, 2, 32);
		LocalClusCount += __shfl_down_sync(0xFFFFFFFF, LocalClusCount, 1, 32);
#pragma unroll
		for (size_t i = 0; i < COORDS; i++)
		{
			LocalClus[i] += __shfl_down_sync(0xFFFFFFFF, LocalClus[i], 16, 32);
			LocalClus[i] += __shfl_down_sync(0xFFFFFFFF, LocalClus[i], 8, 32);
			LocalClus[i] += __shfl_down_sync(0xFFFFFFFF, LocalClus[i], 4, 32);
			LocalClus[i] += __shfl_down_sync(0xFFFFFFFF, LocalClus[i], 2, 32);
			LocalClus[i] += __shfl_down_sync(0xFFFFFFFF, LocalClus[i], 1, 32);
		}
		__syncthreads();
		if ((threadIdx.x & 0x1F) == 0 )
		{
			smClusterCount[threadIdx.x >> 5] = LocalClusCount;
#pragma unroll
			for (size_t i = 0; i < COORDS; i++)
			{
				smLCluster[(threadIdx.x >> 5) * COORDS + i] = LocalClus[i];
			}
		}
		__syncthreads();
		
		//TODO:reduce warp
		{
			if (0 == threadIdx.x)
			{ 
#pragma unroll
				for (size_t i = 1; i < (BLOCKSIZE>>5); i++)
				{
					LocalClusCount += smClusterCount[i];
#pragma unroll
					for (size_t idim = 0; idim < COORDS; idim++)
					{
						LocalClus[idim] += smLCluster[i * COORDS + idim];
					}
				}
				atomicAdd(d_pMemberCount + itcluster, LocalClusCount);
#pragma unroll
				for (size_t idim = 0; idim < COORDS; idim++)
				{
					atomicAdd((d_pDst + itcluster * Coords + idim), LocalClus[idim]);
				}
			}
		}
	}
}

执行情况
在这里插入图片描述
SM和内存使用率为70%左右
在这里插入图片描述
执行时间下降到86us
在这里插入图片描述
Stall MIO Throttle指标得到改善。

总结

本篇基于CUDA实现了Kmeans算法的归约计算部分，通过6个版本的迭代，执行速度从17200us下降到86us，得到约20倍的加速比。

USST_Coder
关注
  • 0
    点赞
  • 踩
  • 3
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
K-means算法Cuda实现
weixin_40720438的博客
03-03 1695
#include "cuda_runtime.h" #include "device_launch_parameters.h" #include &lt;stdio.h&gt; #include &lt;stdlib.h&gt; #include &lt;iostream&gt; #include &lt;fstream&gt; #include &lt;string&gt; #inclu
代码优化和内存屏障（barrier、mb、rmb、wmb）
09-04 6553
内存屏障主要解决的问题是编译器的优化和CPU的乱序执行。 编译器在优化的时候，生成的汇编指令可能和c语言程序的执行顺序不一样，在需要程序严格按照c语言顺序执行时，需要显式的告诉编译不需要优化，这在linux下是通过barrier()宏完成的，它依靠volatile关键字和memory关键字： /* Optimization barrier */ /* The "volatile...
基于Python——Kmeans聚类算法的实现
01-20
1、概述 本篇博文为数据挖掘算法系列的第一篇。现在对于Kmeans算法进行简单的介绍，Kmeans算法是属于无监督的学习的算法，并且是最基本、最简单的一种基于距离的聚类算法。 下面简单说一下Kmeans算法的步骤： 选随机选取K的簇中心（注意这个K是自己选择的） 计算每个数据点离这K个簇中心的距离，然后将这个点划分到距离最小的簇中 重新计算簇中心，即将每个簇的所有数据点相加求均值，将这个均值作为对应簇的新簇中心。 重复2、3步，直到满足了你设置的停止算法迭代的条件 注意：停止算法迭代的条件一般有三个： 没有（或最小数目）对象被重新分配给不同的聚类。 没有（或最小数目）聚类中心再发生变化。 误差
CUDA性能优化系列——Kmeans算法调优（一）
yan31415的博客
03-09 3257
__global__ void k_kmeans(float* d_src, int srcsize, int dim, float* d_cluster, int clustersize, float* d_dst) { extern __shared__ float sm_cluster[]; float regData[4] = { 0.0f,0.0f, 0.0f, 0.0f };//维度不确定的情况下how to use registers???????? int tid = threadId
【CUDA并行编程之八】Cuda实现Kmeans算法
Marshall的专栏
06-24 1611
本文主要介绍如何使用CUDA并行计算框架编程实现机器学习中的Kmeans算法，Kmeans算法的详细介绍在这里，本文重点在并行实现的过程。 当然还是简单的回顾一下kmeans算法的串行过程： 伪代码： [cpp] view plaincopy 创建k个点作为起始质心(经常是随机选择)   当任意一个点的簇分配结果发生改变时    
k-means聚类torch_gpu版本
helloworld_Fly的博客
12-09 1936
class KMEANS: def __init__(self, n_clusters=10, max_iter=None, verbose=True, device=torch.device("cpu")): # self.n_cluster = n_clusters self.n_clusters = n_clusters self.labels = None self.dists = None # shape: [x.shap
kmeans-master_基于cuda的kmeans算法_cuda加速kmeans_源码
10-01
使用cuda将kmeans算法进行并行化，在linux下运行
实验报告——Kmeans聚类方法.docx
01-22
实验报告——Kmeans聚类方法.docx
粒子群算法、粒子群优化kmeans、fcm
06-29
智能算法对传统聚类算法的优化，亲测可用
PostgreSQL数据聚类——kmeans
01-20
在数据挖掘中， k-Means 算法是一种 cluster analysis 的算法，其主要是来计算数据聚集的算法，主要通过不断地取离种子点最近均值的算法。 介绍K-Means算法前，我们需要先弄清楚分类和聚类这两个概念： 分类：类别是...
K-means CUDA实现
05-13
CUDA实现K-means算法，学习GPU编程入门级算法需要了解
kmcuda：NVIDIA GPU CUDA上的大规模K-means和K-nn实施
02-03
使用NVIDIA CUDA的“阴阳” K均值和K-nn K均值的实现基于 。 尽管它引入了一些开销和许多对CUDA不利的条件子句，但与Lloyd算法相比，它仍然显示出1.6到2倍的加速。 K近邻采用相同的三角形不等式思想，并且需要预先计算的质心和聚类分配，类似于展平的球树。 sklearn KMeans KMeansRex KMeansRex OpenMP 塞班 克库达 kmcuda 2 GPU 速度 1倍 4.5倍 8.2倍 15.5倍 17.8倍 29.8倍 记忆 1倍 2倍 2倍 0.6倍 0.6倍 0.6倍 从技术上讲，该项目是一个共享库，可导出kmcuda.h定义的两个函数： kmeans_cuda和knn_cuda 。 它具有内置的Python3和R本机扩展支持，因此您可以from libKMCUDA import kmeans_cuda或dyn.load("libKMCUDA.so") 。 目录 K均值 该项目与其他项目之间的主要区别在于，kmcuda已针对内存消耗低和大量群集进行了优化。 例如，kmcuda可以将480个维度中的4M个样本分类为40000个
cuda-kmeans:K-means 算法在图像上的并行实现
06-06
Cuda K-Means 图像聚类算法由 Andrea Toscano, Università degli Studi di Milano (Informatica) 在 NVidia Cuda 中实现。 这个小项目展示了如何实现应用于图像的 K-Means 聚类算法以减少其颜色。 一些预处理是使用 python 脚本计算的，以便以更好的方式表示图像并轻松找到适合算法的初始质心。 在 Cuda K-Means 例程中涉及全局内存和常量内存。 未来的工作还将包括包含整个图片的纹理内存，从而提高算法的性能。
两种方法利用CUDA实现矩阵乘法
最新发布
Goodness2020的博客
06-01 497
CUDA编程
基于GPU的K-Means聚类算法
weixin_34097242的博客
12-16 393
基于GPU的K-Means聚类算法 　　聚类是信息检索、数据挖掘中的一类重要技术，是分析数据并从中发现有用信息的一种有效手段。它将数据对象分组成为多个类或簇，使得在同一个簇中的对象之间具有较高的相似度，而不同簇中的对象差别很大。作为统计学的一个分支和一种无监督的学习方法，聚类从数学分析的角度提供了一种准确、细致的分析工具。而k-means算法是...
内存屏障（Memory Barrier）究竟是个什么鬼？
黑帽子技术的博客
03-16 1756
3 月，跳不动了？>>> 读者朋友你好hello ...
CUDA线程屏障(thread barrier)解析
xfjlionking的博客
09-10 1361
CUDA线程屏障(thread barrier)解析 CUDA Runtime API中提供了一个__syncthreads()接口用来同步线程块中的线程，此接口在官方文档中说得比较笼统，本文通过实际测试对该接口使用方法、使用陷阱做了探索。 新的改变 我们对Markdown编辑器进行了一些功能拓展与语法支持，除了标准...
CUDA编程：矩阵乘运算从CPU到GPU
深度学习技术前沿
11-02 942
作者丨kaiyuan@知乎来源丨https://zhuanlan.zhihu.com/p/573271688本文主要介绍用CUDA实现矩阵乘法运算（C = A x B）的几个基本方法，帮助大家理解矩阵在GPU上面的运算与CPU上的有何异同，通过实践上手CUDA的优化计算，相比基础方法，能提速10倍以上。本文内容涉及到CUDA矩阵1D运算、2D运算、共享内存、CUBLAS的使用。文中的全部code：...
Python——Kmeans聚类算法
04-29
Kmeans聚类算法是一种无监督学习算法，用于将数据集划分为不同的簇。它是一个迭代算法，通过计算每个数据点与簇中心的距离，将数据点分配到最近的簇中心。然后，根据分配的数据点更新簇中心。重复这个过程，直到簇...

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值