2

I am not a CUDA beginner; I've written a handful of CUDA methods for processing radar data (2D). I've read through most of the CUDA Programming Guide, CUDA By Example, and lots of posts here (Thank you stackoverflow contributors).

I mostly use pitched-linear memory. I've recently gotten into textures, and enjoyed the speed-up.

My question is: how do I get cudaFilterModeLinear to work with a texture object based on a signed 16-bit short?

Minimal reproducible code:

#include <helper_cuda.h> // checkCudaErrors

int main(int argc, char * argv[]){

const unsigned int Nr = 4096;
const unsigned int Na = 1024;

void * ptr;
size_t pitch;
const size_t width  = Nr*sizeof( short );
const size_t height = Na;

checkCudaErrors( cudaMallocPitch( &ptr, &pitch, width, height) );

struct cudaResourceDesc resDesc;
struct cudaTextureDesc texDesc;
cudaTextureObject_t FrameTex;

memset( &resDesc, 0, sizeof(resDesc) );
resDesc.resType = cudaResourceTypePitch2D;
//resDesc.res.pitch2D.desc = cudaCreateChannelDesc<float>();
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<short>();

resDesc.res.pitch2D.devPtr       = ptr;
resDesc.res.pitch2D.pitchInBytes = pitch;
resDesc.res.pitch2D.width        = width;
resDesc.res.pitch2D.height       = height;

// Specify texture object parameters
memset( &texDesc, 0, sizeof(texDesc) );
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
// filter modes: Point, Linear
texDesc.filterMode     = cudaFilterModeLinear;
// read modes: NormalizedFloat, ElementType
texDesc.readMode       = cudaReadModeElementType;
texDesc.normalizedCoords = 0;

// Create texture object
checkCudaErrors( cudaCreateTextureObject( &FrameTex, &resDesc, &texDesc, NULL ));

cudaDeviceReset();
return 0;
}

This will throw

 CUDA error at ... code=26(cudaErrorInvalidFilterSetting) "cudaCreateTextureObject( &FrameTex, &resDesc, &texDesc, NULL )"

The bottom of page 42 of CUDA Programming Guide v8.0 says "Linear texture filtering may be done only for textures that are configured to return floating-point data."

I have no problem if the return value is floating point. But how to base the texture on a 16-bit short?

This post demonstrates cudaFilterModeLinear with uchar, so surely it must be possible. The difference is the code is for Texture References, whereas I desire Texture Objects.

Community
  • 1
  • 1
Tyson Hilmer
  • 741
  • 7
  • 25

1 Answers1

4

This post has the answer.

Although it was not immediately clear on searching through the posts, since the title does not indicate the issue is related to integral texel types.

Specifically, if the texel is integral type, then cudaReadModeNormalizedFloat must be used in conjunction with cudaFilterModeLinear. That is the implied meaning of "... configured to return floating-point data" from the Programming Guide. Geez, why can't Nvidia explicitely state that?

AFAIK there is no dependency on the addressMode or normalizedCoords enumerations.

Edit: cudaReadModeNormalizedFloat does not work with int32 texels. More generally, there is no hardware interpolation for int32.

Community
  • 1
  • 1
Tyson Hilmer
  • 741
  • 7
  • 25