views:

143

answers:

1

Hi,

I can inprove this function under CUDA?

What this function does is:

Given a min and max, ELM1 and ELM, check if any three numbers of array ans[6] are found in any row, from min to max, in array D1,D2,D3,D4,D5,D6, if found return 1

I tried any other way, like looping, or-ing, and-ing, replacing goto with flag etc. etc. but this seems to be the fastest way.

 __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;


}
+1  A: 

remove if statements by converting them to Boolean expressions.

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

make sure your arrays are in registers or shared memory rather than global

also, on such simple algorithm you got way too complicated logic. change layout of D arrays to be D[N][6] as it will simplify many things

by the way, you may want to crop your post a little bit, way too much to read

3 x3 example

     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

basically you are setting matrix cells to true if match A is in array D. on each iteration you reduce columns to true if entire column is true. then you count number of true.

aaa
well, the D1...D6 arrays are constant, which seems to be faster than shared. With D[N][6] would be more complicated, I tried it and it is slower.
Mark
with flag += (DN[ELM1] == ans[0]) takes 700 ms more
Mark
@Mark how many threads are you using?
aaa
more than 13000000
Mark
@Mark I assume each thread does its own work? how many threads per block?if I were you I would change algorithm to produce 6x6 matrix of true/false which you check at the end of each iteration to see if any three rows is all true. if does not make sense, I will post picture
aaa
yes, please post the picture
Mark
@Mark I need to know how many threads per block you using
aaa
256 threads per block
Mark
I tried again with D[N][6], but the kernel gives me the out of memory error, also I did not understand how to implement the 3x3 example
Mark