1

複素数（float2）での「標準」CUDA実装と「テクスチャベース」CUDA実装を使用した1D線形補間を比較しています。

「標準」のCUDA実装は、次の行で構成されています。

/*************************************/
/* LINEAR INTERPOLATION KERNEL - GPU */
/*************************************/
__device__ float linear_kernel_GPU(float in)
{
    float d_y;
    return 1.-abs(in);
}

/**********************************************/
/* LINEAR INTERPOLATION KERNEL FUNCTION - GPU */
/**********************************************/
__global__ void linear_interpolation_kernel_function_GPU(float2* result_d, float2* data_d, float* x_in_d, float* x_out_d, int M, int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N)
    {
        result_d[j].x = 0.;
        result_d[j].y = 0.;
        for(int k=0; k<M; k++)
        {
            if (fabs(x_out_d[j]-x_in_d[k])<1.) {
                result_d[j].x = result_d[j].x + linear_kernel_GPU(x_out_d[j]-x_in_d[k])*data_d[k].x;
                result_d[j].y = result_d[j].y + linear_kernel_GPU(x_out_d[j]-x_in_d[k])*data_d[k].y; }
        }  
    } 
}

extern "C" void linear_interpolation_function_GPU(cuComplex* result_d, cuComplex* data_d, float* x_in_d, float* x_out_d, int M, int N){

    dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1),1);
    linear_interpolation_kernel_function_GPU<<<dimGrid,dimBlock>>>(result_d, data_d, x_in_d, x_out_d, M, N);

}

「テクスチャベース」のCUDA実装は、次の行で構成されています。

texture<float2, 1, cudaReadModeElementType> data_d_texture;

// ********************************************************/
// * LINEAR INTERPOLATION KERNEL FUNCTION - GPU - TEXTURE */
// ********************************************************/
__global__ void linear_interpolation_kernel_function_GPU_texture(cuComplex* result_d, float* x_out_d, int M, int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N) result_d[j] = tex1D(data_d_texture,float(x_out_d[j]+M/2+0.5));

}

// *************************************************/
// * LINEAR INTERPOLATION FUNCTION - GPU - TEXTURE */
// *************************************************/
extern "C" void linear_interpolation_function_GPU_texture(float2* result_d, float2* data, float* x_in_d, float* x_out_d, int M, int N){

    cudaArray* data_d = NULL; cudaMallocArray (&data_d, &data_d_texture.channelDesc, M, 1); 
    cudaMemcpyToArray(data_d, 0, 0, data, sizeof(float2)*M, cudaMemcpyHostToDevice); 
    cudaBindTextureToArray(data_d_texture, data_d); 
    data_d_texture.normalized = false; 
    data_d_texture.filterMode = cudaFilterModeLinear;

    dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1),1);
    linear_interpolation_kernel_function_GPU_texture<<<dimGrid,dimBlock>>>(result_d, x_out_d, M, N);

}

「テクスチャベース」の補間は、「標準」の補間よりも20倍以上高速です。ただし、結果に不一致があり、2つの実装間の二乗平均平方根誤差が約0.07%。

CUDA Cプログラミングガイドによると、補間係数は8ビットの分数値を持つ9ビットの固定小数点形式で格納されるため、この不一致の原因となる可能性があります。

次に、2つの質問があります。

1）「テクスチャベース」の補間の精度を高めるための「トリック」はありますか？

2）この9ビット表現は、float4に移動しても、ここで得られる精度に制限されると思いますよね？言い換えれば、float2からfloat4に数値表現の精度を向上させる意味はありませんか？

前もって感謝します。

4

1 に答える 1

2

テクスチャを「事前補間」して解像度を上げることができます。つまり、初期テクスチャが100x100の場合、事前補間して200x200にすることができ、カーネル内補間の解像度が2倍になります。

于 2013-01-22T15:09:21.403 に答える