在用于matlab的CUDA内核中的2D数组上的操作

Question

假设我有以下序列C:

int add(int* a, int* b, int n)
{
    for(i=0; i<n; i++)
    {
        for(j=0; j<n; j++)
        {
            a[i][j]+=b[i][j];
        }
    }

    return 0;
}

我认为对其进行并列化的最佳方法是实现它是一个2D问题并根据CUDA内核使用2D线程块- 嵌套for循环

考虑到这一点,我开始写这样的cuda kernal:

__global__ void calc(int **A, int **B, int n)
{

    int i= blockIdx.x * blockDim.x + threadIdx.x;
    int j= blockIdx.y * blockDim.y + threadIdx.y;


    if (i>=n || j>=n)
        return;

    A[i][j]+=B[i][j];


}

nvcc告诉我:

./addm.cu(13): Warning: Cannot tell what pointer points to, assuming global memory space
./addm.cu(13): Warning: Cannot tell what pointer points to, assuming global memory space
./addm.cu(13): Warning: Cannot tell what pointer points to, assuming global memory space  

1)我的哲学是正确的吗？2)我认为我理解块,线程等但我不明白是什么

    int i= blockIdx.x * blockDim.x + threadIdx.x;
    int j= blockIdx.y * blockDim.y + threadIdx.y;

不

3)这是一般在2D阵列上执行操作的最有效/最快的方式吗？即不仅仅是矩阵加法,它可以是任何"逐个元素"操作.

4)我可以从matlab调用它吗？通常,当原型是这种形式时,它会变得怪异type** var

多谢你们

Answer 1

你得到的编译器警告来自于在旧GPU上,内存结构不是"平坦"的事实.编译器无法知道内核所使用的指针数组所拥有的地址是什么内存空间.所以它警告你它假设操作正在全局内存中进行.如果您编译Fermi卡的代码(sm_20或sm_21架构),您将看不到该警告,因为这些卡上的内存模型是"平坦的",并且硬件在运行时正确解释了指针.编译器不需要在编译时处理它.

回答你的每个问题:

1. 是.和不.总体思路大约是90%正确,但有几个实施问题将从后面的答案中变得明显.

2. CUDA C内置了变量,允许每个线程确定它正在运行的执行网格中的"坐标",以及每个块的尺寸和它的网格.threadIdx.{xyz}提供块内的线程坐标,blockIdx.{xyz}块与网格坐标.blockDim.{xyz}并分别gridDim.{xyz}提供块和网格的尺寸(注意并非所有硬件都支持3D网格).CUDA使用列主要顺序对每个块中的线程进行编号,并在每个网格中进行块处理.您要查询的计算是{i,j}使用线程和块坐标以及块大小计算2D网格中的等效坐标.这在CUDA编程指南的"编程模型"一章的前几页中有详细讨论.

3. 不,我说这有两个原因.

   首先,在CUDA中使用指针数组进行内存访问并不是一个好主意.指针间接的两个级别极大地增加了获取数据的延迟惩罚.与现代CPU架构相比,典型GPU架构的主要区别在于内存系统.GPU具有惊人的高峰值内存带宽,但访问延迟非常高,而CPU的设计延迟时间最短.因此,必须读取和间接两个指针从内存中获取值是一个非常大的性能损失.将2D数组或矩阵存储在线性存储器中.无论如何,这就是BLAS,LAPACK和Matlab所做的.

   其次,代码中的每个线程都为每个"生产"整数运算(添加)执行四个整数算术运算的设置开销(索引计算).有减少这种情况的策略,通常涉及让每个线程处理多个数组元素.

   如果我要为该操作编写内核,我会像答案底部的代码那样做.这使用线性内存和一维网格.适当占用GPU的线程数适当处理整个输入数组,每个线程处理许多输入.

4. 不,正如我之前在答案中提到的,Matlab使用线性内存来存储矩阵,而不是指针数组.这与您的内核代码所期望的布局不匹配.

示例代码:

__global__ void calc(int *A, int *B, int N)
{

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int s = blockDim.x * gridDim.x;

    for( ; i<N; i+=s) {
        A[i] += B[i];
    }
}