13

ボーダー アドレッシング モードで CUDA テクスチャを使用しています ( cudaAddressModeBorder)。を使用してテクスチャ座標を読み取っていtex2D<float>()ます。テクスチャ座標がテクスチャの外にある場合、 をtex2D<float>()返します0。

この返された境界値を別のものに変更するにはどうすればよい0ですか? テクスチャ座標を手動で確認し、境界値を自分で設定できました。そのような境界値を設定できるCUDA APIがあるかどうか疑問に思っていました。

4

2 に答える 2

22

sgarizvi が述べたように、CUDA は、セクション 3.2.11.1 で説明されている、カスタマイズできない 4 つのアドレス モード、つまり、 clamp、border、wrap、およびmirrorのみをサポートします。CUDA プログラミング ガイドの

前者の 2 つは正規化されていない座標と正規化された座標の両方で機能しますが、後者の 2 つは正規化された座標でのみ機能します。

最初の 2 つを説明するために、簡単にするために、正規化されていない座標のケースと 1D 信号を考えてみましょう。この場合、入力シーケンスはc[k], withk=0,...,M-1です。

cudaAddressModeClamp

信号c[k]は外部に続くk=0,...,M-1ため、c[k] = c[0]for k < 0、c[k] = c[M-1]for k >= M.

cudaAddressModeBorder

信号c[k]は外部に続くk=0,...,M-1ため、c[k] = 0fork < 0と for k >= M.

ここで、最後の 2 つのアドレス モードを説明するために、正規化された座標を考慮する必要があるため、1D 入力信号サンプルは と仮定されc[k / M]ますk=0,...,M-1。

cudaAddressModeWrap

信号はに等しい周期で周期的であるようにc[k / M]外部に継続されます。つまり、任意の (正、負、または消失) integer に対して。k=0,...,M-1Mc[(k + p * M) / M] = c[k / M]p

cudaAddressModeMirror

信号はに等しい周期で周期的であるようにc[k / M]外部に継続されます。言い換えれば、そのようなものに対して。k=0,...,M-12 * M - 2c[l / M] = c[k / M]lk(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
于 2014-12-10T22:02:42.880 に答える