二維數組CUDA 2D紋理CUDA

Question

我想將Nx3數組傳遞給內核，並從紋理內存中讀取數據並寫入第二個數組。這是我用N簡化代碼= 8：

#include <cstdio> 
#include "handle.h" 
using namespace std; 

texture<float,2> tex_w; 

__global__ void kernel(int imax, float(*w)[3], float (*f)[3]) 
{ 
    int i = threadIdx.x; 
    int j = threadIdx.y; 

    if(i<imax) 
     f[i][j] = tex2D(tex_w, i, j); 
} 

void print_to_stdio(int imax, float (*w)[3]) 
{ 
    for (int i=0; i<imax; i++) 
    { 
     printf("%2d %3.6f\t %3.6f\t %3.6f\n",i, w[i][0], w[i][1], w[i][2]); 
    } 
} 

int main(void) 
{ 
    int imax = 8; 
    float (*w)[3]; 
    float (*d_w)[3], (*d_f)[3]; 
    dim3 grid(imax,3); 

    w = (float (*)[3])malloc(imax*3*sizeof(float)); 

    for(int i=0; i<imax; i++) 
    { 
     for(int j=0; j<3; j++) 
     { 
      w[i][j] = i + 0.01f*j; 
     } 
    } 

    cudaMalloc((void**) &d_w, 3*imax*sizeof(float)); 
    cudaMalloc((void**) &d_f, 3*imax*sizeof(float)); 

    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>(); 
    HANDLE_ERROR(cudaBindTexture2D(NULL, tex_w, d_w, desc, imax, 3, sizeof(float)*imax)); 

    cudaMemcpy(d_w, w, 3*imax*sizeof(float), cudaMemcpyHostToDevice); 

    // just use threads for simplicity                 
    kernel<<<1,grid>>>(imax, d_w, d_f); 

    cudaMemcpy(w, d_f, 3*imax*sizeof(float), cudaMemcpyDeviceToHost); 

    cudaUnbindTexture(tex_w); 
    cudaFree(d_w); 
    cudaFree(d_f); 

    print_to_stdio(imax, w); 

    free(w); 
    return 0; 
} 

運行這段代碼，我期望得到：

0 0.000000 0.010000 0.020000 
1 1.000000 1.010000 1.020000 
2 2.000000 2.010000 2.020000 
3 3.000000 3.010000 3.020000 
4 4.000000 4.010000 4.020000 
5 5.000000 5.010000 5.020000 
6 6.000000 6.010000 6.020000 
7 7.000000 7.010000 7.020000 

而是我得到：

0 0.000000 2.020000 5.010000 
1 0.010000 3.000000 5.020000 
2 0.020000 3.010000 6.000000 
3 1.000000 3.020000 6.010000 
4 1.010000 4.000000 6.020000 
5 1.020000 4.010000 7.000000 
6 2.000000 4.020000 7.010000 
7 2.010000 5.000000 7.020000 

我覺得這有什麼與我給予cudaBindTexture2D的音高參數有關，但使用較小的值會導致無效的參數錯誤。

在此先感謝！

Answer 1

在布蘭諾的迴應和更多關於球場如何運作後，我會回答我自己的問題。下面是修改後的代碼：

#include <cstdio> 
#include <iostream> 
#include "handle.cu" 

using namespace std; 

texture<float,2,cudaReadModeElementType> tex_w; 

__global__ void kernel(int imax, float (*f)[3]) 
{ 
    int i = threadIdx.x; 
    int j = threadIdx.y; 
    // width = 3, height = imax                   
    // but we have imax threads in x, 3 in y                
    // therefore height corresponds to x threads (i)              
    // and width corresponds to y threads (j)               
    if(i<imax) 
    { 
     // linear filtering looks between indices              
     f[i][j] = tex2D(tex_w, j+0.5f, i+0.5f); 
    } 
} 

void print_to_stdio(int imax, float (*w)[3]) 
{ 
    for (int i=0; i<imax; i++) 
    { 
     printf("%2d %3.3f %3.3f %3.3f\n",i, w[i][0], w[i][1], w[i][2]); 
    } 
    printf("\n"); 
} 

int main(void) 
{ 
    int imax = 8; 
    float (*w)[3]; 
    float (*d_f)[3], *d_w; 
    dim3 grid(imax,3); 

    w = (float (*)[3])malloc(imax*3*sizeof(float)); 

    for(int i=0; i<imax; i++) 
    { 
     for(int j=0; j<3; j++) 
     { 
      w[i][j] = i + 0.01f*j; 
     } 
    } 

    print_to_stdio(imax, w); 

    size_t pitch; 
    HANDLE_ERROR(cudaMallocPitch((void**)&d_w, &pitch, 3*sizeof(float), imax)); 

    HANDLE_ERROR(cudaMemcpy2D(d_w,    // device destination         
          pitch,   // device pitch (calculated above)      
          w,    // src on host           
          3*sizeof(float), // pitch on src (no padding so just width of row)  
          3*sizeof(float), // width of data in bytes        
          imax,   // height of data          
          cudaMemcpyHostToDevice)); 

    HANDLE_ERROR(cudaBindTexture2D(NULL, tex_w, d_w, tex_w.channelDesc, 3, imax, pitch)); 

    tex_w.normalized = false; // don't use normalized values           
    tex_w.filterMode = cudaFilterModeLinear; 
    tex_w.addressMode[0] = cudaAddressModeClamp; // don't wrap around indices       
    tex_w.addressMode[1] = cudaAddressModeClamp; 

    // d_f will have result array                  
    cudaMalloc(&d_f, 3*imax*sizeof(float)); 

    // just use threads for simplicity                 
    kernel<<<1,grid>>>(imax, d_f); 

    cudaMemcpy(w, d_f, 3*imax*sizeof(float), cudaMemcpyDeviceToHost); 

    cudaUnbindTexture(tex_w); 
    cudaFree(d_w); 
    cudaFree(d_f); 

    print_to_stdio(imax, w); 

    free(w); 
    return 0; 
} 

代替使用的memcpy（），並且具有處理在主機上的間距，使用memcpy2D（）的接受兩者的設備數據和主機的數據的音調參數。由於我們在主機上使用的是簡單分配的數據，因此我的理解是音高只是行寬，或者3 * sizeof（float）。

Answer 2

我可以給你一個完整的解決方案，但你可能不會學到：D 所以這裏有一些提示，也許你可以自己修復其餘的問題。

提示1.
當使用cudaBindTexture2D時，它會請求偏移和音高。兩個參數都有一定的硬件依賴對齊限制。如果您使用cudaMalloc（..），則偏移保證爲0。音高通過使用cudaMallocPitch（..）來檢索。您還需要確保主機內存以相同的方式傾斜，否則您的memcpy將無法按預期工作。

提示2.
瞭解2D中的索引。當訪問W [i] [j]中的元素時，需要知道元素W [i] [j + 1]是內存中的下一個元素，而不是W [i + 1] [j]。

提示3.
使用一維數組並計算您的自身的二維索引。這會給你更好的控制。