Texture vs target architecture

Accelerated Computing CUDA CUDA Programming and Performance
1

The problem is that it is not possible (means “I cannot”) to bind a texture to a device memory if I compile a code with sm_12 or sm_13 flags.

I cut the code below from the simplePitchLinearTexture SDK example.

If I run this code compiled with sm_10 or sm_11 flags it is OK.

But, if I compile with sm_10 or sm_11 flags it FAILS with cudaError = 18, which means “invalud texture reference”.

Does anyone of you, guys, knows what do I do wrong?

This code just allocates a input data on a host, then allocates two types (linear and 2D array) of device memory, copies input data from host to device (for both types), and then binds textures to device memory (again for both types of memory)

Environment: WinXP sp2, CUDA 2.3

Compilation flags are default for simplePitchLinearTexture SDK sample.

[codebox]#include <stdio.h>

texture<float, 2, cudaReadModeElementType> texRefPL;

texture<float, 2, cudaReadModeElementType> texRefArray;

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

// ----------------------

// Setup

// ----------------------

const int nx = 64;

const int ny = 64;

float h_idata = (float) malloc(sizeof(float)nxny);

for(int i = 0; i < nx*ny; ++i) h_idata[i] = (float) i;

// ----------------------

// Allocate input data

// ----------------------

// linear input data

float *d_idataPL;

size_t d_pitchBytes;

cudaMallocPitch((void**) &d_idataPL, &d_pitchBytes, nx*sizeof(float), ny);

// array input data

cudaArray *d_idataArray;

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc();

cudaMallocArray(&d_idataArray, &channelDesc, nx, ny);

// ------------------------

// copy host data to device

// ------------------------

// linear memory

size_t h_pitchBytes = nx*sizeof(float);

cudaMemcpy2D(d_idataPL, d_pitchBytes, h_idata, h_pitchBytes,

		     nx*sizeof(float), ny, cudaMemcpyHostToDevice);

// array

cudaMemcpyToArray(d_idataArray, 0, 0, h_idata,

			  nx*ny*sizeof(float), cudaMemcpyHostToDevice);

printf(“%d\n”, cudaGetLastError()); // error code = 0

// ----------------------

// Bind texture to memory

// ----------------------

// linear memory

cudaBindTexture2D(0, &texRefPL, d_idataPL, &channelDesc, nx, ny, d_pitchBytes);

printf(“%d\n”, cudaGetLastError()); // in case of sm_12 (or sm_13) error code = 18, which means “invalud texture reference”

// array

cudaBindTextureToArray(texRefArray, d_idataArray, channelDesc);

printf(“%d\n”, cudaGetLastError()); // in case of sm_12 (or sm_13) error code = 18, which means “invalud texture reference”

// ----------------------

// Cleanup

// ----------------------

free(h_idata);

cudaUnbindTexture(texRefPL);

cudaUnbindTexture(texRefArray);

cudaFree(d_idataPL);

cudaFreeArray(d_idataArray);

return 0;

}

[/codebox]