How to properly cast a global memory array using the uint4 vector in CUDA to increase memory throughput?

c, cuda, optimization

Solution

If I have understood what you are trying to do, the logical approach is to use the C++ `reinterpret_cast` mechanism to make the compiler generate the correct vector load instruction, then use the CUDA built in byte sized vector type `uchar4` to access each byte within each of the four 32 bit words loaded from global memory. Using this approach, you are really putting trust in the compiler knowing the optimal way to do byte wise access within each 32 bit register.

A completely contrived example could look like this:

#include <cstdio>
#include <cstdlib>

__global__
void kernel(unsigned int *in, unsigned char* out)
{
    int tid = threadIdx.x;

    uint4* p = reinterpret_cast<uint4*>(in);
    uint4  i4 = p[tid]; // vector load here

    uchar4 c0 = *reinterpret_cast<uchar4 *>(&i4.x);
    uchar4 c4 = *reinterpret_cast<uchar4 *>(&i4.y);
    uchar4 c8 = *reinterpret_cast<uchar4 *>(&i4.z);
    uchar4 c12 = *reinterpret_cast<uchar4 *>(&i4.w);

    out[tid*4+0] = c0.x;
    out[tid*4+1] = c4.y;
    out[tid*4+2] = c8.z;
    out[tid*4+3] = c12.w;
}

int main(void)
{
    unsigned int c[8] = { 
        2021161062, 2021158776, 2020964472, 1920497784, 
        2021161058, 2021161336, 2020898936, 1702393976 };

    unsigned int * _c;
    cudaMalloc((void **)&_c, sizeof(int)*size_t(8));
    cudaMemcpy(_c, c, sizeof(int)*size_t(8), cudaMemcpyHostToDevice);
    unsigned char * _m;
    cudaMalloc((void **)&_m, sizeof(unsigned char)*size_t(8));

    kernel<<<1,2>>>(_c, _m);

    unsigned char m[8];
    cudaMemcpy(m, _m, sizeof(unsigned char)*size_t(8), cudaMemcpyDeviceToHost);

    for(int i=0; i<8; i++)
        fprintf(stdout, "%d %c\n", i, m[i]);

    return 0;
}

which should produce a readable string of characters embedded in the array of unsigned integers supplied to the kernel.

One caveat is that the open64 compiler used for compute 1.x targets would often defeat this strategy of trying to generate vector loads if it could detect that not all of the words in the vector were actually used. So make sure that you touch all of the input words in the input vector type to ensure the compiler plays nicely.

Problem

There are generally two techniques to increase the memory throughput of the global memory on a CUDA kernel on compute capability 1.3 GPUs; memory accesses coalescence and accessing words of at least 4 bytes. With the first technique accesses to the same memory segment by threads of the same half-warp are coalesced to fewer transactions while be accessing words of at least 4 bytes this memory segment is effectively increased from 32 bytes to 128. Update: solution based on talonmies answer. To access 16-byte instead of 1-byte words when there are unsigned chars stored in the global memory, the uint4 vector is commonly used by casting the memory array to uint4. To get the values from the uint4 vector, it can be recasted to uchar4 as below: ``` #include <cuda.h> #include <stdio.h> #include <stdlib.h> __global__ void kernel ( unsigned char *d_text, unsigned char *d_out ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; extern __shared__ unsigned char s_array[]; uint4 *uint4_text = reinterpret_cast<uint4 *>(d_text); uint4 uint4_var; //memory transaction uint4_var = uint4_text[0]; //recast data to uchar4 uchar4 c0 = *reinterpret_cast<uchar4 *>(&uint4_var.x); uchar4 c4 = *reinterpret_cast<uchar4 *>(&uint4_var.y); uchar4 c8 = *reinterpret_cast<uchar4 *>(&uint4_var.z); uchar4 c12 = *reinterpret_cast<uchar4 *>(&uint4_var.w); d_out[idx] = c0.y; } int main ( void ) { unsigned char *d_text, *d_out; unsigned char *h_out = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) ); unsigned char *h_text = ( unsigned char * ) malloc ( 16 * sizeof ( unsigned char ) ); int i; for ( i = 0; i < 16; i++ ) h_text[i] = 65 + i; cudaMalloc ( ( void** ) &d_text, 16 * sizeof ( unsigned char ) ); cudaMalloc ( ( void** ) &d_out, 16 * sizeof ( unsigned char ) ); cudaMemcpy ( d_text, h_text, 16 * sizeof ( unsigned char ), cudaMemcpyHostToDevice ); kernel<<<1,16>>>(d_text, d_out ); cudaMemcpy ( h_out, d_out, 16 * sizeof ( unsigned char ), cudaMemcpyDeviceToHost ); for ( i = 0; i < 16; i++ ) printf("%c\n", h_out[i]); return 0; } ```

Original source