CUDA设备指针操作

Question

我使用过：

float *devptr;
//...
cudaMalloc(&devptr, sizeofarray);
cudaMemcpy(devptr, hostptr, sizeofarray, cudaMemcpyHostToDevice);

在CUDA C中分配和填充数组。 现在我正在尝试运行一个CUDA内核，例如：

__global__ void kernelname(float *ptr)
{
   //...
}

在数组中加上一个偏移量的话，在C/C++中应该是这样的：

kernelname<<<dimGrid, dimBlock>>>(devptr+offset);

然而，这似乎不起作用。

有没有一种方法可以在不将偏移值作为单独参数发送到内核并在内核代码中使用该偏移的情况下完成此操作？有什么想法吗？

- pmcr

Answer 1

在CUDA中，指针算术运算可以正常工作。您可以在主机代码中为CUDA指针添加偏移量，它将正确地工作（请记住，偏移量不是字节偏移量，而是普通的单词或元素偏移量）。

编辑：一个简单的工作示例：

#include <cstdio>
int main(void)
{

    const int na = 5, nb = 4;
    float a[na] = { 1.2, 3.4, 5.6, 7.8, 9.0 };
    float *_a, b[nb];

    size_t sza = size_t(na) * sizeof(float);
    size_t szb = size_t(nb) * sizeof(float);

    cudaFree(0);

    cudaMalloc((void **)&_a, sza );
    cudaMemcpy( _a, a, sza, cudaMemcpyHostToDevice);
    cudaMemcpy( b, _a+1, szb, cudaMemcpyDeviceToHost);

    for(int i=0; i<nb; i++)
        printf("%d %f\n", i, b[i]);

    cudaThreadExit();
}

在这里，您可以看到对设备指针进行了单词/元素偏移，以便在第二个cudaMemcpy调用中从第二个单词而不是第一个单词开始复制。

- talonmies

---

谢谢你的回答。我知道它不是字节偏移量，但是我不知道什么是字长。我只是像在C++中一样添加了我需要指针的索引号码。这样正确吗？ - pmcr

谢谢您的时间。非常有帮助。现在它可以工作了。这是一个愚蠢的错误。我真的很抱歉。 - pmcr

嗨，talonmies。你所说的“word size offset”是什么意思？这不取决于你要分配的数组类型吗？如果你的示例中的数组是char类型，那么_a+1指向a[1]，而不是a[3]，对吗？ - Luc

问题：如果我在GPU上有一个char *，偏移量是以字节还是以单词（4个字节）为单位？ - XapaJIaMnu

偏移量始终是类型的元素数量，而不是字节数。如果您想按一个数组元素进行偏移，则偏移量为+1，与类型无关。 - talonmies

Answer 2

指针算术可以在主机端代码上工作，在nvidia提供的示例代码中经常使用。

“线性内存存在于设备的40位地址空间中，因此可以通过指针引用单独分配的实体，例如在二叉树中。”

更多信息请阅读：http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz4KialMz00 从性能基元（npp）文档中可以找到完美的指针算术示例。

“4.5.1 选择通道源图像指针 这是指向源图像第一个像素中感兴趣的通道的指针。例如，如果pSrc是三通道图像ROI内第一个像素的指针。使用适当的选择通道复制原语，可以通过将指针偏移一个来将此源图像的第二个通道复制到由pDst给出的目标图像的第一个通道中： nppiCopy_8u_C3CR(pSrc + 1, nSrcStep, pDst, nDstStep, oSizeROI);”

*注意：这可以正常工作而不需要乘以数据元素的字节数，因为编译器知道指针的数据类型，并相应地计算地址。

在C和C ++中，可以通过上述方式或符号 &ptr [offset]（返回数据的设备内存地址而不是值，值不能在主机端代码中使用）完成指针算术。当使用任一符号时，数据类型的大小会自动处理，并且偏移量以数据元素数量而非字节数指定。

- Jonathan Olson