title: setting up a CUDA 2D "unsigned char" texture for linear interpolation
body: <p>I have a linear array of unsigned chars representing a 2D array.  I would like to place it into a CUDA 2D texture and perform (floating point) linear interpolation on it, i.e., have the texture call fetch the 4 nearest unsigned char neighbors, internally convert them to float, interpolate between them, and return the resulting floating point value.</p>
  
  <p>I am having some difficulty setting up the texture and binding it to a texture reference.  I have been through the CUDA reference manual &amp; appendices, but I'm just not having any luck.</p>
  
  <p>Below is runnable code to set up and bind 1) a floating point texture and 2) an unsigned char texture.  The floating point code runs just fine.  However, if you uncomment the two commented unsigned char lines toward the bottom, an "invalid argument" error is thrown.</p>
  
  <pre><code>#include &lt;cstdio&gt;
  #include &lt;cuda_runtime.h&gt;
  
  typedef unsigned char uchar;
  
  // Define (global) texture references; must use "cudaReadModeNormalizedFloat"
  // for ordinal textures
  texture&lt;float, cudaTextureType2D, cudaReadModeNormalizedFloat&gt; texRefFloat;
  texture&lt;uchar, cudaTextureType2D, cudaReadModeNormalizedFloat&gt; texRefUChar;
  
  // Define size of (row major) textures
  size_t const WIDTH  = 1000;
  size_t const HEIGHT = 1000;
  size_t const TOT_PIX = WIDTH*HEIGHT;
  
  int main(void)
  {
     // Set texel formats
     cudaChannelFormatDesc descFloat = cudaCreateChannelDesc&lt;float&gt;();
     cudaChannelFormatDesc descUChar = cudaCreateChannelDesc&lt;uchar&gt;();
  
     // Choose to perform texture 2D linear interpolation
     texRefFloat.filterMode = cudaFilterModeLinear;
     texRefUChar.filterMode = cudaFilterModeLinear;
  
     // Allocate texture device memory
     float * d_buffFloat; cudaMalloc(&amp;d_buffFloat, sizeof(float)*TOT_PIX);
     uchar * d_buffUChar; cudaMalloc(&amp;d_buffUChar, sizeof(uchar)*TOT_PIX);
  
     // Bind texture references to textures
     cudaError_t errFloat = cudaSuccess;
     cudaError_t errUChar = cudaSuccess;
  
     errFloat = cudaBindTexture2D(0, texRefFloat, d_buffFloat, descFloat,
                    WIDTH, HEIGHT, sizeof(float)*WIDTH);
     // Uncomment the following two lines for an error
     //errUChar = cudaBindTexture2D(0, texRefUChar, d_buffUChar, descUChar,
     //               WIDTH, HEIGHT, sizeof(uchar)*WIDTH);
  
     // Check for errors during binding
     if (errFloat != cudaSuccess)
     {
        printf("Error binding float texture reference: %s\n",
            cudaGetErrorString(errFloat));
        exit(-1);
     }
  
     if (errUChar != cudaSuccess)
     {
        printf("Error binding unsigned char texture reference: %s\n",
            cudaGetErrorString(errUChar));
        exit(-1);
     }
  
     return 0;
  }
  </code></pre>
  
  <p>Any help/insight would be most appreciated!</p>
  
  <p>Aaron</p>