CUDA 는 CUDAArray 의 텍 스 처 를 사용 합 니 다.

http://www.cnblogs.com/apapaxionga/archive/2012/03/28/2421621.html
단순 범례
다음은 실례 를 들 어 보 자 ∼ 여 기 는 CUDA Array 의 texture 로 transpose 동작 을 한다.[프로그램 원본 코드 다운로드]
우선 main () 의 내용 은 다음 과 같다.

void main( int argc, char** argv )

{

    int w    = 1920,

         h    = 1200;

    //Setup test data     unsignedchar  *aSrc = new unsigned char[ w * h ],                   *aRS1 = new unsigned char[ w * h ],                   *aRS2 = new unsigned char[ w * h ];     for( int i =0; i < w * h ; ++ i )        aSrc[i] = i % 256;
    //CPU code    Transpose_CPU( aSrc, aRS1, w, h );
    //GPU Code    Transpose_GPU( aSrc, aRS2, w, h );
    //check     for( int i =0; i < w * h; ++ i )        if( aRS1[i] != aRS2[i] )        {            printf( "Error!!!!" );            break;        } }
마찬가지 로 간단 하 다. 먼저 원시 자 료 를 선고 한 것 이다. aSrc, 그리고 전 치 된 자료 aRS1 화해시키다 aRS 2; 그리고 원본 자료 에서 aSrc 값 을 입력 하 십시오. 1920 * 1200, aRS 1 일 것 같 습 니 다. 화해시키다 aRS2 1200 * 1920 일 것 입 니 다. 그러나 1 차원 배열 을 선언 할 때 차이 가 없 기 때문에 특별히 수정 하지 않 았 습 니 다.)
그 다음은 CPU 버 전과 GPU 버 전의 프로그램 을 각각 달리 고 이들 의 결 과 를 비교 하 는 것 이다 ∼ CPU 버 전의 함 식 Transpose CPU () 내용 은 다음 과 같다.

void Transpose_CPU( unsigned char* sImg, unsigned char *tImg,

                    int w, int h )

{

    int x, y, idx1, idx2;

    for( y = 0; y < h; ++ y )

        for( x = 0; x < w; ++ x )

        {

            idx1 = y * w + x;

            idx2 = x * h + y;

            tImg[idx2] = sImg[idx1];

        }

}

내용 은 더 설명 할 필요 가 없 을 것 같 습 니 다 ∼ 한 마디 로 방향 에 따라 다른 방법 으로 계산 해 내 는 것 입 니 다 idx1 화해시키다 idx 2 두 메모리 공간의 색인 값 으로 자 료 를 sImg 다음으로 복사 tImg, 이 를 통 해 전 치 된 동작 을 합 니 다.
그리고 Transpose_GPU() 있 는. cu 파일 의 내용 은 다음 과 같 습 니 다.

#define BLOCK_DIM 16  texture<unsigned char, 2, cudaReadModeElementType> rT;  extern "C" 

void Transpose_GPU( unsigned char* sImg, unsigned char *tImg,

                             int w, int h );



 

__global__ 

void Transpose_Texture( unsigned char* aRS, int w, int h )

{

     int idxX = blockIdx.x * blockDim.x + threadIdx.x,

          idxY = blockIdx.y * blockDim.y + threadIdx.y;

     if( idxX < w && idxY < h )

         aRS[ idxX * h + idxY ] = tex2D( rT, idxX, idxY );

}

 

void Transpose_GPU( unsigned char* sImg, unsigned char *tImg,

                             int w, int h )

{

     // compute the size of data

     int data_size = sizeof(unsigned char) * w * h;

 

    // part1a. prepare the result data

     unsigned char *dImg;

     cudaMalloc( (void**)&dImg, data_size );

 

     // part1b. prepare the source data

     cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<unsigned char>();

     cudaArray* cuArray;

     cudaMallocArray(&cuArray, &chDesc, w, h);

     cudaMemcpyToArray( cuArray, 0, 0, sImg, data_size,

                                   cudaMemcpyHostToDevice );

     cudaBindTextureToArray( rT, cuArray );

 

     // part2. run kernel

     dim3 block( BLOCK_DIM, BLOCK_DIM ),

            grid( ceil( (float)w / BLOCK_DIM), ceil( (float)h / BLOCK_DIM) );

     Transpose_Texture<<< grid, block>>>( dImg, w, h );

 

     // part3. copy the data from device

     cudaMemcpy( tImg, dImg, data_size, cudaMemcpyDeviceToHost );

 

     // par4. release data

     cudaUnbindTexture( rT );

     cudaFreeArray( cuArray );

     cudaFree( dImg ); } 

우선, 전에 도 말씀 드 렸 듯 이 현재 의 CUDA 는 texture file - scope 에 선언 되 었 기 때문에 처음부터 2D texture 를 입력 자료 로 선언 해 야 합 니 다. 솔직히 Heresy 는 이 점 에 대해 불편 합 니 다.
다음 main() 호출 된 Transpose_GPU() 그 가 한 내용 은 다음 과 같다.
먼저 필요 한 메모리 크기 를 계산 해 보 세 요.
[part1a] 선고 하 다. dImg, 메모리 주 소 를 할당 합 니 다. dImg 계산 한 결 과 를 저장 합 니 다.
[part1b] CUDAarray 구축 cuArray, 메모리 주소 할당, hostmemory (sImg) 에서 devicememory (cuArray) 로 자 료 를 복사 하고 cudaBindTextureToArray () 를 통 해 장차 rT 화해시키다 cuArray 연락 해.
[part2] 커 널 함수 호출: Transpose Texture () 여기 서 threadblock 의 크기 는 BLOCK DIM * BLOCK DIM (16 * 16) 으로 정의 되 며, grid 의 크기 는 너비 와 높이 에 따라 block 의 크기 로 나 뉜 다.
[part3] 결 과 를 device memory (dImg) 에서 hostmemory (tImg) 로 복사 합 니 다.
[part4] 투과 cudaUnbindTexture() 장차 rT 화해시키다 sImg 간 의 연락 이 해제 되 고 cudaFreeArray (), cudaFree () 를 사용 합 니 다. device memory 를 방출 합 니 다.
이 프로그램의 kernefunction Transpose_Texture() 블록 Idx, 블록 Dim, threadIdx 를 직접 통 해 이 세 가지 변 수 는 2 차원 의 위 치 를 계산 하고 x, y 가 범 위 를 초과 하지 않 았 을 때 자 료 를 복사 하여 (idxx, idxY) 의 자료, 투과 tex2D() 꺼 내 서 aRS [idxx * h + idxY] 에 저장 합 니 다.
 
지금까지 CUDA 2D texture 를 사용 하 는 가장 기본 적 인 방법 인 것 같 습 니 다. part.1 시 에서 언급 한 것 처럼 CUDA Array 를 사용 하 는 texture 는 사실 사용 할 수 있 는 추가 기능 이 있 습 니 다! 하 이 레벨 의 사용 외 에 도 low - level, 더 디 테 일 한 기능 을 사용 할 수 있 습 니 다 ∼ 하지만 여 기 는 잠시 언급 하지 않 겠 습 니 다 ∼ 나중에 시간 이 있 으 면 말씀 드 리 겠 습 니 다.