CUDA 纹理插值对于归一化坐标不正确

Question

我正在尝试在浮点查找表上使用 CUDA 纹理插值。当使用归一化坐标时，结果不正确；当使用非归一化坐标时，它们是正确的。这是为什么呢？

可编译示例：

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

texture<float, cudaTextureType1D, cudaReadModeElementType> table_tex;

const int N_table = 6;

// y = 2*x for x in [0, 1)
float table[N_table] = {0, 0.4, 0.8, 1.2, 1.6, 2.0};

__global__ void hw_linear_interpolation(const float* inputs,
                                        float* interpolated,
                                        const unsigned int n_inputs)
{
    int tid = threadIdx.x;

    if (tid < n_inputs)
    {
        float val = inputs[tid];

#ifdef NORMALIZED
        float interp = tex1D(table_tex, val);
#else
        float interp = tex1D(table_tex, (N_table-1)*val+0.5f);
#endif
        interpolated[tid] = interp;
    }
}

int main(void)
{
    int N_inputs = 11;

    thrust::host_vector<float> h_inputs(N_inputs);
    thrust::device_vector<float> d_outputs(N_inputs);
    thrust::host_vector<float> h_outputs(N_inputs);

    // Allocate CUDA array in device memory to bind table_tex to.
    cudaChannelFormatDesc channelDesc =
                             cudaCreateChannelDesc<float>();
    cudaArray* cuArray_table;
    cudaMallocArray(&cuArray_table, &channelDesc, N_table, 0);

    // Copy to device memory some data located at address h_data
    // in host memory
    cudaMemcpyToArray(cuArray_table, 0, 0, table, N_table*sizeof(float),
                      cudaMemcpyHostToDevice);

    // Initialize input values to interpolate from the table for.
    for (int i=0; i<N_inputs; i++) {
        h_inputs[i] = i*0.1f;
    }

    thrust::device_vector<float> d_inputs = h_inputs;

    // Set up texture for linear interpolation with normalized inputs.
    table_tex.addressMode[0] = cudaAddressModeClamp;
    table_tex.filterMode = cudaFilterModeLinear;
#ifdef NORMALIZED
    table_tex.normalized = true;
#else
    table_tex.normalized = false;
#endif

    cudaBindTextureToArray(table_tex, cuArray_table);
    hw_linear_interpolation<<<1, 128>>>(
        thrust::raw_pointer_cast(d_inputs.data()),
        thrust::raw_pointer_cast(d_outputs.data()),
        N_inputs);
    cudaUnbindTexture(table_tex);
    h_outputs = d_outputs;

    std::cout << "     x     |   interp. y   |   actual y  ";
    std::cout << std::endl;
    std::cout << "-----------------------------------------";
    std::cout << std::endl;

    std::cout.setf(std::ios::fixed, std::ios::floatfield);
    for (int i=0; i<N_inputs; i++)
    {
        std::cout << "    ";
        std::cout.precision(1);
        std::cout.width(3);
        std::cout << h_inputs[i];
        std::cout << "    |";

        std::cout << "    ";
        std::cout.precision(5);
        std::cout.width(7);
        std::cout << h_outputs[i];
        std::cout << "    |";

        std::cout << "   ";
        std::cout.width(7);
        std::cout << 2*(i*0.1f);
        std::cout << std::endl;
    }

    return 0;
}

编译为nvcc -arch=sm_20 interpolation_so.cu 给出

$ ./a.out 
     x     |   interp. y   |   actual y  
-----------------------------------------
    0.0    |    0.00000    |   0.00000
    0.1    |    0.20000    |   0.20000
    0.2    |    0.40000    |   0.40000
    0.3    |    0.60000    |   0.60000
    0.4    |    0.80000    |   0.80000
    0.5    |    1.00000    |   1.00000
    0.6    |    1.20000    |   1.20000
    0.7    |    1.40000    |   1.40000
    0.8    |    1.60000    |   1.60000
    0.9    |    1.80000    |   1.80000
    1.0    |    2.00000    |   2.00000

但编译为nvcc -arch=sm_20 interpolation_so.cu -DNORMALIZED 给出

$ ./a.out
     x     |   interp. y   |   actual y  
-----------------------------------------
    0.0    |    0.00000    |   0.00000
    0.1    |    0.04063    |   0.20000
    0.2    |    0.27969    |   0.40000
    0.3    |    0.52031    |   0.60000
    0.4    |    0.75938    |   0.80000
    0.5    |    1.00000    |   1.00000
    0.6    |    1.24063    |   1.20000
    0.7    |    1.47969    |   1.40000
    0.8    |    1.72031    |   1.60000
    0.9    |    1.95938    |   1.80000
    1.0    |    2.00000    |   2.00000

编辑：基于Dithermaster's answer的部分修复

Answer 1

因为你在这里加了 0.5:tex1D(table_tex, val+0.5f); 归一化坐标从 0.0 到 1.0，因此将 0.5 偏移量添加到整个范围的一半。