2D 배열의 2D 배열 CUDA

Question

커널에 Nx3 배열을 전달하고 텍스쳐 메모리에서 읽은 다음 두 번째 배열에 쓰려고합니다.

#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; 
} 

이 코드를 실행 내가 얻을 기대 : 여기 = 8 N 내 간단한 코드입니다

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

브 라우 오의 반응과 피치 작동 방식에 대해 자세히 살펴본 후 제 자신의 질문에 답변 해 드리겠습니다.

() 대신 memcpy2D을 사용하여 방어 적이기()를 사용하여 호스트 시스템 피치를 다루는 데

#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; 
} 

장치 데이터 및 호스트 데이터 모두 피치 인수를 받아 다음은 수정 된 코드이다. 우리는 단순히 호스트에 할당 된 데이터를 사용하기 때문에 피치가 단순히 행 너비 또는 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.
1D 어레이를 사용하여 2D 색인을 계산합니다. 이렇게하면 더 잘 제어 할 수 있습니다.