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

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

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

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

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

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


person Recker    schedule 25.08.2012    source источник
comment
можете ли вы просто объяснить простыми словами, что ваш код предназначен делать? тогда у нас больше шансов помочь вам с оптимизацией   -  person    schedule 25.08.2012
comment
@asm ... пытаюсь найти массив JC результирующей разреженной матрицы при умножении двух разреженных матриц ....   -  person Recker    schedule 26.08.2012
comment
хорошо, вы реализуете что-то вроде массивов JC в Matlab. Я постараюсь придумать некоторые идеи   -  person    schedule 26.08.2012
comment
@asm.... У меня есть одна матрица в формате CSR, а другая в формате CSC... вот почему я должен проверить, будут ли в результирующем столбце ненулевые значения или нет... Я делаю это, умножая матрицу B каждого столбца на все строки матрицы A (*** предполагается, что операция выполняется как C = A * B ***)... если есть не нули, я добавлю 1 к каждому местоположению результирующего массива Cjc, чтобы после thrust::inclusive_scan на стороне хоста у меня был правильный массив JC что я хочу....   -  person Recker    schedule 26.08.2012


Ответы (1)


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

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

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

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

person Programmer    schedule 25.08.2012
comment
нет, я не думаю, что будут проблемы с объединением, если у вас нет старой доброй GTX 8800 )) Но я согласен, что код, предоставленный abhinode, не имеет особого смысла .. - person ; 25.08.2012