2D纹理的倾斜对齐

Question

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

设备的实际对齐要求由 cudaDeviceProp :: texturePitchAlignment 确定，该值在我的设备上为32字节。

我的问题是：

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

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

我在以下系统上获得此行为：

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

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

- sgarizvi

---

这是在什么硬件上？我一直发现cudaMallocPitch遵守报告的纹理对齐方式。在我现在可以访问的唯一设备上，以字节为单位报告的对齐方式为256，并且我总是得到256字节的倍数来进行内存分配。 - talonmies

我已经更新了问题。在问题中添加了详细的系统配置。 - sgarizvi

Answer 1

这是一个稍微推测性的答案，但请记住，对于纹理，分配的音高必须满足两个对齐属性，一个用于纹理指针，另一个用于纹理行。我怀疑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 数组。

- talonmies

---

我想你是对的。在我的两个系统上，我得到了 cudaDeviceProp::textureAlignment == 512。 - sgarizvi

如talonmies所说，纹理对齐有一个要求（textureAlignment），我记得在旧硬件上是256字节，在当前硬件上是512字节，每行还有一个对齐要求（texturePitchAlignment）。通常来说，texturePitchAlignment <= textureAlignment。CUDA中的malloc函数返回适合纹理对齐的内存。 - njuffa

Answer 2

在进行一些内存分配方面的实验后，最终我找到了一个能够节省内存的解决方案。如果我强制对由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;
}

- sgarizvi

---

它按照你所说的运行良好，但它是否遵守了对内存进行合并访问的要求？ - Mohsen