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

最新推荐文章于 2021-09-24 16:30:00 发布
最新推荐文章于 2021-09-24 16:30:00 发布
阅读量1.8k 收藏 2
点赞数 1
分类专栏： # CUDA性能调优 CUDA 文章标签： cuda gpgpu
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/yan31415/article/details/115486722
版权

本篇对调度方式进行优化，实现内存拷贝和计算overlap。

单流同步调用

/*
单流同步
*/
void CallKmeansSync()
{
	//TODO:init host memory
	float* h_Src/*[Coords][SrcCount]*/,
		* h_Clusters/*[ClusterCount][Coords]*/;
	int* h_MemberShip/*[kSrcCount]*/,
		* h_MemberCount/*[kClusterCount]*/;
	h_Src = (float*)malloc(kSrcCount * kCoords * sizeof(float));
	h_Clusters = (float*)malloc(kClusterCount * kCoords * sizeof(float));/*[kClusterCount][kCoords]*/
	for (size_t i = 0; i < kClusterCount; i++)
	{
		for (int j = 0; j < kCoords; ++j)
		{
			h_Clusters[i * kCoords + j] = float(100 * i + 10);
		}
	}
	for (size_t i = 0; i < kClusterCount; i++)
	{
		for (size_t j = 0; j < (kSrcCount / kClusterCount); j++)
		{
			for (size_t iDim = 0; iDim < kCoords; iDim++)
			{
				h_Src[iDim * kSrcCount + i * (kSrcCount / kClusterCount) + j] = i * 100 + 0.5f;// -(float)((rand() % 100) / 100);
			}
		}
	}
	h_MemberShip = (int*)malloc(kSrcCount * sizeof(int));
	memset(h_MemberShip, 9, kSrcCount * sizeof(int));
	h_MemberCount = (int*)malloc(kClusterCount * sizeof(int));
	memset(h_MemberCount, 0, kClusterCount * sizeof(int));

	//TODO:init device mempry
	float* d_pSrc/*[Coords][SrcCount]*/,
		* d_pClusters/*[ClusterCount][Coords]*/;
	int* d_pChanged/*[1]*/,
		* d_pMemberShip/*[kSrcCount]*/,
		* d_pMemberCount/*[kClusterCount]*/;
	cudaMalloc(&d_pSrc, kSrcCount * kCoords * sizeof(float));
	cudaMemcpy(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);

	cudaMalloc(&d_pClusters, kClusterCount * kCoords * sizeof(float));


	cudaMalloc(&d_pMemberShip, kSrcCount * sizeof(int));
	cudaMemcpy(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice);

	cudaMalloc(&d_pMemberCount, kClusterCount * sizeof(int));
	cudaMemcpy(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice);

	cudaMalloc(&d_pChanged, sizeof(int));

	//TODO:find the points
	int itCount = 0;
	int iChanged = 0;
	cudaStream_t sMember, sNewCluster;
	cudaStreamCreate(&sMember);
	cudaStreamCreate(&sNewCluster);
	cudaEvent_t eMember, eNewCluster;
	cudaEventCreate(&eMember);
	cudaEventCreate(&eNewCluster);
	do
	{
		{
			const int UnrollScale = 8;
			cudaMemcpy(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);//s1
			cudaMemset(d_pMemberCount, 0, kClusterCount * sizeof(int));//s1
			kKmeansClassifyMembershipMath<kCoords, UnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / UnrollScale, kClusterCount* kCoords * sizeof(float) >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);
			checkCUDARtn(cudaDeviceSynchronize());
			cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);//s1
			cudaMemset(d_pChanged, 0, sizeof(int));//s1
			cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));//s2
			const int kUnrollAdvUNROLLSCALE = 16;
			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvUNROLLSCALE, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvUNROLLSCALE, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
			checkCUDARtn(cudaDeviceSynchronize());
			cudaMemcpy(h_MemberCount, d_pMemberCount, kClusterCount * sizeof(int), cudaMemcpyDeviceToHost);//s2
			cudaMemcpy(h_Clusters, d_pClusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyDeviceToHost);//s2
			for (size_t i = 0; i < kClusterCount; i++)//s2
			{
				for (int j = 0; j < kCoords; ++j)
				{
					h_Clusters[i * kCoords + j] = (h_MemberCount[i] == 0) ? 0 : (h_Clusters[i * kCoords + j] / h_MemberCount[i]);
				}
			}
		}
	} while (/*(0 != iChanged) &&*/ itCount++ < 8);
	std::cout << "it count " << itCount << std::endl;
	cudaEventDestroy(eMember);
	cudaEventDestroy(eNewCluster);
	cudaStreamDestroy(sMember);
	cudaStreamDestroy(sNewCluster);
	free(h_Src);
	free(h_Clusters);
	free(h_MemberShip);
	free(h_MemberCount);
	cudaFree(d_pSrc);
	cudaFree(d_pClusters);
	cudaFree(d_pChanged);
	cudaFree(d_pMemberShip);
	cudaFree(d_pMemberCount);
}

void CallFunckmeans(int FuncIndex)
{
	const int kSingleStm = 0;
	const int kMultStm = 1;
	const int kPrevCopy = 2;
	const int kNonSync = 3;
	const int kSync = 4;
	if (FuncIndex == kSingleStm)
	{
		CallKmeansSingleStm();
	}
	else if (FuncIndex == kMultStm)
	{
		CallKmeansMultOverlap();
	}
	else if (FuncIndex == kPrevCopy)
	{
		CallKmeansMultPrevCopy();
	}
	else if (FuncIndex == kNonSync)
	{
		CallKmeansMultNonSync();
	}
	else if (FuncIndex == kSync)
	{
		CallKmeansSync();
	}
	return;
}

调度情况
在这里插入图片描述
总耗时7ms，可以看到核函数之间由于同步没有连续执行

单流异步调用

在这里插入代码片/*
单流异步
*/
void CallKmeansSingleStm()
{

	//TODO:init host memory
	float* h_Src/*[Coords][SrcCount]*/,
		* h_Clusters/*[ClusterCount][Coords]*/;
	int* h_MemberShip/*[kSrcCount]*/,
		* h_MemberCount/*[kClusterCount]*/;
	//h_Src = (float*)malloc(kSrcCount * kCoords * sizeof(float));
	cudaMallocHost(&h_Src, kSrcCount * kCoords * sizeof(float));
	h_Clusters = (float*)malloc(kClusterCount * kCoords * sizeof(float));/*[kClusterCount][kCoords]*/
	for (size_t i = 0; i < kClusterCount; i++)
	{
		for (int j = 0; j < kCoords; ++j)
		{
			h_Clusters[i * kCoords + j] = float(100 * i + 10);
		}
	}
	for (size_t i = 0; i < kClusterCount; i++)
	{
		for (size_t j = 0; j < (kSrcCount / kClusterCount); j++)
		{
			for (size_t iDim = 0; iDim < kCoords; iDim++)
			{
				h_Src[iDim * kSrcCount + i * (kSrcCount / kClusterCount) + j] = i * 100 + 0.5f;// -(float)((rand() % 100) / 100);
			}
		}
	}
	//h_MemberShip = (int*)malloc(kSrcCount * sizeof(int));
	cudaMallocHost(&h_MemberShip, kSrcCount * sizeof(int));
	memset(h_MemberShip, 9, kSrcCount * sizeof(int));
	h_MemberCount = (int*)malloc(kClusterCount * sizeof(int));
	memset(h_MemberCount, 0, kClusterCount * sizeof(int));

	//TODO:init stream
	cudaStream_t stm;
	cudaStreamCreate(&stm);
	const int EventNum = 10;
	cudaEvent_t event[EventNum];
	for (size_t i = 0; i < EventNum; i++)
	{
		cudaEventCreate(&event[i]);
	}

	//TODO:init device mempry
	float* d_pSrc/*[Coords][SrcCount]*/,
		* d_pClusters/*[ClusterCount][Coords]*/;
	int* d_pChanged/*[1]*/,
		* d_pMemberShip/*[kSrcCount]*/,
		* d_pMemberCount/*[kClusterCount]*/;
	cudaMalloc(&d_pSrc, kSrcCount * kCoords * sizeof(float));
	cudaMalloc(&d_pClusters, kClusterCount * kCoords * sizeof(float));
	cudaMalloc(&d_pMemberShip, kSrcCount * sizeof(int));
	cudaMalloc(&d_pMemberCount, kClusterCount * sizeof(int));
	cudaMalloc(&d_pChanged, sizeof(int));

	//cudaMemcpy(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpyAsync(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm);
	//cudaMemcpy(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpyAsync(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice, stm);
	//cudaMemcpy(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpyAsync(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice, stm);
	//cudaMemcpy(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);//默认流导致和stm流之间同步的时间消耗几百微秒
	cudaMemcpyAsync(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm);
	//TODO:find the points
	int itCount = 0;
	int* changed;
	cudaMallocHost(&changed, sizeof(int));
	*changed = 0;
	int iChanged = 0;
	const int kUnrollAdvExecuteScale = 8;
	const int kUnrollAdvkUnrollScale = 16;
	const int preCalcCount = 10;
	for (; itCount < preCalcCount; itCount++)
	{
		kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float) >> >
			(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);
		//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
		cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm);
		//cudaMemset(d_pChanged, 0, sizeof(int));
		cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm);
		cudaEventRecord(event[itCount % 10], stm);
		//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
		cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm);
		kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
			//kKmeansSumAtomic << <kGridCount* 1, kBlockSize / 1, (kBlockSize* kCoords / 1 + kClusterCount * kCoords + kClusterCount) * sizeof(float) >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);
		kCalcNewClusterCenter << <1, kClusterCount, 0, stm >> > (d_pClusters, d_pMemberCount);
		cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm);
	}
	cudaEventSynchronize(event[preCalcCount - 5]);
	while ((0 != *changed) && itCount++ < 100)
	{
		kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float) >> >
			(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);
		//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
		cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm);
		cudaEventRecord(event[itCount % 10], stm);
		//cudaMemset(d_pChanged, 0, sizeof(int));
		cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm);
		//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
		cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm);
		kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount* kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize* kCoords + kBlockSize) * sizeof(float) / kExecuteScale >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
		kCalcNewClusterCenter << <1, kClusterCount, 0, stm >> > (d_pClusters, d_pMemberCount);
		cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm);
	}
	std::cout << "it count " << itCount << std::endl;
	cudaMemcpy(h_Clusters, d_pClusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyDeviceToHost);
	for (size_t i = 0; i < kClusterCount; i++)
	{
		std::cout << h_Clusters[i] << std::endl;
	}
	//free(h_Src);
	cudaFreeHost(h_Src);
	free(h_Clusters);
	//free(h_MemberShip);
	cudaFreeHost(h_MemberShip);
	cudaFreeHost(changed);
	free(h_MemberCount);
	cudaFree(d_pSrc);
	cudaFree(d_pClusters);
	cudaFree(d_pChanged);
	cudaFree(d_pMemberShip);
	cudaFree(d_pMemberCount);
	cudaStreamDestroy(stm);
	for (size_t i = 0; i < EventNum; i++)
	{
		cudaEventDestroy(event[i]);
	}
}

调度情况
在这里插入图片描述
执行时间5.6ms，由上图可以该方式可以使核函数连续执行，但是由于主机端调用事件同步函数，导致流中核函数执行发生1.7ms的间断。

多流执行

多流执行，通过OpenMP开启4个线程，每个线程中一个流。同时在最开始的内存拷贝阶段通过流中添加事件同步，保证依赖关系。

/*
4个线程各自持有一个流
流中最开始的内存拷贝之间通过事件同步添加依赖
*/
void CallKmeansMultOverlap()
{
	const int ThreadNum = 4;
	//TODO:init host memory
	float* h_SrcBase/*[ThreadNum][Coords][SrcCount]*/,
		* h_ClustersBase/*[ThreadNum][ClusterCount][Coords]*/;
	int* h_MemberShipBase/*[ThreadNum][kSrcCount]*/,
		* h_MemberCountBase/*[ThreadNum][kClusterCount]*/;
	//h_Src = (float*)malloc(kSrcCount * kCoords * sizeof(float));
	cudaMallocHost(&h_SrcBase, ThreadNum * kSrcCount * kCoords * sizeof(float));
	h_ClustersBase = (float*)malloc(ThreadNum * kClusterCount * kCoords * sizeof(float));
	int* changedBase;
	cudaMallocHost(&changedBase, ThreadNum * sizeof(int));
	for (size_t tid = 0; tid < ThreadNum; tid++)
	{
		float* h_LocalClusters = h_ClustersBase + kClusterCount * kCoords * tid;
		for (size_t i = 0; i < kClusterCount; i++)
		{
			for (int j = 0; j < kCoords; ++j)
			{
				h_LocalClusters[i * kCoords + j] = float(100 * i + 10);
			}
		}
		float* h_LocalSrc = h_SrcBase + kSrcCount * kCoords * tid;
		for (size_t i = 0; i < kClusterCount; i++)
		{
			for (size_t j = 0; j < (kSrcCount / kClusterCount); j++)
			{
				for (size_t iDim = 0; iDim < kCoords; iDim++)
				{
					h_LocalSrc[iDim * kSrcCount + i * (kSrcCount / kClusterCount) + j] = i * 100 + 0.5f;// -(float)((rand() % 100) / 100);
				}
			}
		}
		changedBase[tid] = 0;
	}

	//h_MemberShip = (int*)malloc(kSrcCount * sizeof(int));
	cudaMallocHost(&h_MemberShipBase, ThreadNum * kSrcCount * sizeof(int));
	memset(h_MemberShipBase, 9, ThreadNum * kSrcCount * sizeof(int));
	h_MemberCountBase = (int*)malloc(ThreadNum * kClusterCount * sizeof(int));
	memset(h_MemberCountBase, 0, ThreadNum * kClusterCount * sizeof(int));

	//TODO:init stream
	cudaStream_t stm[ThreadNum];
	for (size_t i = 0; i < ThreadNum; i++)
	{
		//cudaStreamCreate(&stm[i]);
		cudaStreamCreateWithFlags(&stm[i], cudaStreamNonBlocking);
	}
	const int StmEventCount = 10;
	cudaEvent_t event[StmEventCount * ThreadNum];
	for (size_t i = 0; i < StmEventCount * ThreadNum; i++)
	{
		cudaEventCreate(&event[i]);
	}

	//TODO:init device mempry
	float* d_pSrcBase/*[ThreadNum][Coords][SrcCount]*/,
		* d_pClustersBase/*[ThreadNum][ClusterCount][Coords]*/;
	int* d_pChangedBase/*[ThreadNum][1]*/,
		* d_pMemberShipBase/*[ThreadNum][kSrcCount]*/,
		* d_pMemberCountBase/*[ThreadNum][kClusterCount]*/;
	cudaMalloc(&d_pSrcBase, ThreadNum * kSrcCount * kCoords * sizeof(float));
	cudaMalloc(&d_pClustersBase, ThreadNum * kClusterCount * kCoords * sizeof(float));
	cudaMalloc(&d_pMemberShipBase, ThreadNum * kSrcCount * sizeof(int));
	cudaMalloc(&d_pMemberCountBase, ThreadNum * kClusterCount * sizeof(int));
	cudaMalloc(&d_pChangedBase, ThreadNum * sizeof(int));


	const int kUnrollAdvExecuteScale = 8;
	const int kUnrollAdvkUnrollScale = 16;
	const int preCalcCount = 8;
#pragma omp parallel for num_threads(ThreadNum)
	for (size_t itstm = 0; itstm < ThreadNum; itstm++)
	{
		float* h_Src = h_SrcBase + itstm * kSrcCount * kCoords;
		float* d_pSrc = d_pSrcBase + itstm * kSrcCount * kCoords;
		int* h_MemberShip = h_MemberShipBase + itstm * kSrcCount;
		int* d_pMemberShip = d_pMemberShipBase + itstm * kSrcCount;
		int* h_MemberCount = h_MemberCountBase + itstm * kClusterCount;
		int* d_pMemberCount = d_pMemberCountBase + itstm * kClusterCount;
		float* h_Clusters = h_ClustersBase + itstm * kClusterCount * kCoords;
		float* d_pClusters = d_pClustersBase + itstm * kClusterCount * kCoords;
		int* changed = changedBase + itstm;
		int* d_pChanged = d_pChangedBase + itstm;
		if (itstm != 0)
		{
			cudaStreamWaitEvent(stm[itstm], event[(itstm - 1) * StmEventCount], 0);
		}
		//cudaMemcpy(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm[itstm]);
		//cudaMemcpy(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice, stm[itstm]);
		//cudaMemcpy(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice, stm[itstm]);
		//cudaMemcpy(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);//默认流导致和stm流之间同步的时间消耗几百微秒
		cudaMemcpyAsync(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm[itstm]);
		cudaEventRecord(event[itstm * StmEventCount], stm[itstm]);
		//TODO:find the points
		int itCount = 1;


		for (; itCount < preCalcCount; itCount++)
		{
			kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float), stm[itstm] >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);

			//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm[itstm]);
			//cudaMemset(d_pChanged, 0, sizeof(int));
			cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm[itstm]);

			cudaEventRecord(event[(itCount % 10) + itstm * StmEventCount], stm[itstm]);

			//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
			cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm[itstm]);

			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale, stm[itstm] >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
				//kKmeansSumAtomic << <kGridCount* 1, kBlockSize / 1, (kBlockSize* kCoords / 1 + kClusterCount * kCoords + kClusterCount) * sizeof(float) >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);
			kCalcNewClusterCenter << <1, kClusterCount, 0, stm[itstm] >> > (d_pClusters, d_pMemberCount);

			cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm[itstm]);

		}
		cudaEventSynchronize(event[(itCount % 10) + itstm * StmEventCount - 4]);
		while ((0 != *changed) && itCount++ < 100)
		{
			kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float), stm[itstm] >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);
			//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm[itstm]);
			cudaEventRecord(event[itCount % 10], stm[itstm]);
			//cudaMemset(d_pChanged, 0, sizeof(int));
			cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm[itstm]);
			//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
			cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm[itstm]);
			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale, stm[itstm] >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
			kCalcNewClusterCenter << <1, kClusterCount, 0, stm[itstm] >> > (d_pClusters, d_pMemberCount);
			cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm[itstm]);
		}
		printf("tid %d and count %d \n", itstm, itCount);
		cudaMemcpyAsync(h_Clusters, d_pClusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyDeviceToHost, stm[itstm]);
		for (size_t i = 0; i < kClusterCount; i++)
		{
			printf("tid %d and value %f \n", itstm, h_Clusters[i]);
		}
	}


	//free(h_Src);
	cudaFreeHost(h_SrcBase);
	free(h_ClustersBase);
	//free(h_MemberShip);
	cudaFreeHost(h_MemberShipBase);
	cudaFreeHost(changedBase);

	free(h_MemberCountBase);
	cudaFree(d_pSrcBase);
	cudaFree(d_pClustersBase);
	cudaFree(d_pChangedBase);
	cudaFree(d_pMemberShipBase);
	cudaFree(d_pMemberCountBase);

	for (size_t i = 0; i < ThreadNum; i++)
	{
		cudaStreamDestroy(stm[i]);
	}

	for (size_t i = 0; i < StmEventCount * ThreadNum; i++)
	{
		cudaEventDestroy(event[i]);
	}

}

执行情况
在这里插入图片描述
总执行时间8.6ms，这种多流的调度方式，由主机到设备端的内存拷贝没有连续执行。同时第二次内存拷贝开始，核函数的执行和内存拷贝完成overlap。

提前执行内存拷贝

/*
提前进行异步内存拷贝
4个线程，每个执行各自的流通过事件同步触发核函数启动
*/
void CallKmeansMultPrevCopy()
{
	const int ThreadNum = 4;
	//TODO:init host memory
	float* h_SrcBase/*[ThreadNum][Coords][SrcCount]*/,
		* h_ClustersBase/*[ThreadNum][ClusterCount][Coords]*/;
	int* h_MemberShipBase/*[ThreadNum][kSrcCount]*/,
		* h_MemberCountBase/*[ThreadNum][kClusterCount]*/;
	//h_Src = (float*)malloc(kSrcCount * kCoords * sizeof(float));
	cudaMallocHost(&h_SrcBase, ThreadNum * kSrcCount * kCoords * sizeof(float));
	h_ClustersBase = (float*)malloc(ThreadNum * kClusterCount * kCoords * sizeof(float));
	int* changedBase;
	cudaMallocHost(&changedBase, ThreadNum * sizeof(int));
	for (size_t tid = 0; tid < ThreadNum; tid++)
	{
		float* h_LocalClusters = h_ClustersBase + kClusterCount * kCoords * tid;
		for (size_t i = 0; i < kClusterCount; i++)
		{
			for (int j = 0; j < kCoords; ++j)
			{
				h_LocalClusters[i * kCoords + j] = float(100 * i + 10);
			}
		}
		float* h_LocalSrc = h_SrcBase + kSrcCount * kCoords * tid;
		for (size_t i = 0; i < kClusterCount; i++)
		{
			for (size_t j = 0; j < (kSrcCount / kClusterCount); j++)
			{
				for (size_t iDim = 0; iDim < kCoords; iDim++)
				{
					h_LocalSrc[iDim * kSrcCount + i * (kSrcCount / kClusterCount) + j] = i * 100 + 0.5f;// -(float)((rand() % 100) / 100);
				}
			}
		}
		changedBase[tid] = 0;
	}

	//h_MemberShip = (int*)malloc(kSrcCount * sizeof(int));
	cudaMallocHost(&h_MemberShipBase, ThreadNum * kSrcCount * sizeof(int));
	memset(h_MemberShipBase, 9, ThreadNum * kSrcCount * sizeof(int));
	h_MemberCountBase = (int*)malloc(ThreadNum * kClusterCount * sizeof(int));
	memset(h_MemberCountBase, 0, ThreadNum * kClusterCount * sizeof(int));

	//TODO:init stream
	cudaStream_t stm[ThreadNum + 1];
	for (size_t i = 0; i < ThreadNum + 1; i++)
	{
		//cudaStreamCreate(&stm[i]);
		cudaStreamCreateWithFlags(&stm[i], cudaStreamNonBlocking);
	}
	const int StmEventCount = 10;
	cudaEvent_t event[(1 + StmEventCount) * ThreadNum];
	for (size_t i = 0; i < (1 + StmEventCount) * ThreadNum; i++)
	{
		cudaEventCreate(&event[i]);
	}

	//TODO:init device mempry
	float* d_pSrcBase/*[ThreadNum][Coords][SrcCount]*/,
		* d_pClustersBase/*[ThreadNum][ClusterCount][Coords]*/;
	int* d_pChangedBase/*[ThreadNum][1]*/,
		* d_pMemberShipBase/*[ThreadNum][kSrcCount]*/,
		* d_pMemberCountBase/*[ThreadNum][kClusterCount]*/;
	cudaMalloc(&d_pSrcBase, ThreadNum * kSrcCount * kCoords * sizeof(float));
	cudaMalloc(&d_pClustersBase, ThreadNum * kClusterCount * kCoords * sizeof(float));
	cudaMalloc(&d_pMemberShipBase, ThreadNum * kSrcCount * sizeof(int));
	cudaMalloc(&d_pMemberCountBase, ThreadNum * kClusterCount * sizeof(int));
	cudaMalloc(&d_pChangedBase, ThreadNum * sizeof(int));


	const int kUnrollAdvExecuteScale = 8;
	const int kUnrollAdvkUnrollScale = 16;
	const int preCalcCount = 8;

	for (size_t itstm = 0; itstm < ThreadNum; itstm++)
	{
		float* h_Src = h_SrcBase + itstm * kSrcCount * kCoords;
		float* d_pSrc = d_pSrcBase + itstm * kSrcCount * kCoords;
		int* h_MemberShip = h_MemberShipBase + itstm * kSrcCount;
		int* d_pMemberShip = d_pMemberShipBase + itstm * kSrcCount;
		int* h_MemberCount = h_MemberCountBase + itstm * kClusterCount;
		int* d_pMemberCount = d_pMemberCountBase + itstm * kClusterCount;
		float* h_Clusters = h_ClustersBase + itstm * kClusterCount * kCoords;
		float* d_pClusters = d_pClustersBase + itstm * kClusterCount * kCoords;
		int* changed = changedBase + itstm;
		int* d_pChanged = d_pChangedBase + itstm;
		//cudaMemcpy(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm[ThreadNum]);
		//cudaMemcpy(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice, stm[ThreadNum]);
		//cudaMemcpy(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice, stm[ThreadNum]);
		//cudaMemcpy(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);//默认流导致和stm流之间同步的时间消耗几百微秒
		cudaMemcpyAsync(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm[ThreadNum]);
		cudaEventRecord(event[ThreadNum * StmEventCount + itstm], stm[ThreadNum]);
	}
#pragma omp parallel for num_threads(ThreadNum)
	for (size_t itstm = 0; itstm < ThreadNum; itstm++)
	{
		float* h_Src = h_SrcBase + itstm * kSrcCount * kCoords;
		float* d_pSrc = d_pSrcBase + itstm * kSrcCount * kCoords;
		int* h_MemberShip = h_MemberShipBase + itstm * kSrcCount;
		int* d_pMemberShip = d_pMemberShipBase + itstm * kSrcCount;
		int* h_MemberCount = h_MemberCountBase + itstm * kClusterCount;
		int* d_pMemberCount = d_pMemberCountBase + itstm * kClusterCount;
		float* h_Clusters = h_ClustersBase + itstm * kClusterCount * kCoords;
		float* d_pClusters = d_pClustersBase + itstm * kClusterCount * kCoords;
		int* changed = changedBase + itstm;
		int* d_pChanged = d_pChangedBase + itstm;
		//TODO:find the points
		cudaStreamWaitEvent(stm[itstm], event[ThreadNum * StmEventCount + itstm], 0);
		int itCount = 1;


		for (; itCount < preCalcCount; itCount++)
		{
			kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float), stm[itstm] >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);

			//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm[itstm]);
			//cudaMemset(d_pChanged, 0, sizeof(int));
			cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm[itstm]);

			cudaEventRecord(event[(itCount % 10) + itstm * StmEventCount], stm[itstm]);

			//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
			cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm[itstm]);

			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale, stm[itstm] >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
				//kKmeansSumAtomic << <kGridCount* 1, kBlockSize / 1, (kBlockSize* kCoords / 1 + kClusterCount * kCoords + kClusterCount) * sizeof(float) >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);
			kCalcNewClusterCenter << <1, kClusterCount, 0, stm[itstm] >> > (d_pClusters, d_pMemberCount);

			cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm[itstm]);

		}
		cudaEventSynchronize(event[(itCount % 10) + itstm * StmEventCount - 4]);
		while ((0 != *changed) && itCount++ < 100)
		{
			kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float), stm[itstm] >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);
			//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm[itstm]);
			cudaEventRecord(event[itCount % 10], stm[itstm]);
			//cudaMemset(d_pChanged, 0, sizeof(int));
			cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm[itstm]);
			//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
			cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm[itstm]);
			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale, stm[itstm] >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
			kCalcNewClusterCenter << <1, kClusterCount, 0, stm[itstm] >> > (d_pClusters, d_pMemberCount);
			cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm[itstm]);
		}
		printf("tid %d and count %d \n", itstm, itCount);
		cudaMemcpyAsync(h_Clusters, d_pClusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyDeviceToHost, stm[itstm]);
		for (size_t i = 0; i < kClusterCount; i++)
		{
			printf("tid %d and value %f \n", itstm, h_Clusters[i]);
		}
	}


	//free(h_Src);
	cudaFreeHost(h_SrcBase);
	free(h_ClustersBase);
	//free(h_MemberShip);
	cudaFreeHost(h_MemberShipBase);
	cudaFreeHost(changedBase);

	free(h_MemberCountBase);
	cudaFree(d_pSrcBase);
	cudaFree(d_pClustersBase);
	cudaFree(d_pChangedBase);
	cudaFree(d_pMemberShipBase);
	cudaFree(d_pMemberCountBase);

	for (size_t i = 0; i < ThreadNum + 1; i++)
	{
		cudaStreamDestroy(stm[i]);
	}

	for (size_t i = 0; i < (1 + StmEventCount) * ThreadNum; i++)
	{
		cudaEventDestroy(event[i]);
	}

}

调度情况
在这里插入图片描述

执行时间9.4ms，这种方式内存拷贝连续执行，但是事件同步触发流中核函数执行，并没有很好的overlap。

多流不添加事件同步

/*
4个线程，每个执行各自的流
不同流之间拷贝数据未添加同步操作
*/
void CallKmeansMultNonSync()
{
	const int ThreadNum = 4;
	//TODO:init host memory
	float* h_SrcBase/*[ThreadNum][Coords][SrcCount]*/,
		* h_ClustersBase/*[ThreadNum][ClusterCount][Coords]*/;
	int* h_MemberShipBase/*[ThreadNum][kSrcCount]*/,
		* h_MemberCountBase/*[ThreadNum][kClusterCount]*/;
	//h_Src = (float*)malloc(kSrcCount * kCoords * sizeof(float));
	cudaMallocHost(&h_SrcBase, ThreadNum * kSrcCount * kCoords * sizeof(float));
	h_ClustersBase = (float*)malloc(ThreadNum * kClusterCount * kCoords * sizeof(float));
	int* changedBase;
	cudaMallocHost(&changedBase, ThreadNum * sizeof(int));
	for (size_t tid = 0; tid < ThreadNum; tid++)
	{
		float* h_LocalClusters = h_ClustersBase + kClusterCount * kCoords * tid;
		for (size_t i = 0; i < kClusterCount; i++)
		{
			for (int j = 0; j < kCoords; ++j)
			{
				h_LocalClusters[i * kCoords + j] = float(100 * i + 10);
			}
		}
		float* h_LocalSrc = h_SrcBase + kSrcCount * kCoords * tid;
		for (size_t i = 0; i < kClusterCount; i++)
		{
			for (size_t j = 0; j < (kSrcCount / kClusterCount); j++)
			{
				for (size_t iDim = 0; iDim < kCoords; iDim++)
				{
					h_LocalSrc[iDim * kSrcCount + i * (kSrcCount / kClusterCount) + j] = i * 100 + 0.5f;// -(float)((rand() % 100) / 100);
				}
			}
		}
		changedBase[tid] = 0;
	}

	//h_MemberShip = (int*)malloc(kSrcCount * sizeof(int));
	cudaMallocHost(&h_MemberShipBase, ThreadNum * kSrcCount * sizeof(int));
	memset(h_MemberShipBase, 9, ThreadNum * kSrcCount * sizeof(int));
	h_MemberCountBase = (int*)malloc(ThreadNum * kClusterCount * sizeof(int));
	memset(h_MemberCountBase, 0, ThreadNum * kClusterCount * sizeof(int));

	//TODO:init stream
	cudaStream_t stm[ThreadNum + 1];
	for (size_t i = 0; i < ThreadNum + 1; i++)
	{
		//cudaStreamCreate(&stm[i]);
		cudaStreamCreateWithFlags(&stm[i], cudaStreamNonBlocking);
	}
	const int StmEventCount = 10;
	cudaEvent_t event[(1 + StmEventCount) * ThreadNum];
	for (size_t i = 0; i < (1 + StmEventCount) * ThreadNum; i++)
	{
		cudaEventCreate(&event[i]);
	}

	//TODO:init device mempry
	float* d_pSrcBase/*[ThreadNum][Coords][SrcCount]*/,
		* d_pClustersBase/*[ThreadNum][ClusterCount][Coords]*/;
	int* d_pChangedBase/*[ThreadNum][1]*/,
		* d_pMemberShipBase/*[ThreadNum][kSrcCount]*/,
		* d_pMemberCountBase/*[ThreadNum][kClusterCount]*/;
	cudaMalloc(&d_pSrcBase, ThreadNum * kSrcCount * kCoords * sizeof(float));
	cudaMalloc(&d_pClustersBase, ThreadNum * kClusterCount * kCoords * sizeof(float));
	cudaMalloc(&d_pMemberShipBase, ThreadNum * kSrcCount * sizeof(int));
	cudaMalloc(&d_pMemberCountBase, ThreadNum * kClusterCount * sizeof(int));
	cudaMalloc(&d_pChangedBase, ThreadNum * sizeof(int));


	const int kUnrollAdvExecuteScale = 8;
	const int kUnrollAdvkUnrollScale = 16;
	const int preCalcCount = 8;
#pragma omp parallel for num_threads(ThreadNum)
	for (size_t itstm = 0; itstm < ThreadNum; itstm++)
	{
		float* h_Src = h_SrcBase + itstm * kSrcCount * kCoords;
		float* d_pSrc = d_pSrcBase + itstm * kSrcCount * kCoords;
		int* h_MemberShip = h_MemberShipBase + itstm * kSrcCount;
		int* d_pMemberShip = d_pMemberShipBase + itstm * kSrcCount;
		int* h_MemberCount = h_MemberCountBase + itstm * kClusterCount;
		int* d_pMemberCount = d_pMemberCountBase + itstm * kClusterCount;
		float* h_Clusters = h_ClustersBase + itstm * kClusterCount * kCoords;
		float* d_pClusters = d_pClustersBase + itstm * kClusterCount * kCoords;
		int* changed = changedBase + itstm;
		int* d_pChanged = d_pChangedBase + itstm;
		//cudaMemcpy(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pSrc, h_Src, kSrcCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm[itstm]);
		//cudaMemcpy(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pMemberShip, h_MemberShip, kSrcCount * sizeof(int), cudaMemcpyHostToDevice, stm[itstm]);
		//cudaMemcpy(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice);
		cudaMemcpyAsync(d_pMemberCount, h_MemberCount, kClusterCount * sizeof(int), cudaMemcpyHostToDevice, stm[itstm]);
		//cudaMemcpy(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice);//默认流导致和stm流之间同步的时间消耗几百微秒
		cudaMemcpyAsync(d_pClusters, h_Clusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyHostToDevice, stm[itstm]);
		cudaEventRecord(event[itstm* StmEventCount], stm[itstm]);
		//TODO:find the points
		cudaStreamWaitEvent(stm[itstm], event[ThreadNum * StmEventCount + itstm], 0);
		int itCount = 1;


		for (; itCount < preCalcCount; itCount++)
		{
			kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float), stm[itstm] >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);

			//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm[itstm]);
			//cudaMemset(d_pChanged, 0, sizeof(int));
			cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm[itstm]);

			cudaEventRecord(event[(itCount % 10) + itstm * StmEventCount], stm[itstm]);

			//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
			cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm[itstm]);

			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale, stm[itstm] >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
				//kKmeansSumAtomic << <kGridCount* 1, kBlockSize / 1, (kBlockSize* kCoords / 1 + kClusterCount * kCoords + kClusterCount) * sizeof(float) >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);
			kCalcNewClusterCenter << <1, kClusterCount, 0, stm[itstm] >> > (d_pClusters, d_pMemberCount);

			cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm[itstm]);

		}
		cudaEventSynchronize(event[(itCount % 10) + itstm * StmEventCount - 4]);
		while ((0 != *changed) && itCount++ < 100)
		{
			kKmeansClassifyMembershipMath<kCoords, kUnrollAdvkUnrollScale, kClusterCount, kClusterUnroll> << <kGridCount * 1, kBlockSize / kUnrollAdvkUnrollScale, kClusterCount* kCoords * sizeof(float), stm[itstm] >> >
				(kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pChanged);
			//cudaMemcpy(&iChanged, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost);
			cudaMemcpyAsync(changed, d_pChanged, sizeof(int), cudaMemcpyDeviceToHost, stm[itstm]);
			cudaEventRecord(event[itCount % 10], stm[itstm]);
			//cudaMemset(d_pChanged, 0, sizeof(int));
			cudaMemsetAsync(d_pChanged, 0, sizeof(int), stm[itstm]);
			//cudaMemset(d_pClusters, 0, kClusterCount * kCoords * sizeof(float));
			cudaMemsetAsync(d_pClusters, 0, kClusterCount * kCoords * sizeof(float), stm[itstm]);
			kkKmeansUnrollAdv<kBlockSize / kExecuteScale, kUnrollAdvkUnrollScale, kCoords, kClusterCount> << <kGridCount * kExecuteScale / kUnrollAdvkUnrollScale, kBlockSize / kExecuteScale, (kBlockSize * kCoords + kBlockSize) * sizeof(float) / kExecuteScale, stm[itstm] >> > (kSrcCount, kCoords, kClusterCount, d_pSrc, d_pMemberShip, d_pClusters, d_pMemberCount);//s2
			kCalcNewClusterCenter << <1, kClusterCount, 0, stm[itstm] >> > (d_pClusters, d_pMemberCount);
			cudaMemsetAsync(d_pMemberCount, 0, kClusterCount * sizeof(int), stm[itstm]);
		}
		printf("tid %d and count %d \n", itstm, itCount);
		cudaMemcpyAsync(h_Clusters, d_pClusters, kClusterCount * kCoords * sizeof(float), cudaMemcpyDeviceToHost, stm[itstm]);
		for (size_t i = 0; i < kClusterCount; i++)
		{
			printf("tid %d and value %f \n", itstm, h_Clusters[i]);
		}
	}


	//free(h_Src);
	cudaFreeHost(h_SrcBase);
	free(h_ClustersBase);
	//free(h_MemberShip);
	cudaFreeHost(h_MemberShipBase);
	cudaFreeHost(changedBase);

	free(h_MemberCountBase);
	cudaFree(d_pSrcBase);
	cudaFree(d_pClustersBase);
	cudaFree(d_pChangedBase);
	cudaFree(d_pMemberShipBase);
	cudaFree(d_pMemberCountBase);

	for (size_t i = 0; i < ThreadNum + 1; i++)
	{
		cudaStreamDestroy(stm[i]);
	}

	for (size_t i = 0; i < (1 + StmEventCount) * ThreadNum; i++)
	{
		cudaEventDestroy(event[i]);
	}

}

调度情况
在这里插入图片描述
执行时间8.3ms，同时从第二次内存拷贝开始，实现了overlap。但是经过多次测试，这种调度方式的执行结果并不稳定。

总结

对比五个执行结果，通过添加流中事件同步，保证执行顺序，调度结果相对更加稳定，执行时间相对较快，可以作为优先考虑的使用方式。

USST_Coder
关注
  • 1
    点赞
  • 踩
  • 2
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
关于 Nsight Monitor 启动时出现port8000 怎么解决
ZIV555的博客
05-31 1598
因为需要打开Nsight Monitor调整TDR，可是每次打开总是会粗线，下面情况。 首先我们得找出占用这个端口的值。 1、首先我们打开CMD，用管理员身份打开，然后输入netstat -ano，列出所有端口的情况。 2、查看被占用端口对应的PID，输入命令：netstat -aon|findstr "8000"，回车，记下最后一位数字，即PID,如果是2839。 3、（可用可不用）
K-means CUDA实现
05-13
CUDA实现K-means算法，学习GPU编程入门级算法需要了解
【CUDA并行编程之八】Cuda实现Kmeans算法
懂幸福，爱生活
02-01 9613
本文主要介绍如何使用CUDA并行计算框架编程实现机器学习中的Kmeans算法，Kmeans算法的详细介绍在这里，本文重点在并行实现的过程。
python多线程实现kmeans,3种方式实现python多线程并发处理
weixin_33758343的博客
03-26 323
标签： python奇淫技巧最优线程数Ncpu=CPU的数量Ucpu=目标CPU使用率W/C=等待时间与计算时间的比率为保持处理器达到期望的使用率，最优的线程池的大小等于$$Nthreads=Ncpu*Ucpu*(1+W/C$$cpu密集型任务，即$W<如果希望CPU利用率为100%，则$Nthreads=Ncpu$IO密集型任务，即系统大部分时间在跟I/O交互，而这个时间线程不会占用CPU...
redis cluster 操作命令
u012693823的博客
07-23 2627
1：查看集群信息： cluster info > cluster info cluster_state:ok #集群状态，ok：在线；fail：槽未分配或则无法联系到多数master节点 cluster_slots_assigned:16384 #已分配到集群节点的哈希槽数量（不是没有被绑定的数量）。 cluster_slots_ok:16384 #正确分配的槽位，哈希槽状态不是FAIL 和 PFAIL 的数量. ...
基于Python——Kmeans聚类算法的实现
01-20
本篇博文为数据挖掘算法系列的第一篇。现在对于Kmeans算法进行简单的介绍，Kmeans算法是属于无监督的学习的算法，并且是最基本、最简单的一种基于距离的聚类算法。 下面简单说一下Kmeans算法的步骤： 选随机选取K的...
cuda-使用cuda并行加速实现之kmeans聚类算法的实现.zip
最新发布
04-24
cuda cuda_使用cuda并行加速实现之kmeans聚类算法的实现
实验报告——Kmeans聚类方法.docx
01-22
实验报告——Kmeans聚类方法 K-means 聚类方法是最为经典的基于划分的聚类方法之一，其基本思想是以空间中 k 个点为中心进行聚类，对最靠近他们的对象归类。通过迭代的方法，逐次更新各聚类中心的值，直至得到最好...
kmeans-master_基于cuda的kmeans算法_cuda加速kmeans_源码
10-01
使用cuda将kmeans算法进行并行化，在linux下运行
粒子群算法、粒子群优化kmeans、fcm
06-29
智能算法对传统聚类算法的优化，亲测可用
clusterLayer点聚合进阶
c17854254215的博客
08-08 4958
在arcgis api for js中有一个clusterlayer，可以实现点聚合显示，以解决点过于密集导致的视觉效果差。clusterLayer是重写的graphicLayer。官方demo效果如下： 注意:已知api版本3.18--3.25均有bug，表现为点击单个点时属性信息不变，一直显示同一个点的属性信息，目前已知3.9版本没有该问题。 如果我们要将该方法应用到项目中，显然仅有官...
CUDA基本优化方法
nn1997729的博客
06-24 3105
一、基于线程的优化方法 1.选取合适的gridDim和blockDim blockDim最好为32的整数倍：因为执行指令的基本单位为线程束，线程束内的所有线程统一执行广播下来的命令，而线程束的线程数量基本为32。当block被分到SM中去，其会被划分为多个线程束，若blockDim!=线程束内线程数整数倍，则会造成线程的浪费。 2.减少存在分支的if 因为线程束中的所有线程执行同一条的指令，若出现存在出现分支的if，即线程束内的部分线程符合if条件，部分符合else条件，因cuda编译不具备分支预测能力，则
win10下VS15环境下Nsight Debug的Connection to the Nsight monitor on 'localhost' failed.的解决
把你挖出来的博客
03-31 3601
在win10下VS15环境下，对CUDA程序GPU部分进行断点调试，出现问题如下： 解决方法如下： 1)右击Nsight图标，选中Options… 2)修改两项，修改好如下图所示： 3)重启Nsight软件，接下来就好调试了！
k-means算法优化
qq_42182596的博客
05-26 1688
一、k-means算法介绍 k-means 算法是首先从含有n个数据对象的数据集中随机选择K个数据对象作为初始中心。 然后计算每个数据对象到各中心的距离，根据最近邻原则，所有数据对象将会被划分到离它最近的那个中心所代表的簇中。 接着分别计算新生成的各个簇中数据对象的均值作为各簇新的中心，比较新的中心和上一次得到的中心，如果没有发生变化，则算法收敛，输出结果； 如果新的中心和上一次的中心相比发生变化，则要以新的中心对所有数据对象重新进行划分。直到满足算法的收敛条件为止。 二、缺点和优化方法 K-means聚类
python 调用gpu算力_KMCUDA： GPU/CUDA 实现Kmeans
weixin_39647977的博客
12-09 1159
最近项目使用到Kmeans算法，考虑到CPU实现速度上的限制，需要使用GPU加速，因此查到libKMCUDA库。记录安装使用过程中遇到的一些问题。1、kmcuda简介项目内容：Large scale K-means and K-nn implementation on NVIDIA GPU / CUDA该项目具体的介绍可参照github上的说明。性能如下：从技术上来讲，该项目是一个共享库，可导出k...
MegEngine| CUDA 矩阵乘法终极优化
旷视开发者技术社区
09-24 1136
前言单精度矩阵乘法（SGEMM）几乎是每一位学习 CUDA 的同学绕不开的案例，这个经典的计算密集型案例可以很好地展示 GPU 编程中常用的优化技巧，而能否写出高效率的 SGEMM Ker...
CUDA 学习（二十）、优化策略5： 算法
10-31 754
一、概述        如果要考虑算法，我们需要先思考如下问题：       （1）、如何将问题分解成块或片，然后如何将这些块分解成线程；       （2）、线程如何访问数据以及产生什么样的内存模式；       （3）、算法总共要执行多少工作以及串行化的实现方法有何显著不同； 二、排序         可用的排序算法有很多，其中有一些可以轻松、高效地在GPU 上实现，但也有
算法优化系列笔记——cuda入门
Lo_Bamboo的博客
05-05 1015
1 引言 算法在工程化过程中，最躲不开就是算法的优化问题。优化分很多个方向，最简单的实现方式是并行化加速。如：一个向量相加，在cpu中你是串行一个元素一个元素的加减，如果采用并行化加速，你可以一次操作，可以将向量加法完成。当然，并行化的实现方式大概分两类：(1) cpu多线程的方法，如：openmp，(2)异构计算的方法。如gpu加速，fpga加速，NPU加速等。cpu多线程的方法的必要条件就是...
Python——Kmeans聚类算法
04-29
Kmeans聚类算法是一种无监督学习算法，用于将数据集划分为不同的簇。它是一个迭代算法，通过计算每个数据点与簇中心的距离，将数据点分配到最近的簇中心。然后，根据分配的数据点更新簇中心。重复这个过程，直到簇...

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

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

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值