Search code examples
cudatexturestexture-packing

CUDA: how to read 4 (or 16) chars in one transaction per thread using textures and char4 (or int4)?


I have a large character array in the device global memory that is accessed in a coalescent manner by threads. I've read somewhere that I could speed up memory access by reading 4 or 16 chars in one memory transaction per thread. I believe I would have to use textures and the char4 or int4 structs. However, I can't find any documentation or examples on this. Could anyone here please provide a simple example or pointers to where I can learn more about this?

In my code I define the char array as

char *database = NULL;
cudaMalloc( (void**) &database, SIZE * sizeof(char) );

What would the definition be if I want to use textures and char4 (or int4)?

Thanks very much.


Solution

  • I finally figured out the answer to my own question. The definition with char4 would be

    char4 *database = NULL;
    cudaMalloc( (void**) &database, SIZE * sizeof(char4)/4 );
    

    Don't need textures for this. The kernel does speedup by a factor of three with char4 but reduces to two if I do loop unrolling. For the sake of completeness my kernel is

    __global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {
    
      unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
      float A = 0; int i; char ch;
      if(id < jobs_todo) {
        for(i = 0; i < 1000; i += 1){
         ch = database[jobs_todo*i + id];
         if(ch == 'A') A++;
        }
        results[id] = A;
      }
    }
    

    And with char4 it is

    __global__ void kernel4(unsigned int jobs_todo, char4* database, float* results ) {
    
      unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
      float A = 0; int i; char4 ch4;
      if(id < jobs_todo) {
        for(i = 0; i < 1000/4; i += 1){
         ch4 = database[jobs_todo*i + id];
         if(ch4.x == 'A') A++;
         if(ch4.y == 'A') A++;
         if(ch4.z == 'A') A++;
         if(ch4.w == 'A') A++;
        }
        results[id] = A;
      }
    }
    

    I also tried int4 but it's just .0002 seconds faster than the char4 time.