__global__ 内核中的棘手数组算法（CUDA 示例）答案

【问题标题】：Tricky array arithmetics inside a __global__ kernel (CUDA samples)__global__ 内核中的棘手数组算法（CUDA 示例）
【发布时间】：2015-01-30 06:09:35
【问题描述】：

我对 CUDA 示例 "CUDA Separable Convolution" 中的代码有疑问。为了进行行卷积，此代码首先将数据加载到共享内存中。使用指针算法，每个线程将输入指针移动到它们自己的位置，然后将一些全局内存写入共享内存。这是让我感到困惑的一段代码：

__global__ void convolutionRowsKernel(
    float *d_Dst,
    float *d_Src,
    int imageW,
    int imageH,
    int pitch
    )
{
    __shared__ float s_Data[ROWS_BLOCKDIM_Y][(ROWS_RESULT_STEPS + 2 *     ROWS_HALO_STEPS) * ROWS_BLOCKDIM_X];

    //Offset to the left halo edge
    const int baseX = (blockIdx.x * ROWS_RESULT_STEPS - ROWS_HALO_STEPS) * ROWS_BLOCKDIM_X + threadIdx.x;
    const int baseY = blockIdx.y * ROWS_BLOCKDIM_Y + threadIdx.y;

    d_Src += baseY * pitch + baseX;
    d_Dst += baseY * pitch + baseX;

    //Load main data
#pragma unroll

    for (int i = ROWS_HALO_STEPS; i < ROWS_HALO_STEPS + ROWS_RESULT_STEPS; i++)
    {
        s_Data[threadIdx.y][threadIdx.x + i * ROWS_BLOCKDIM_X] = d_Src[i * ROWS_BLOCKDIM_X];
    }
...

据我了解这段代码，每个线程都会计算自己的baseX和baseY的值，之后所有活动线程将开始增加指针d_Src和@ 987654325@同时。

所以，据我所知，如果数组 d_Src 和 d_Dst 在本地内存中（例如，每个线程都有自己的数组副本），这将是正确的。但是这个数组在全局设备内存中！那么会发生什么呢，所有活动的线程都会增加指针，结果就会不正确。有人可以解释一下，为什么这有效？

谢谢

【问题讨论】：

标签： cuda

【解决方案1】：

之所以有效，是因为每个线程都有自己的指针副本。

void foo(float* bar){
    bar++;
}

float* test = 0;
foo(test);
cout<<test<<endl; //will print 0

【讨论】：

那么，是不是真的，任何输入指针都会自动放入本地内存，例如每个线程都有自己的副本，但是指向的数据指针在全局设备内存中，就像在内核调用之前一样？
是的，数据（指针指向的地方）仍在全局内存中。您的问题实际上与 cuda 或全局内存无关，它只是指针在 c 中的工作方式。