cudacreatetexture.cc

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

#include "cuda_runtime.h"
#include "helper_cuda.h" // from cuda_samples

// Copy of texture descriptor from CUDA-11.7, so we can build the sample 
// in a way that simulates compilation with an older CUDA version.
struct __device_builtin__ cudaTextureDesc_117
{
    /**
     * Texture address mode for up to 3 dimensions
     */
    enum cudaTextureAddressMode addressMode[3];
    /**
     * Texture filter mode
     */
    enum cudaTextureFilterMode  filterMode;
    /**
     * Texture read mode
     */
    enum cudaTextureReadMode    readMode;
    /**
     * Perform sRGB->linear conversion during texture read
     */
    int                         sRGB;
    /**
     * Texture Border Color
     */
    float                       borderColor[4];
    /**
     * Indicates whether texture reads are normalized or not
     */
    int                         normalizedCoords;
    /**
     * Limit to the anisotropy ratio
     */
    unsigned int                maxAnisotropy;
    /**
     * Mipmap filter mode
     */
    enum cudaTextureFilterMode  mipmapFilterMode;
    /**
     * Offset applied to the supplied mipmap level
     */
    float                       mipmapLevelBias;
    /**
     * Lower end of the mipmap level range to clamp access to
     */
    float                       minMipmapLevelClamp;
    /**
     * Upper end of the mipmap level range to clamp access to
     */
    float                       maxMipmapLevelClamp;
    /**
     * Disable any trilinear filtering optimizations.
     */
    int                         disableTrilinearOptimization;

    // New field was added by CUDA-11.7. It will be garbage if descriptor is 
    // supplied by the code built with cuda before 11.7
    int poison;  
};


void runTest(int argc, char **argv) {
  // Allocate device memory for result
  int size = 15;
  float *dData = NULL;
  checkCudaErrors(cudaMalloc((void **)&dData, size*sizeof(float)));
  
  cudaTextureObject_t tex;
  cudaResourceDesc texRes;
  memset(&texRes, 0, sizeof(cudaResourceDesc));

  texRes.resType = cudaResourceTypeLinear;
  texRes.res.linear.devPtr = dData;
  texRes.res.linear.sizeInBytes = sizeof(float) * size;
  texRes.res.linear.desc = cudaCreateChannelDesc<float>();

  cudaTextureDesc_old texDescr;
  memset(&texDescr, 0, sizeof(cudaTextureDesc));

  texDescr.filterMode = cudaFilterModePoint;
  texDescr.addressMode[0] = cudaAddressModeBorder;
  texDescr.addressMode[1] = cudaAddressModeBorder;
  texDescr.addressMode[2] = cudaAddressModeBorder;
  texDescr.readMode = cudaReadModeElementType;

  // pretend that descriptor was created by an older CUDA version which didn't 
  // know anything about the new field and it ended up carrying some garbage.
  texDescr.poison = -1;

  checkCudaErrors(cudaCreateTextureObject(&tex, &texRes, (cudaTextureDesc*)&texDescr, NULL));

  checkCudaErrors(cudaDestroyTextureObject(tex));
  checkCudaErrors(cudaFree(dData));
}

int main(int argc, char **argv) {
  runTest(argc, argv); // Will print an error if texture creation failed.
}