2D紋理的間距對齊

2D紋理是CUDA在圖像處理應用中的一個有用功能。要將間距線性內存綁定到2D紋理，內存必須對齊。 cudaMallocPitch對齊內存分配是一個很好的選擇。在我的設備上，由cudaMallocPitch返回的音高是512的倍數，即內存是512字節對齊的。2D紋理的間距對齊

設備的實際對齊要求由我的設備上的32字節的cudaDeviceProp::texturePitchAlignment決定。

我的問題是：

如果2D紋理實際對齊要求是32個字節，那麼爲什麼cudaMallocPitch返回512字節對齊的內存？

這不是浪費內存嗎？例如，如果我創建一個尺寸爲513 x 100的8位圖像，它將佔用1024 x 100個字節。

我得到以下系統這種行爲：

1：華碩G53JW +的Windows 8的x64 +的GeForce GTX 460M + CUDA 5 +酷睿i7 740QM + 4GB RAM

2：戴爾Inspiron N5110 +的Windows 7 64 +的GeForce GT525M + CUDA 4.2 + Corei7 2630QM + 6GB RAM

來源

2012-09-23 sgarizvi

這是什麼硬件？我一直髮現cudaMallocPitch會尊重報告的紋理對齊。在我現在可以訪問的唯一設備上，報告的字節對齊方式是256，我總是獲得256個字節的倍數。 – talonmies

我已更新該問題。在問題中添加了詳細的系統配置。 – sgarizvi

這是一個稍微投機的答案，但要記住，有2種取向性，其分配的間距必須滿足的紋理，一個用於textutr指針和一個紋理行。我懷疑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適用於陣列。

來源

2012-09-23 10:35:47 talonmies

我想你是對的。在我的兩個系統上，我得到'cudaDeviceProp :: textureAlignment == 512'。 – sgarizvi

正如talonmies所說，紋理的排列有一個要求（textureAlignment，我記得舊版硬件上的256個字節，當前硬件上512個字節），加上每行對齊要求（texturePitchAlignment）。一般來說，texturePitchAlignment <= textureAlignment。 CUDA返回內存中的malloc函數適合紋理對齊。 – njuffa

在做了一些內存分配的實驗後，終於找到了一個節省內存的工作解決方案。如果我強制對齊由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-09 09:54:09 sgarizvi

2D紋理的間距對齊

回答

相關問題