【问题标题】:How can I improve this function under CUDA?如何在 CUDA 下改进此功能?
【发布时间】:2011-03-02 08:22:02
【问题描述】:

CUDA下面的功能可以改进吗?

函数的作用是,

给定一个minmaxELM1ELM,检查数组@987654329的任意一行中是否存在任意三个数组ans[6],从minmax,在数组@987654329中@、D2D3D4D5D6,如果找到返回 1。

我尝试使用loopsOR-ing、AND-ing,用标志等替换goto。但这似乎是最快的方法。

 __device__ bool THREEA(unsigned int n0, unsigned int n,unsigned int* ST1,unsigned int* D1, unsigned int* D2,unsigned int* D3,unsigned int* D4,unsigned int* D5,unsigned int* D6,unsigned int* ans)
{
     unsigned int ELM, ELM1,flag;
     ELM = ST1[n0]+n;  //local.37

     ELM1 = n;       //local.33
     while (ELM1 < ELM)
     {

         flag = 0;
         if (D1[ELM1] == ans[0])
         {
          flag++;
         }
         if (D2[ELM1] == ans[0])
         {
          flag++;
         }
         if (D3[ELM1] == ans[0])
         {
          flag++;
         }
         if (D4[ELM1] == ans[0])
         {
          flag++;
         }
         if (D5[ELM1] == ans[0])
         {
          flag++;
         }
         if (D6[ELM1] == ans[0])
         {
          flag++;
         }
         if (flag != 1)
          goto onethreefour;
         if (D1[ELM1] == ans[1])
         {
          flag++;
         }
         if (D2[ELM1] == ans[1])
         {
          flag++;
         }
         if (D3[ELM1] == ans[1])
         {
          flag++;
         }
         if (D4[ELM1] == ans[1])
         {
          flag++;
         }
         if (D5[ELM1] == ans[1])
         {
          flag++;
         }
         if (D6[ELM1] == ans[1])
         {
          flag++;
         }
         if (flag != 2)
          goto onethreefour;
         if (D1[ELM1] == ans[2])
         {
          return 1;
         }
         if (D2[ELM1] == ans[2])
         {
          return 1;
         }
         if (D3[ELM1] == ans[2])
         {
          return 1;
         }
         if (D4[ELM1] == ans[2])
         {
          return 1;
         }
         if (D5[ELM1] == ans[2])
         {
          return 1;
         }
         if (D6[ELM1] == ans[2])
         {
          return 1;
         }
         if (D1[ELM1] == ans[3])
         {
          return 1;
         }
         if (D2[ELM1] == ans[3])
         {
          return 1;
         }
         if (D3[ELM1] == ans[3])
         {
          return 1;
         }
         if (D4[ELM1] == ans[3])
         {
          return 1;
         }
         if (D5[ELM1] == ans[3])
         {
          return 1;
         }
         if (D6[ELM1] == ans[3])
         {
          return 1;
         }
         if (D1[ELM1] == ans[4])
         {
          return 1;
         }
         if (D2[ELM1] == ans[4])
         {
          return 1;
         }
         if (D3[ELM1] == ans[4])
         {
          return 1;
         }
         if (D4[ELM1] == ans[4])
         {
          return 1;
         }
         if (D5[ELM1] == ans[4])
         {
          return 1;
         }
         if (D6[ELM1] == ans[4])
         {
          return 1;
         }
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }

onethreefour:
         flag = 0;
         if (D1[ELM1] == ans[0])
         {
          flag++;
         }
         if (D2[ELM1] == ans[0])
         {
          flag++;
         }
         if (D3[ELM1] == ans[0])
         {
          flag++;
         }
         if (D4[ELM1] == ans[0])
         {
          flag++;
         }
         if (D5[ELM1] == ans[0])
         {
          flag++;
         }
         if (D6[ELM1] == ans[0])
         {
          flag++;
         }
         if (flag != 1)
          goto onefourfive;
         if (D1[ELM1] == ans[2])
         {
          flag++;
         }
         if (D2[ELM1] == ans[2])
         {
          flag++;
         }
         if (D3[ELM1] == ans[2])
         {
          flag++;
         }
         if (D4[ELM1] == ans[2])
         {
          flag++;
         }
         if (D5[ELM1] == ans[2])
         {
          flag++;
         }
         if (D6[ELM1] == ans[2])
         {
          flag++;
         }
         if (flag != 2)
          goto onefourfive;
         if (D1[ELM1] == ans[3])
         {
          return 1;
         }
         if (D2[ELM1] == ans[3])
         {
          return 1;
         }
         if (D3[ELM1] == ans[3])
         {
          return 1;
         }
         if (D4[ELM1] == ans[3])
         {
          return 1;
         }
         if (D5[ELM1] == ans[3])
         {
          return 1;
         }
         if (D6[ELM1] == ans[3])
         {
          return 1;
         }
         if (D1[ELM1] == ans[4])
         {
          return 1;
         }
         if (D2[ELM1] == ans[4])
         {
          return 1;
         }
         if (D3[ELM1] == ans[4])
         {
          return 1;
         }
         if (D4[ELM1] == ans[4])
         {
          return 1;
         }
         if (D5[ELM1] == ans[4])
         {
          return 1;
         }
         if (D6[ELM1] == ans[4])
         {
          return 1;
         }
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }


onefourfive:
         flag = 0;
         if (D1[ELM1] == ans[0])
         {
          flag++;
         }
         if (D2[ELM1] == ans[0])
         {
          flag++;
         }
         if (D3[ELM1] == ans[0])
         {
          flag++;
         }
         if (D4[ELM1] == ans[0])
         {
          flag++;
         }
         if (D5[ELM1] == ans[0])
         {
          flag++;
         }
         if (D6[ELM1] == ans[0])
         {
          flag++;
         }
         if (flag != 1)
          goto onefivesix;
         if (D1[ELM1] == ans[3])
         {
          flag++;
         }
         if (D2[ELM1] == ans[3])
         {
          flag++;
         }
         if (D3[ELM1] == ans[3])
         {
          flag++;
         }
         if (D4[ELM1] == ans[3])
         {
          flag++;
         }
         if (D5[ELM1] == ans[3])
         {
          flag++;
         }
         if (D6[ELM1] == ans[3])
         {
          flag++;
         }
         if (flag != 2)
          goto onefivesix;
         if (D1[ELM1] == ans[4])
         {
          return 1;
         }
         if (D2[ELM1] == ans[4])
         {
          return 1;
         }
         if (D3[ELM1] == ans[4])
         {
          return 1;
         }
         if (D4[ELM1] == ans[4])
         {
          return 1;
         }
         if (D5[ELM1] == ans[4])
         {
          return 1;
         }
         if (D6[ELM1] == ans[4])
         {
          return 1;
         }
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
         return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
         return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }

onefivesix:
         flag = 0;
         if (D1[ELM1] == ans[0])
         {
          flag++;
         }
         if (D2[ELM1] == ans[0])
         {
          flag++;
         }
         if (D3[ELM1] == ans[0])
         {
          flag++;
         }
         if (D4[ELM1] == ans[0])
         {
          flag++;
         }
         if (D5[ELM1] == ans[0])
         {
          flag++;
         }
         if (D6[ELM1] == ans[0])
         {
          flag++;
         }
         if (flag != 1)
          goto twothreefour;
         if (D1[ELM1] == ans[4])
         {
          flag++;
         }
         if (D2[ELM1] == ans[4])
         {
          flag++;
         }
         if (D3[ELM1] == ans[4])
         {
          flag++;
         }
         if (D4[ELM1] == ans[4])
         {
          flag++;
         }
         if (D5[ELM1] == ans[4])
         {
          flag++;
         }
         if (D6[ELM1] == ans[4])
         {
          flag++;
         }
         if (flag != 2)
          goto twothreefour;
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }
twothreefour:
         flag = 0;
         if (D1[ELM1] == ans[1])
         {
          flag++;
         }
         if (D2[ELM1] == ans[1])
         {
          flag++;
         }
         if (D3[ELM1] == ans[1])
         {
          flag++;
         }
         if (D4[ELM1] == ans[1])
         {
          flag++;
         }
         if (D5[ELM1] == ans[1])
         {
          flag++;
         }
         if (D6[ELM1] == ans[1])
         {
          flag++;
         }
         if (flag != 1)
          goto twofourfive;
         if (D1[ELM1] == ans[2])
         {
          flag++;
         }
         if (D2[ELM1] == ans[2])
         {
          flag++;
         }
         if (D3[ELM1] == ans[2])
         {
          flag++;
         }
         if (D4[ELM1] == ans[2])
         {
          flag++;
         }
         if (D5[ELM1] == ans[2])
         {
          flag++;
         }
         if (D6[ELM1] == ans[2])
         {
          flag++;
         }
         if (flag != 2)
          goto twofourfive;
         if (D1[ELM1] == ans[3])
         {
          return 1;
         }
         if (D2[ELM1] == ans[3])
         {
          return 1;
         }
         if (D3[ELM1] == ans[3])
         {
          return 1;
         }
         if (D4[ELM1] == ans[3])
         {
          return 1;
         }
         if (D5[ELM1] == ans[3])
         {
          return 1;
         }
         if (D6[ELM1] == ans[3])
         {
          return 1;
         }
         if (D1[ELM1] == ans[4])
         {
          return 1;
         }
         if (D2[ELM1] == ans[4])
         {
          return 1;
         }
         if (D3[ELM1] == ans[4])
         {
          return 1;
         }
         if (D4[ELM1] == ans[4])
         {
          return 1;
         }
         if (D5[ELM1] == ans[4])
         {
          return 1;
         }
         if (D6[ELM1] == ans[4])
         {
          return 1;
         }
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }
twofourfive:
         flag = 0;
         if (D1[ELM1] == ans[1])
         {
          flag++;
         }
         if (D2[ELM1] == ans[1])
         {
          flag++;
         }
         if (D3[ELM1] == ans[1])
         {
          flag++;
         }
         if (D4[ELM1] == ans[1])
         {
          flag++;
         }
         if (D5[ELM1] == ans[1])
         {
          flag++;
         }
         if (D6[ELM1] == ans[1])
         {
          flag++;
         }
         if (flag != 1)
          goto twofivesix;
         if (D1[ELM1] == ans[3])
         {
          flag++;
         }
         if (D2[ELM1] == ans[3])
         {
          flag++;
         }
         if (D3[ELM1] == ans[3])
         {
          flag++;
         }
         if (D4[ELM1] == ans[3])
         {
          flag++;
         }
         if (D5[ELM1] == ans[3])
         {
          flag++;
         }
         if (D6[ELM1] == ans[3])
         {
          flag++;
         }
         if (flag != 2)
          goto twofivesix;
         if (D1[ELM1] == ans[4])
         {
          return 1;
         }
         if (D2[ELM1] == ans[4])
         {
          return 1;
         }
         if (D3[ELM1] == ans[4])
         {
          return 1;
         }
         if (D4[ELM1] == ans[4])
         {
          return 1;
         }
         if (D5[ELM1] == ans[4])
         {
          return 1;
         }
         if (D6[ELM1] == ans[4])
         {
          return 1;
         }
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }
twofivesix:
         flag = 0;
         if (D1[ELM1] == ans[1])
         {
          flag++;
         }
         if (D2[ELM1] == ans[1])
         {
          flag++;
         }
         if (D3[ELM1] == ans[1])
         {
          flag++;
         }
         if (D4[ELM1] == ans[1])
         {
          flag++;
         }
         if (D5[ELM1] == ans[1])
         {
          flag++;
         }
         if (D6[ELM1] == ans[1])
         {
          flag++;
         }
         if (flag != 1)
          goto threefourfive;
         if (D1[ELM1] == ans[4])
         {
          flag++;
         }
         if (D2[ELM1] == ans[4])
         {
          flag++;
         }
         if (D3[ELM1] == ans[4])
         {
          flag++;
         }
         if (D4[ELM1] == ans[4])
         {
          flag++;
         }
         if (D5[ELM1] == ans[4])
         {
          flag++;
         }
         if (D6[ELM1] == ans[4])
         {
          flag++;
         }
         if (flag != 2)
          goto threefourfive;
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         } 
threefourfive:
         flag = 0;
         if (D1[ELM1] == ans[2])
         {
          flag++;
         }
         if (D2[ELM1] == ans[2])
         {
          flag++;
         }
         if (D3[ELM1] == ans[2])
         {
          flag++;
         }
         if (D4[ELM1] == ans[2])
         {
          flag++;
         }
         if (D5[ELM1] == ans[2])
         {
          flag++;
         }
         if (D6[ELM1] == ans[2])
         {
          flag++;
         }
         if (flag != 1)
          goto threefivesix;
         if (D1[ELM1] == ans[3])
         {
          flag++;
         }
         if (D2[ELM1] == ans[3])
         {
          flag++;
         }
         if (D3[ELM1] == ans[3])
         {
          flag++;
         }
         if (D4[ELM1] == ans[3])
         {
          flag++;
         }
         if (D5[ELM1] == ans[3])
         {
          flag++;
         }
         if (D6[ELM1] == ans[3])
         {
          flag++;
         }
         if (flag != 2)
          goto threefivesix;
         if (D1[ELM1] == ans[4])
         {
          return 1;
         }
         if (D2[ELM1] == ans[4])
         {
          return 1;
         }
         if (D3[ELM1] == ans[4])
         {
          return 1;
         }
         if (D4[ELM1] == ans[4])
         {
          return 1;
         }
         if (D5[ELM1] == ans[4])
         {
          return 1;
         }
         if (D6[ELM1] == ans[4])
         {
          return 1;
         }
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }
threefivesix:
         flag = 0;
         if (D1[ELM1] == ans[2])
         {
          flag++;
         }
         if (D2[ELM1] == ans[2])
         {
          flag++;
         }
         if (D3[ELM1] == ans[2])
         {
          flag++;
         }
         if (D4[ELM1] == ans[2])
         {
          flag++;
         }
         if (D5[ELM1] == ans[2])
         {
          flag++;
         }
         if (D6[ELM1] == ans[2])
         {
          flag++;
         }
         if (flag != 1)
          goto fourfivesix;
         if (D1[ELM1] == ans[4])
         {
          flag++;
         }
         if (D2[ELM1] == ans[4])
         {
          flag++;
         }
         if (D3[ELM1] == ans[4])
         {
          flag++;
         }
         if (D4[ELM1] == ans[4])
         {
          flag++;
         }
         if (D5[ELM1] == ans[4])
         {
          flag++;
         }
         if (D6[ELM1] == ans[4])
         {
          flag++;
         }
         if (flag != 2)
          goto fourfivesix;
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }
fourfivesix:
         flag = 0;
         if (D1[ELM1] == ans[3])
         {
          flag++;
         }
         if (D2[ELM1] == ans[3])
         {
          flag++;
         }
         if (D3[ELM1] == ans[3])
         {
          flag++;
         }
         if (D4[ELM1] == ans[3])
         {
          flag++;
         }
         if (D5[ELM1] == ans[3])
         {
          flag++;
         }
         if (D6[ELM1] == ans[3])
         {
          flag++;
         }
         if (flag != 1)
          goto increasecounter;
         if (D1[ELM1] == ans[4])
         {
          flag++;
         }
         if (D2[ELM1] == ans[4])
         {
          flag++;
         }
         if (D3[ELM1] == ans[4])
         {
          flag++;
         }
         if (D4[ELM1] == ans[4])
         {
          flag++;
         }
         if (D5[ELM1] == ans[4])
         {
          flag++;
         }
         if (D6[ELM1] == ans[4])
         {
          flag++;
         }
         if (flag != 2)
          goto increasecounter;
         if (D1[ELM1] == ans[5])
         {
          return 1;
         }
         if (D2[ELM1] == ans[5])
         {
          return 1;
         }
         if (D3[ELM1] == ans[5])
         {
          return 1;
         }
         if (D4[ELM1] == ans[5])
         {
          return 1;
         }
         if (D5[ELM1] == ans[5])
         {
          return 1;
         }
         if (D6[ELM1] == ans[5])
         {
          return 1;
         }



increasecounter:

         ELM1++;
     } 
          //If it is Three min

          return 0;


}

【问题讨论】:

  • 用 perl 编写生成器(选择您的语言)并收工? ;)

标签: c optimization cuda


【解决方案1】:

通过将 if 语句转换为布尔表达式来删除它们。

flag += (DN[ELM1] == ans[0])

确保您的数组位于寄存器或共享内存中,而不是全局中

另外,在如此简单的算法上,您的逻辑太复杂了。将 D 数组的布局更改为 D[N][6],因为它会简化很多事情

顺便说一句,你可能想稍微剪裁一下你的帖子,阅读量太大了

3 x3 示例

     A
  |0 0 0|           |x x 0
D |0 0 0| -> ... -> |x x 0 -> reduce down -> |x x o| -> reduce across -> 2x
  |0 0 0|           |x x x

基本上,如果匹配 A 在数组 D 中,则您将矩阵单元设置为 true。 如果整个列为真,则在每次迭代中将列减少为真。 然后你数数为真。

【讨论】:

  • 好吧,D1...D6 数组是恒定的,这似乎比共享更快。使用 D[N][6] 会更复杂,我试过了,速度更慢。
  • 带标志 += (DN[ELM1] == ans[0]) 需要 700 毫秒以上
  • @Mark 我假设每个线程都有自己的工作?每个块有多少个线程?如果我是你,我会更改算法以生成 6x6 真/假矩阵,您在每次迭代结束时检查该矩阵是否有任何三行都为真。如果没有意义,我会发图片
  • 我再次尝试使用 D[N][6],但内核给了我内存不足的错误,我也不明白如何实现 3x3 示例
猜你喜欢
  • 2023-02-08
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 2014-10-20
  • 1970-01-01
  • 1970-01-01
相关资源
最近更新 更多