Ядро Cuda - возможные оптимизации

Question

Вот ядро, которое я запускаю для вычисления некоторого массива параллельно.

__device__ bool mult(int colsize,int rowsize,int *Aj,int *Bi) 
    {  
     for(int j = 0; j < rowsize;j++) 
     {   
      for(int k = 0;k < colsize;k++) 
      { 
       if(Aj[j] == Bi[k]) 
       {  
       return true; 
       }        
      }   
     } 
      return false;  
    } 


__global__ void kernel(int *Aptr,int *Aj,int *Bptr,int *Bi,int rows,int cols,int *Cjc) 
    { 
     int tid = threadIdx.x + blockIdx.x * blockDim.x; 
     int i; 
     if(tid < cols) 
     { 
      int beg = Bptr[tid]; 
      int end = Bptr[tid+1]; 
      for(i = 0;i < rows;i++) 
      { 
       int cbeg = Aptr[i]; 
       int cend = Aptr[i+1]; 
       if(mult(end - beg,cend - cbeg,Aj+cbeg,Bi+beg)) 
       {             
        Cjc[tid+1] += 1; 
        //atomicAdd(Cjc+tid+1,1);   
       } 
      }     
     }    
    } 

Мои конфигурации запуска и вызов ядра следующие.

int numBlocks,numThreads; 

     if(q % 32 == 0) 
     { 
      numBlocks = q/32; 
      numThreads = 32; 
     } 
     else 
     { 
      numBlocks = (q+31)/32; 
      numThreads = 32; 
     } 
findkernel<<<numBlocks,numThreads>>>(devAptr,devAcol,devBjc,devBir,m,q,d_Cjc); 

Я должен признать, что это ядро ​​работает довольно slow.Once я получаю массив обратно в принимающей стороне, я использую thrust::inclusive_scan, чтобы найти мой результирующий массив. Мой вопрос: есть ли место для улучшения/оптимизации для моего ядра? Я пытался использовать разделяемую память, но это порождало неправильные ответы или бросало исключения во время выполнения.

Также как распределяется динамически распределенная разделяемая память (которая выделяется третьим параметром при запуске ядра) среди блоков?

Любая помощь/подсказки/инсинуации будут оценены по достоинству. Спасибо заранее.

Answer 1

Что касается разделяемой памяти, выделенной с использованием kernel<<<blocks,threads,mem>>> mem, это объем памяти, выделенный каждому блоку. Таким образом, каждый блок получает mem объем памяти.

Для вашего кода я не понимаю, почему существуют 2 цикла для функции mult. Просто хочу указать, что каждый поток будет выполнять эти 2 для циклов. Более того, поскольку у вас также есть цикл for в функции kernel, это означает, что каждый поток будет выполнять несколько циклов 2 for в мультифункции. Это медленно. Кроме того, занятие

int beg = Bptr[tid]; 
int end = Bptr[tid+1]; 

не является точно объединенным доступом. Не объединенный доступ медленный.