CUDA 텍스처의 다른 어드레싱 모드

Question

국경 주소 지정 모드 (cudaAddressModeBorder)에서 CUDA 텍스처를 사용하고 있습니다. tex2D<float>()을 사용하여 텍스처 좌표를 읽습니다. 텍스처 좌표가 텍스처 외부에있을 때 tex2D<float>()은 0을 반환합니다.

어떻게이 반환 된 테두리 값을 0에서 다른 것으로 변경할 수 있습니까? 텍스처 좌표를 수동으로 확인하고 테두리 값을 직접 설정할 수 있습니다. CUDA API와 같은 경계 값을 설정할 수 있는지 궁금합니다.

Answer 1

는 CUDA 즉 네 개의 비 맞춤형 주소 모드, 3.2.11.1 절에서 설명 클램프, 테두리, 랩와 미러을 지원한다. CUDA 프로그래밍 가이드

앞의 두 개는 정규화되지 않은 좌표와 정규화 된 좌표에서 모두 작동하지만 후자의 두 개는 정규화 된 좌표에서만 작동합니다.

첫 번째 두 가지를 설명하려면 단순화를 위해 비표준 좌표의 경우를 고려하고 1D 신호를 고려해 보겠습니다. 이 경우 입력 시퀀스는 이고 k=0,...,M-1입니다.

cudaAddressModeClamp 신호 c[k]는 k < 0 위해 이렇게 k=0,...,M-1c[k] = c[0] 그 밖에 지속하고 k >= M위한 c[k] = c[M-1]된다

.

cudaAddressModeBorder 신호 c[k]는 k < 0과 k >= M 밖에 k=0,...,M-1되도록 c[k] = 0 계속

.

이제 마지막 두 주소 모드를 설명하기 위해 정규화 된 좌표를 고려해야하므로 1D 입력 신호 샘플은 k=0,...,M-1 인 c[k/M]으로 가정합니다.

cudaAddressModeWrap 신호 c[k/M] 그것이 M 동일한 주기로 주기적이되도록 외부 k=0,...,M-1을 계속

. 즉, 임의의 (양수, 음수 또는 소실) 정수 p에 대해서는 c[(k + p * M)/M] = c[k/M]입니다.

cudaAddressModeMirror 신호 c[k/M] 그것이 2 * M - 2 동일한 주기로 주기적이되도록 외부 k=0,...,M-1을 계속

.즉, l 및 k과 같은 c[l/M] = c[k/M]은 (l + k)mod(2 * M - 2) = 0입니다.

다음 코드는 4 개 개의 모든 가능한 주소 모드

#include <stdio.h> 

texture<float, 1, cudaReadModeElementType> texture_clamp; 
texture<float, 1, cudaReadModeElementType> texture_border; 
texture<float, 1, cudaReadModeElementType> texture_wrap; 
texture<float, 1, cudaReadModeElementType> texture_mirror; 

/********************/ 
/* CUDA ERROR CHECK */ 
/********************/ 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

/******************************/ 
/* CUDA ADDRESS MODE CLAMPING */ 
/******************************/ 
__global__ void Test_texture_clamping(const int M) { 

    printf("Texture clamping - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_clamp, -(float)threadIdx.x)); 
    printf("Texture clamping - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_clamp, (float)(M + threadIdx.x))); 

} 

/****************************/ 
/* CUDA ADDRESS MODE BORDER */ 
/****************************/ 
__global__ void Test_texture_border(const int M) { 

    printf("Texture border - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_border, -(float)threadIdx.x)); 
    printf("Texture border - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_border, (float)(M + threadIdx.x))); 

} 

/**************************/ 
/* CUDA ADDRESS MODE WRAP */ 
/**************************/ 
__global__ void Test_texture_wrap(const int M) { 

    printf("Texture wrap - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_wrap, -(float)threadIdx.x/(float)M)); 
    printf("Texture wrap - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_wrap, (float)(M + threadIdx.x)/(float)M)); 

} 

/****************************/ 
/* CUDA ADDRESS MODE MIRROR */ 
/****************************/ 
__global__ void Test_texture_mirror(const int M) { 

    printf("Texture mirror - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_mirror, -(float)threadIdx.x/(float)M)); 
    printf("Texture mirror - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_mirror, (float)(M + threadIdx.x)/(float)M)); 

} 

/********/ 
/* MAIN */ 
/********/ 
void main(){ 

    const int M = 4; 

    // --- Host side memory allocation and initialization 
    float *h_data = (float*)malloc(M * sizeof(float)); 

    for (int i=0; i<M; i++) h_data[i] = (float)i; 

    // --- Texture clamping 
    cudaArray* d_data_clamping = NULL; gpuErrchk(cudaMallocArray(&d_data_clamping, &texture_clamp.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_clamping, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_clamp, d_data_clamping); 
    texture_clamp.normalized = false; 
    texture_clamp.addressMode[0] = cudaAddressModeClamp; 

    dim3 dimBlock(2 * M, 1); dim3 dimGrid(1, 1); 
    Test_texture_clamping<<<dimGrid,dimBlock>>>(M); 

    printf("\n\n\n"); 

    // --- Texture border 
    cudaArray* d_data_border = NULL; gpuErrchk(cudaMallocArray(&d_data_border, &texture_border.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_border, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_border, d_data_border); 
    texture_border.normalized = false; 
    texture_border.addressMode[0] = cudaAddressModeBorder; 

    Test_texture_border<<<dimGrid,dimBlock>>>(M); 

    printf("\n\n\n"); 

    // --- Texture wrap 
    cudaArray* d_data_wrap = NULL; gpuErrchk(cudaMallocArray(&d_data_wrap, &texture_wrap.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_wrap, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_wrap, d_data_wrap); 
    texture_wrap.normalized = true; 
    texture_wrap.addressMode[0] = cudaAddressModeWrap; 

    Test_texture_wrap<<<dimGrid,dimBlock>>>(M); 

    printf("\n\n\n"); 

    // --- Texture mirror 
    cudaArray* d_data_mirror = NULL; gpuErrchk(cudaMallocArray(&d_data_mirror, &texture_mirror.channelDesc, M, 1)); 
    gpuErrchk(cudaMemcpyToArray(d_data_mirror, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); 
    cudaBindTextureToArray(texture_mirror, d_data_mirror); 
    texture_mirror.normalized = true ; 
    texture_mirror.addressMode[0] = cudaAddressModeMirror; 

    Test_texture_mirror<<<dimGrid,dimBlock>>>(M); 

    printf("\n\n\n"); 
} 

그

있는 출력

index     -7 -6 -5 -4 -3 -2 -1 0 1 2 3 4 5 6 7 8 9 10 11 
clamp     0 0 0 0 0 0 0 0 1 2 3 3 3 3 3 3 3 3 3 
border     0 0 0 0 0 0 0 0 1 2 3 0 0 0 0 0 0 0 0 
wrap     1 2 3 0 1 2 3 0 1 2 3 0 1 2 3 0 1 2 3 
mirror     1 2 3 3 2 1 0 0 1 2 3 3 2 1 0 0 1 2 3 

Answer 2

현재 (CUDA 5.5)부터는 CUDA 텍스처 가져 오기 동작을 사용자 지정할 수 없습니다. 오직 4 자동 기본 모드 1 (즉 테두리, 클램프, 랩와 미러은 ) 범위 질감 중 이용 될 수 페치.

sgarizvi 의해 바와 같이