matlab CUDA 内核中二维数组的操作

假设我有以下序列号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 问题并按照以下方式使用 2D 线程块CUDA 内核 - 嵌套 for 循环

考虑到这一点，我开始像这样编写我的 cuda 内核：

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

does

3）这是否是在 2D 数组上执行操作的最有效/最快的方法？即不仅仅是矩阵加法，它可以是任何“逐个元素”运算。

4）我可以从matlab中调用它吗？通常当原型是以下形式时它会很奇怪type** var

多谢你们

您收到的编译器警告来自以下事实：在较旧的 GPU 上，内存结构不是“扁平”的。编译器无法知道内核正在工作的指针数组所保存的地址是什么内存空间。因此它警告您，它假设该操作正在全局内存中执行。如果您编译 Fermi 卡（sm_20 或 sm_21 架构）的代码，您将不会看到该警告，因为这些卡上的内存模型是“平坦”的，并且硬件在运行时可以正确解释指针。编译器不需要在编译时处理它。

回答您的每个问题：

1. 是的。和不。总体想法大约 90% 正确，但有几个实施问题将从下面的答案中变得显而易见。

2. CUDA C 具有内置变量，允许每个线程确定其在其运行的执行网格中的“坐标”，以及每个块和网格本身的尺寸。threadIdx.{xyz}提供块内的线程坐标，以及blockIdx.{xyz}块与网格的坐标。blockDim.{xyz} and gridDim.{xyz}分别提供块和网格的尺寸（注意并非所有硬件都支持 3D 网格）。 CUDA用途栏目主要顺序用于对每个块内的线程以及每个网格内的块进行编号。您正在查询的计算正在计算等价的{i,j}使用线程和块坐标以及块大小在 2D 网格中进行坐标。 CUDA 编程指南的“编程模型”一章的前几页对此进行了详细讨论。

3. 不，我这么说有两个原因。

   首先，在 CUDA 中使用指针数组进行内存访问并不是一个好主意。两级指针间接寻址极大地增加了获取数据的延迟损失。与现代 CPU 架构相比，典型 GPU 架构的主要区别在于内存系统。 GPU 具有惊人的高峰值内存带宽，但访问延迟非常高，而 CPU 的设计目标是最小延迟。所以必须阅读和间接两个指针从内存中获取值会带来很大的性能损失。将二维数组或矩阵存储在线性存储器中。这就是 BLAS、LAPACK 和 Matlab 所做的事情。

   其次，代码中的每个线程都为每个“生产性”整数运算（加法）执行四个整数算术运算的设置开销（索引计算）。有一些策略可以减少这种情况，通常涉及让每个线程处理多个数组元素。

   如果我要为该操作编写一个内核，我会执行类似于答案底部的代码的操作。这使用线性存储器和1D grid。适当数量的线程可以正确占用 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];
    }
}