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

我不得不承认，这个内核运行得很慢。一旦我将阵列返回到主机端，我使用thrust::inclusive_scan 来查找我的结果阵列。 我的问题是，我的内核是否还有改进/优化的空间？我尝试使用共享内存，但它会产生错误的答案或引发运行时异常。

另外，动态分配的共享内存（由内核启动中的第三个参数分配）是如何在块之间分配的？

任何帮助/提示/暗示将不胜感激。 提前致谢。

Answer 1

至于使用kernel<<<blocks,threads,mem>>>mem分配的共享内存是每个块分配的内存量。所以每个块都会获得mem 的内存量。

对于您的代码，我不明白为什么 mult 函数中有 2 个 for 循环。只想指出每个线程都将执行这 2 个 for 循环。此外，由于kernel 函数中还有一个 for 循环，这意味着每个线程将多次执行 mult 函数中的 2 个for 循环。这很慢。此外，做

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

不完全是合并访问。非合并访问很慢。