0

CUDA 配列にバインドされた CUDA テクスチャ オブジェクトからの正規化されていない線形補間は、正しくない結果を返すようです。補間された値は、0.5予想よりも小さい係数になっているようです。正規化された線形補間は適切に機能しているようです。

このコードに何か問題がありますか? 正規化されていないテクスチャ補間を行う場合、2 倍する必要がありますか?

コード：

#include <iostream>
#include <cstdio>

// simple function to print an array
template <typename T>
void print_array(const T *a, const size_t length) {
  for (size_t i=0; i!=length; i++) {
    std::cout << "a[" << i << "]: " << a[i] << std::endl;
  }
}

// attempt to interpolate linear memory
__global__
void cuda_texture_interpolate(cudaTextureObject_t tex,
                              float start,
                              float stop,
                              int count) {
  if (count < 1) { count = 1; }
  float h = (stop-start)/((float)count);
  float x = start;
  float y;
  for (int i = 0; i != count; i++) {
    y = tex1D<float>(tex,x);
    printf("x: %4g ; y: %4g\n",x,y);
    x = x + h;
  }
  y = tex1D<float>(tex,x);
  printf("x: %4g ; y: %4g\n",x,y);
}

int main(void) {
  // set up host array
  int n = 5;
  float a_host[5] = {3,2,1,2,3};
  printf("printing array on host.\n");
  print_array(a_host,n);

  // allocate and copy to cuda array
  cudaChannelFormatDesc channelDesc =
      cudaCreateChannelDesc(32, 0, 0, 0,
                            cudaChannelFormatKindFloat);
  cudaArray* cuArray;
  cudaMallocArray(&cuArray, &channelDesc, n);

  // Copy to device memory some data located at address h_data
  // in host memory
  cudaMemcpyToArray(cuArray, 0, 0, a_host, n*sizeof(float),
                    cudaMemcpyHostToDevice);

  // create texture object
  cudaResourceDesc resDesc;
  memset(&resDesc, 0, sizeof(resDesc));
  resDesc.resType = cudaResourceTypeArray;
  resDesc.res.array.array = cuArray;

  cudaTextureDesc texDesc;
  memset(&texDesc, 0, sizeof(texDesc));
  texDesc.addressMode[0]   = cudaAddressModeClamp;
  texDesc.filterMode       = cudaFilterModeLinear;
  texDesc.readMode         = cudaReadModeElementType;
  //texDesc.normalizedCoords = 1;
  texDesc.normalizedCoords = 0;


  cudaResourceViewDesc resViewDesc;
  memset(&resViewDesc, 0, sizeof(resViewDesc));
  resViewDesc.format = cudaResViewFormatFloat1;
  resViewDesc.width = n;

  // create texture object
  cudaTextureObject_t tex;
  cudaCreateTextureObject(&tex, &resDesc, &texDesc, &resViewDesc);

  // call interpolation kernel
  printf("interpolate (f(x) -> y).\n");
  //cuda_texture_interpolate<<<1,1>>>(tex,0.0,1.0,10);
  cuda_texture_interpolate<<<1,1>>>(tex,0.0,5.0,10);

  // clean up
  cudaDestroyTextureObject(tex);
  cudaFreeArray(cuArray);

  printf("end of texture_object_interpolation.\n");
  return 0;
}

結果：

$ ./texture_object_interpolation
printing array on host.
a[0]: 3
a[1]: 2
a[2]: 1
a[3]: 2
a[4]: 3
interpolate (f(x) -> y).
x:    0 ; y:  1.5
x:  0.5 ; y:  1.5
x:    1 ; y: 1.25
x:  1.5 ; y:    1
x:    2 ; y: 0.75
x:  2.5 ; y:  0.5
x:    3 ; y: 0.75
x:  3.5 ; y:    1
x:    4 ; y: 1.25
x:  4.5 ; y:  1.5
x:    5 ; y:  1.5
end of texture_object_interpolation.

上記のコード、makefile、および正規化された補間のコードについては、この要点を参照してください。

4

1 に答える 1