global 内核中棘手的数组算法（CUDA 样本）

Question

我对 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和d_Dst同时。

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

谢谢

Answer 1

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

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

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

global 内核中棘手的数组算法（CUDA 样本）

Tricky array arithmetics inside a global kernel (CUDA samples)

cuda

__global__ 内核中棘手的数组算法（CUDA 样本）

Tricky array arithmetics inside a __global__ kernel (CUDA samples)

cuda

global 内核中棘手的数组算法（CUDA 样本）

Tricky array arithmetics inside a global kernel (CUDA samples)