CUDA Textures with structs

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I’m working on an object detector following the Viola & Jones approach by using Haar cascades.

At the moment, the Haar cascades are implemented as huge arrays containing different structs for stages, classifiers and features.

I use the global memory, upload the structs to the device and give the pointer as parameter to the kernels:

// Host (Upload)

cudaMalloc((void**)&deviceFeatures, sizeof(cudaFeature) * totalFeatures);

cudaMemcpy(deviceFeatures, features, sizeof(cudaFeature) * totalFeatures, cudaMemcpyHostToDevice);

checkCudaError("uploading feature array");

// Device

cudaFeature feature = deviceFeatures + offset;

Due to performance reasons, I’m thinking about using textures for the structs:

texture<cudaFeature, 1, cudaReadModeElementType> featureTexture;

// Host (Upload/Binding)

cudaMalloc((void**)&deviceFeatures, sizeof(cudaFeature) * totalFeatures);

cudaMemcpy(deviceFeatures, features, sizeof(cudaFeature) * totalFeatures, cudaMemcpyHostToDevice);

cudaBindTexture(0, featureTexture, deviceFeatures);

checkCudaError("uploading feature array");

// Device

cudaFeature feature = tex1Dfetch(featureTexture, offset);

Unfortunately, the tex1Dfetch methods only support basic types like int, uchar, float etc… and I got the following compiler error:

Is there any possibility to access structs by using cuda textures?

Maybe my approach is completely wrong and somebody has a better solution for my needs!?

Many thanks in advice,

Daniel

2

Hi!

I’m afraid it is just like you said. Since texturing is “cast” into hardware, there’s no way to texture from types different from those implemented in CUDA (and the built-in vector types).

If your arrays aren’t too large you will might be able to copy them into constant memory which is read-only but fully cached.

3

BTW: Is it possible to allocate constant memory dynamically?

4

I’ve tried to use the constant memory, but my arrays are too large. :(

Maybe this is a dirty hack, but what’s about casting my structs for example into uchar which I can bind to a texture.
I know the size of my structs, so on the device I could read out the corresponding number of uchar’s and cast them back into my struct!

Is there a way to split up my structs into a char array which I can bind to a texture?

5

Use a union:

[codebox]union Combi{

  YourStruct struct;

  char bytes;

};[/codebox]

You can assign a value to the struct part of the union and when you read bytes it will contain the byte representation of the struct. Works the other way around as well.

Hellraiser

6

Could you give an example of how to do this?