【发布时间】：2015-05-25 02:22:31
【问题描述】：
我是一个相当新的 cuda 用户。我正在练习我的第一个 cuda 应用程序，我尝试使用 GPU（GTX 670）加速 kmeans 算法。
简而言之，每个线程都在一个点上工作，该点与所有集群中心进行比较，并将一个点分配给具有最小距离的中心（内核代码可以在下面使用 cmets 看到）。
根据 Nsight Visual Studio，我的占用率为 99.61%（1024 个块，每个块 1024 个线程），99.34% 的流式多处理器活动，79.98% 的扭曲问题效率，没有共享内存库冲突，18.4GFLOPs Single MUL 和 55.2 GFLOPs Single ADD（使用给定参数完成 kmeans 内核大约需要 14.5 毫秒）。
根据维基百科，GTX670 的峰值性能为 2460 GFLOPs。我离它还很近。除此之外，一些论文声称它们可以达到峰值性能的一半以上。我看不出我可以进一步优化这个内核代码。有什么可以应用于内核的优化吗？任何建议或帮助表示赞赏，我可以根据需要提供任何其他信息。
提前致谢。
#define SIZE 1024*1024 //number of points
#define CENTERS 32 //number of cluster centroids
#define DIM 8 //dimension of each point and center
#define cudaTHREADSIZE 1024 //threads per block
#define cudaBLOCKSIZE SIZE/cudaTHREADSIZE //number of blocks for kernel
__global__ void kMeans(float *dp, float *dc,int *tag, int *membershipChangedPerBlock)
{
//TOTAL NUMBER OF THREADS SHOULD BE EQUAL TO THE NUMBER OF POINTS, BECAUSE EACH THREAD WORKS ON A SINGLE POINT
__shared__ unsigned char membershipChanged[cudaTHREADSIZE];
__shared__ float dc_shared[CENTERS*DIM];
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int threadID = threadIdx.x;
membershipChanged[threadIdx.x] = 0;
//move centers to shared memory, because each and every thread will call it(roughly + %10 performance here)
while(threadID < CENTERS*DIM){
dc_shared[threadID] = dc[threadID];
threadID += blockDim.x;
}
__syncthreads();
while(tid < SIZE){
int index,prevIndex;
float dist, min_dist;
index = 0;//all initial point indices(centroid number) are assigned to 0.
prevIndex = 0;
dist = 0;
min_dist = 0;
//euclid distance for center 0
for(int dimIdx = 0; dimIdx < DIM; dimIdx++){
min_dist += (dp[tid + dimIdx*SIZE] - dc_shared[dimIdx*CENTERS])*(dp[tid + dimIdx*SIZE] - dc_shared[dimIdx*CENTERS]);
}
//euclid distance for other centers with distance comparison
for(int centerIdx = 1; centerIdx < CENTERS; centerIdx++){
dist = 0;
for(int dimIdx = 0; dimIdx < DIM; dimIdx++){
dist += (dp[tid + dimIdx*SIZE] - dc_shared[centerIdx + dimIdx*CENTERS])*(dp[tid + dimIdx*SIZE] - dc_shared[centerIdx + dimIdx*CENTERS]);
}
//compare distances, if found a shorter one, change index to that centroid number
if(dist < min_dist){
min_dist = dist;
index = centerIdx;
}
}
if (tag[tid] != index) {//if a point's cluster membership changes, flag it as changed in order to compute total membership changes later on
membershipChanged[threadIdx.x] = 1;
}
tag[tid] = index;
__syncthreads();//sync before applying sum reduction to membership changes
//sum reduction
for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
if (threadIdx.x < s) {
membershipChanged[threadIdx.x] +=
membershipChanged[threadIdx.x + s];
}
__syncthreads();
}
if (threadIdx.x == 0) {
membershipChangedPerBlock[blockIdx.x] = membershipChanged[0];
}
tid += blockDim.x * gridDim.x;
}
}
【问题讨论】：
-
我建议提供一个其他人可以编译和运行的完整代码。
-
您在问题的任何地方都没有提到内存带宽。我假设您发布的代码是全局内存带宽有限的。
-
我用 nsight 查找了内存统计报告，这是imgur.com/fa45zfi 的结果。 GTX670 拥有 192.256 GB/s 的全局内存带宽。根据这些结果，L2 缓存使用 137.63 GB/s 带宽。这是否意味着我实际上使用了 0.7 的峰值带宽？
-
@talonmies 在查看 Nsight nvreport 的 Experiment Results 的 Issue Efficiency 选项卡后，我发现 84.55% 的问题停滞原因是执行依赖。每个线程都在一个点上工作，因此一个线程无法计算下一个维度距离并将其与总距离相加而不计算前一个维度距离。我的目标是通过使用线程寄存器来增加并行维度距离计算的数量，最后我将对这些半总距离求和。希望这将通过使用更多的带宽和更少的执行依赖停顿来增加我的 FLOP。
标签： cuda parallel-processing gpgpu k-means nsight