二重循环，如何变成正确的cuda程序？

system · 2012 年12 月 9 日 14:58

仿照《CUDA by example》中“任意向量加法”的处理，我改写了CPUMatches函数中的两重循环，变成了GPUMatches，但是运行之后，为什么显示的2个结果不一致啊？多谢啦


运行输出是
cpu matched=1057707
gpu matched=258048

程序如下


#define _CRT_SECURE_NO_DEPRECATE
#include <stdio.h>
#include <stdlib.h>
void CPUMatches(int picnum, int TotalPtL, int *host_match){
   int PointIdxL=0, j=0;
   for (PointIdxL=0; PointIdxL<TotalPtL; PointIdxL++)
   {
   for (j=0; j<picnum; j++)                
   {
   host_match[j+PointIdxL*picnum]=12;                        
   }        
   }    
}
__global__ void GPUMatches(int picnum, int TotalPtL, int *dev_match){
   int PointIdxL=threadIdx.x+blockIdx.x*blockDim.x;
   int j=threadIdx.y+blockIdx.y*blockDim.y;
   while (PointIdxL<TotalPtL)
   {
   while (j<picnum)                
   {
   dev_match[j+PointIdxL*picnum]=12;            
   j+=blockDim.y*gridDim.y;
   }
   PointIdxL+=blockDim.x*gridDim.x;
   }    
}
int CheckMatch(int picnumR, int TotalPtL, int *match){
   int matched=0, i=0, j=0;
   for (i=0;i<TotalPtL;i++)
   for(j=0;j<picnumR;j++)
   if (match[j+i*picnumR]==12)
   matched++;
   return matched;
}

int main(void){
   int i=0, j=0, TotalPtL=16789, picnumR=63, matched=0;
   
   int *host_match=NULL;
   int *dev_match=NULL;    
   host_match=(int *)calloc(TotalPtL*picnumR, sizeof(int));
   cudaMalloc((void**)&dev_match, TotalPtL*picnumR*sizeof(int));
   
   for (i=0;i<TotalPtL;i++)
   for(j=0;j<picnumR;j++)
   host_match[j+i*picnumR]=-1;
   
   cudaMemcpy(dev_match, host_match, TotalPtL*picnumR*sizeof(int), cudaMemcpyHostToDevice);
   
   CPUMatches(picnumR, TotalPtL, host_match);
   matched=CheckMatch(picnumR, TotalPtL, host_match);
   printf("cpu matched=%i\n\n", matched);
   
   dim3 grids(256, 256);
   dim3 threads(16, 16);
   GPUMatches<<<grids, threads>>> (picnumR, TotalPtL, dev_match);
   
   cudaMemcpy(host_match, dev_match, 
   TotalPtL*picnumR*sizeof(int), 
   cudaMemcpyDeviceToHost);
   
   matched=CheckMatch(picnumR, TotalPtL, host_match);
   printf("gpu matched=%i\n\n", matched);   
   
   free(host_match);
   cudaFree(dev_match);
}

system · 2012 年12 月 9 日 15:45

您好，大致浏览了您的代码，你的CPU代码和GPU代码并不等效。

CPU代码是一个简单的两维情况下的循环赋值，一般来说这种循环转换为GPU的实现，只需要用2维的线程铺开，然后kernel里面得到线程对应的位置后，除了判断是否超出边界以外，保留一句赋值就可以了。
而您在kernel里面写了两重的循环，而且循环的逻辑多少有些让人摸不到头脑。

所以您需要重新改写下您的kernel。

欢迎莅临cudazone，祝您编码愉快~

system · 2012 年12 月 10 日 12:45

感谢ice的回复。你的回复，正是我没有弄懂的地方。可以麻烦你明确指出应该如何改这个程序吗？我已经被绕在这样的程序中好几天了，烦躁啊。

你所说的“让人摸不到头脑的两重混换的逻辑”来自于我对“CUDA by example》中“任意向量加法”（5.3节）的理解（如下）


__global__ void add(int *a, int *b, int *c){
   int tid = threadIdx.x + blockIdx.x *  blockDim.x;
   while (tid<N){ 
   c[tid]=a[tid]+b[tid];
   tid+=blockDim.x*gridDim.x;
   }
}

system · 2012 年12 月 10 日 13:36

[list=1]
[]global void GPUMatches(int picnum, int TotalPtL, int dev_match){
[] int PointIdxL=threadIdx.x+blockIdx.xblockDim.x;
[] int j=threadIdx.y+blockIdx.yblockDim.y;
[] while (PointIdxL<TotalPtL)
[] {
[] while (j<picnum)
[] {
[] dev_match[j+PointIdxLpicnum]=12;
[] j+=blockDim.ygridDim.y;
[] }
[] PointIdxL+=blockDim.xgridDim.x;
[] }
[]}
[]改成
[]global void GPUMatches(int picnum, int TotalPtL, int dev_match){
int PointIdxL=threadIdx.x+blockIdx.xblockDim.x;
int j=threadIdx.y+blockIdx.yblockDim.y;
while (PointIdxL<TotalPtL)
{
[]int index = j;
while (index <picnum)
{
dev_match[j+PointIdxLpicnum]=12;
index+=blockDim.ygridDim.y;
}
PointIdxL+=blockDim.xgridDim.x;
}
}
[*]

[]在你原来的程序中，只要第一次j增加到picnum后，外面的循环就都不干活，因此只有很少一部分的数据值为12（你可以将传回给cpu的数据打印出来看看）。
[]

[*]另外我手头没有机器，没做过实验，请自己验证
[/list]

system · 2012 年12 月 10 日 13:59

不太了解cuda by example的例子，您可以先按照4#风辰大给出的建议修改一下试试。

祝您编码愉快~

system · 2012 年12 月 11 日 02:17

//楼主把内核改成下面这样：
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if(col < picnum && row < TotalPtL)
   dev_match[col+row*picnum]=12;  
//配置内核改成：
dim3 grid;
dim3 block;
grid.x = picnumR % 16 == 0 ? picnumR / 16 : picnumR / 16 + 1;
grid.y = TotalPtL % 16 == 0 ? TotalPtL / 16 : TotalPtL / 16 + 1;
block.x = 16;
block.y = 16;

system · 2012 年12 月 12 日 02:43

感谢大家的回答，dev_match是二维的整型数组，大小是picnumRTotalPtL。我的最初的程序当中，在核函数中，两个while并不是奇怪的循环，通过和下面的“PointIdxL+=blockDim.xgridDim.x”这样的语句连用，是为了避免出现线程格超过65535的问题（《GPU高性能编程》——即翻译版的《cuda by example》——P47“在GPU上对任意长度的矢量求和”提到这个方法）

我的原始程序仍然有错误，也就是说我还有更多没有搞懂的问题，我会逐渐分解每一个问题，在这里请教大家。再次感谢

system · 2012 年12 月 12 日 03:15

在kepler的GPU上，Maximun x-dimension of a grid of thread blocks 可以是2^31-1，或许在新的硬件上您可以简化您的算法，当然老硬件上依然需要自己折腾。

另外，在贴源码之余，如果能简要介绍实现意图和预期实现的步骤，可以方便其他人的理解。

祝您编码愉快。

system · 2012 年12 月 12 日 03:22

这里简化之后的程序，只是为了给dev_match的所有元素赋值为数字12，别无它用

system · 2012 年12 月 12 日 06:25

您好，8#让您简要介绍预期实现的步骤主要是为了让您阐明顶楼中kernel为何会写成那个样子，比如哪个变量是干什么的，每个线程大致要干什么活等。

刚才我又重新回溯了上述各楼，现在我大致总结下，仅供参考：

1：按照您7#给出的叙述，“dev_match是二维的整型数组，大小是picnumR*TotalPtL。”以及顶楼给出的串行实现，这两者是相符合。应该说实现此问题并无难处，按照6#的代码直接用或者根据规模稍作修改即可。您在7#提到，block在一个维度上有65535个的限制，其实这个限制是block自己数量的限制，每个block内部最多可以有1024个threads，这样算下来，是相当大的一个数了，您可以考虑下，如此是否依然会超出block数量的限制，即使超出，也有办法。

2：您在7#提到您在顶楼的kernel是参考 CUDA by Example里面的例子写的，并在3#给出了该例子的代码摘要。目测此代码，大致是一个二维的数组，用一维铺开的线程一行一行地计算，每个线程先通过tid的初始值找到自己对应的起始位置，然后通过tid+=blockDim.x*gridDim.x找到下一行自己需要计算的位置，这里的增量恰是全局二维数组一行的长度。这样，原本一个串行二重循环的问题，现在用线程拆掉了一重循环，每个线程自己维护一个一重循环。

3：回到您顶楼的kernel，您的串行实现本身就是二重循环，而您写的并行实现也是二重循环，（当然二重循环并非一定有问题），根据您的代码，按照4#风辰指出的问题，将内循环不干活的问题解决掉以后，是什么行为呢？假定我们考虑最开始的那个thread，那么此时，PointIdxL==0，j==0，进入循环后，内层循环先执行一次赋值，然后j+=blockDim.y*gridDim.y，一般来说，这种用法是通过线程分布和数组分布对应实现寻址的，但您上面的赋值却不是这种用法，按照您赋值和while循环的写法，这里似乎应该每次增加1；而按照j自增的写法，似乎很快就跳出循环了；或者您完全依靠j的自增完成本线程的任务，似乎外面的循环就不需要了。或者如果picnum和TotalPtL值特别大，而您的整个grid里面的线程数也比这个要算的点数小很多，那么内循环算一个picnum里面的若干个离散的位置，外循环算定位TotalPtL里面若干个离散的位置。相当于将最大的数组分成了若干个等同于线程数量的块，然后每个线程计算每个块里面的对应位置，这样设想似乎是合适的。（因为您没有给出您的规划和如何实现想法的描述，我只能在这里反复猜）。

4：在3里面，最后那种行为是可以的，只是略显折腾，如果您也是如此考虑的，您只需仔细检查您的代码是否遵从了您的想法，以及有没有什么BUG。当然也可以选择不要那么折腾，尽量简单一点实现，比如用一维的block等。

祝您编码愉快~

system · 2012 年12 月 12 日 10:39

我已经验证了我给你的算法是正确的,输出如下:cpu matched=1057707

gpu matched=1057707
上次比较忙,没有给出代码,今天补上.

#include <stdio.h>
#include <stdlib.h>
void CPUMatches(int picnum, int TotalPtL, int *host_match){
   int PointIdxL=0, j=0;
   for (PointIdxL=0; PointIdxL<TotalPtL; PointIdxL++)
   {
   for (j=0; j<picnum; j++)                
   {
   host_match[j+PointIdxL*picnum]=12;                        
   }        
   }    
}
__global__ void GPUMatches(int picnum, int TotalPtL, int *dev_match){
   int PointIdxL=threadIdx.x+blockIdx.x*blockDim.x;
   int j=threadIdx.y+blockIdx.y*blockDim.y;
   while (PointIdxL<TotalPtL)
   { int index = j;
   while (index<picnum)                
   {
   dev_match[j+PointIdxL*picnum]=12;            
   index+=blockDim.y*gridDim.y;
   }
   PointIdxL+=blockDim.x*gridDim.x;
   }    
}
int CheckMatch(int picnumR, int TotalPtL, int *match){
   int matched=0, i=0, j=0;
   for (i=0;i<TotalPtL;i++)
   for(j=0;j<picnumR;j++)
   if (match[j+i*picnumR]==12)
   matched++;
   return matched;
}

int main(void){
   int i=0, j=0, TotalPtL=16789, picnumR=63, matched=0;
   
   int *host_match=NULL;
   int *dev_match=NULL;    
   host_match=(int *)calloc(TotalPtL*picnumR, sizeof(int));
   cudaMalloc((void**)&dev_match, TotalPtL*picnumR*sizeof(int));
   
   for (i=0;i<TotalPtL;i++)
   for(j=0;j<picnumR;j++)
   host_match[j+i*picnumR]=-1;
   
   cudaMemcpy(dev_match, host_match, TotalPtL*picnumR*sizeof(int), cudaMemcpyHostToDevice);
   
   CPUMatches(picnumR, TotalPtL, host_match);
   matched=CheckMatch(picnumR, TotalPtL, host_match);
   printf("cpu matched=%i\n\n", matched);
   
   dim3 grids(256, 256);
   dim3 threads(16, 16);
   GPUMatches<<<grids, threads>>> (picnumR, TotalPtL, dev_match);
   
   cudaMemcpy(host_match, dev_match, 
   TotalPtL*picnumR*sizeof(int), 
   cudaMemcpyDeviceToHost);
   
   matched=CheckMatch(picnumR, TotalPtL, host_match);
   printf("gpu matched=%i\n\n", matched);   
   
   free(host_match);
   cudaFree(dev_match);
}

system · 2012 年12 月 13 日 02:07

感谢 ice 和风辰的耐心指教:handshake
最终的代码是将2D数组某一维的索引值的处理，全部放在花括号里，也就是如下的红色部分

（另，啥级别的普通用户，才可以回复的时候，不填“验证问答”和“验证码”？）

global void GPUMatches(int picnum, int TotalPtL, int dev_match){
int PointIdxL=threadIdx.x+blockIdx.xblockDim.x;
int j=threadIdx.y+blockIdx.yblockDim.y;
while (PointIdxL<TotalPtL)
{
while (j<picnum)
{
dev_match[j+PointIdxLpicnum]=12;
j+=blockDim.ygridDim.y;
}
PointIdxL+=blockDim.xgridDim.x;
}
}

变成了
global void GPUMatches(int picnum, int TotalPtL, int dev_match){
int j=0;
int PointIdxL=threadIdx.x+blockIdx.xblockDim.x;
while (PointIdxL<TotalPtL) {
j=threadIdx.y+blockIdx.yblockDim.y;
while(j<picnum) {
if((j+PointIdxLpicnum)<picnumTotalPtL)
dev_match[j+PointIdxLpicnum]=12;
j+=blockDim.ygridDim.y;
}
PointIdxL+=blockDim.xgridDim.x;
}
}

system · 2012 年12 月 13 日 04:19

问题搞定就好，前一段时间机器人灌水严重，所以后台添加了验证环节，至于多少级别可以无需验证，这个我也不知道，是后台管理员设的…

祝您编码愉快！