CODEWITHSUNDEEP
cudagpu
Thiago Alexandre
Thiago Alexandre

Reputation: 21

How to use arrays of Textures in CUDA

I was trying to use an array of textures in CUDA but I'm getting an "invalid texture reference" error. Anyone has any idea how to implement array of textures?

#include<iostream>
#include<cstdio>
#include<cstdlib>

#define MAX_TEXTURE 134217728
#define MAX_LENGTH 268435456

using namespace std;

texture<char, 1, cudaReadModeElementType> texref[2];

int main(){
// host variable
char *text;
// device variable
char *dev_a;

text = (char*)malloc(MAX_LENGTH*sizeof(char));

// any initial value
for(int i=0; i<MAX_LENGTH; i++)
   text[i] = 'a';

cudaMalloc((void**)&dev_a, MAX_LENGTH*sizeof(char));
cudaMemcpy(dev_a, text, MAX_LENGTH*sizeof(char), cudaMemcpyHostToDevice);

cudaError_t err =  cudaBindTexture(0, texref[0], dev_a, MAX_TEXTURE*sizeof(char));
cout << cudaGetErrorString(err) << endl;

size_t offset = MAX_TEXTURE;
err =  cudaBindTexture(&offset, texref[1], dev_a, MAX_TEXTURE*sizeof(char));
cout << cudaGetErrorString(err) << endl;


cudaFree(dev_a);
free(text);
return 0;
}

Upvotes: 1

Views: 1905

Answers (1)

havogt
havogt

Reputation: 2812

This is not possible.

If available to you (Kepler or Maxwell gpu and cuda 5.0 or later), you should try texture objects (bindless textures).

In my code I faked the array by using a function to access various texture references with something similar to the following (in a header file):

texture<char, 1, cudaReadModeElementType> tex0;
texture<char, 1, cudaReadModeElementType> tex1;
// ...

texture<char, 1, cudaReadModeElementType> getTexture( int id )
{
    if( id == 0 )
        return tex0;
    else if( id == 1 )
        return tex1;
    // ...
}

As far as I know there is no way to avoid listing all the texture<char, 1, cudaReadModeElementType> texX; explicitly and be aware that texture references are declared at file scope.

Probably you can improve my code snippet by using preprocessor macros...

Upvotes: 4

Related Questions