二维数组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.
使用一维数组并计算您的自身的二维索引。这会给你更好的控制。