5

2D 纹理是 CUDA 在图像处理应用中的一个有用特性。要将音高线性内存绑定到 2D 纹理，内存必须对齐。cudaMallocPitch是对齐内存分配的好选择。在我的设备上，返回的音高cudaMallocPitch是 512 的倍数，即内存是 512 字节对齐的。

设备的实际对齐要求cudaDeviceProp::texturePitchAlignment取决于我设备上的 32 字节。

我的问题是：

如果 2D 纹理的实际对齐要求是 32 字节，那么为什么cudaMallocPitch返回 512 字节对齐的内存呢？

这不是浪费内存吗？例如，如果我创建一个大小为 513 x 100 的 8 位图像，它将占用 1024 x 100 字节。

我在以下系统上得到这种行为：

1：华硕 G53JW + Windows 8 x64 + GeForce GTX 460M + CUDA 5 + Core i7 740QM + 4GB RAM

2：戴尔 Inspiron N5110 + Windows 7 x64 + GeForce GT525M + CUDA 4.2 + Corei7 2630QM + 6GB RAM

4

2 回答 2

4

这是一个略微推测的答案，但请记住，分配的间距必须满足纹理的两个对齐属性，一个用于纹理指针，一个用于纹理行。我怀疑这cudaMallocPitch是对前者的尊重，由cudaDeviceProp::textureAlignment. 例如：

#include <cstdio>

int main(void)
{
    const int ncases = 12;
    const size_t widths[ncases] = { 5, 10, 20, 50, 70, 90, 100,
        200, 500, 700, 900, 1000 };
    const size_t height = 10;

    float *vals[ncases];
    size_t pitches[ncases];

    struct cudaDeviceProp p;
    cudaGetDeviceProperties(&p, 0);
    fprintf(stdout, "Texture alignment = %zd bytes\n",
            p.textureAlignment);
    cudaSetDevice(0);
    cudaFree(0); // establish context

    for(int i=0; i<ncases; i++) {
        cudaMallocPitch((void **)&vals[i], &pitches[i], 
            widths[i], height);
        fprintf(stdout, "width = %zd <=> pitch = %zd \n",
                widths[i], pitches[i]);
    }

    return 0;
}

在 GT320M 上给出以下信息：

Texture alignment = 256 bytes
width = 5 <=> pitch = 256 
width = 10 <=> pitch = 256 
width = 20 <=> pitch = 256 
width = 50 <=> pitch = 256 
width = 70 <=> pitch = 256 
width = 90 <=> pitch = 256 
width = 100 <=> pitch = 256 
width = 200 <=> pitch = 256 
width = 500 <=> pitch = 512 
width = 700 <=> pitch = 768 
width = 900 <=> pitch = 1024 
width = 1000 <=> pitch = 1024

我猜这cudaDeviceProp::texturePitchAlignment适用于 CUDA 数组。

于 2012-09-23T10:35:47.517 回答
3

在对内存分配进行了一些实验之后，我终于找到了一个可以节省内存的工作解决方案。如果我强制对齐分配的内存cudaMalloc，则cudaBindTexture2D可以完美运行。

cudaError_t alignedMalloc2D(void** ptr, int width, int height, int* pitch, int alignment = 32)
{       
   if((width% alignment) != 0)
      width+= (alignment - (width % alignment));

   (*pitch) = width;

   return cudaMalloc(ptr,width* height);
}

此函数分配的内存是 32 字节对齐的，这是cudaBindTexture2D. 我的内存使用量现在减少了 16 倍，所有使用 2D 纹理的 CUDA 函数也可以正常工作。

这是一个小实用函数，用于获取当前选择的 CUDA 设备间距对齐要求。

int getCurrentDeviceTexturePitchAlignment()
{
   cudaDeviceProp prop;
   int currentDevice = 0;

   cudaGetDevice(&currentDevice);

   cudaGetDeviceProperties(&prop,currentDevice);

   return prop.texturePitchAlignment;
}
于 2012-10-09T09:54:09.797 回答