# C and CUDA: circular buffer implementation

I have a programme which uses many circular buffers in an identical fashion on a CPU and GPU (C and C/C++ CUDA). I essentially require many queues, however, due to this being run on a GPU, I have limited its length so memory can be set up once at the beginning of the program. I therefore have the code below with a circular buffer/queue. Device and host code is similar with only minor changes for memory efficiency on GPU which is described when necessary. I have implemented this myself so it can be utilized within a hand written kernel on the GPU and so I can compare results between the CPU and GPU for verification purposes. I emit code which allocates/frees memory on the host and device for conciseness.

Worth noting is that, for my purpose, I do not require pop to return the value. Instead I first peek and compare with some condition, and if it is true, pop and discard the data. Furthermore, I can include a peek_tail; however, currently, I do not require this. During a typical program run, I am continuously pushing new data (for all i), and at the same time checking whether the oldest (via peek) is ready to be discarded (again for all i). While this is happening, I am periodically iterating over all items (j) within each (i) buffer. Lastly, when compiling I assume that no pointers alias and compile with strict pointer aliasing flags for both GCC and nvcc when appropriate.

Data structure (note power of 2 for capacity):

typedef struct
{
// Parameters (here only host is shown, identical versions are placed in
// constant memory of the device to access)
int N; // multiple of warp size
int capacity; // power of 2
// Host data
int* size_h;
float* data_h;
// Device data
int* size_d;
float* data_d;
} ring;


Host code (I will explain later why indexing to data is offset) with i={0,...,N-1} and j={0,...,capacity-1} where j is used in a for loop to iterate over the contents of each (i) buffer:

static inline void
Push_ring(ring* in, const int i, const float val)
{
// wrap tail if needed
int x = in->head_h[i] + in->size_h[i];
x &= in->capacity - 1;
in->data_h[(x * in->N) + i] = val;
// increase size
in->size_h[i]++;
if (in->size_h[i] >= in->capacity)
{
#ifdef WITH_WARN
printf("Ring full.\n");
#endif // WITH_WARN
in->size_h[i] = in->capacity;
}
// if first set peek head
if (in->size_h[i] == 1)
}

static inline void
Pop_ring(ring* in, const int i)
{
// if empty
if (in->size_h[i] == 0)
return;
// update size and peek head
in->size_h[i]--;
if (in->size_h[i] < 0)
in->size_h[i] = 0;
if (in->size_h[i] > 0)
else
}

static inline float
Peek_ring(ring* in, const int i)
{
}

static inline int
Size_ring(ring* in, const int i)
{
return in->size_h[i];
}

static inline float
Iterate_ring(ring* in, const int i, const int j)
{
// Wrap i
int x = in->head_h[i] + j;
x &= in->capacity - 1;
// Return pointer to it
return in->data_h[(x * in->N) + i];
}


Device code (same usage of i and j). Here instead of passing a pointer to the structure (which is in host memory), pointers to head, size, peek_head and data are passed (which are in device memory); these refer to ring.XXX_d where XXX is variable name. Lastly, N and capacity are __constant__ variables broadcast to all threads in a warp.

__device__ static __forceinline__ void
float* data, const int i, const int capacity,
const int N, const float val)
{
// use temp variables
int sizeTemp = size[i];
// wrap tail if needed
int x = headTemp + sizeTemp;
x &= capacity - 1;
data[(x * N) + i] = val;
// increase size
sizeTemp++;
if (sizeTemp >= capacity)
{
#ifdef WITH_WARN
printf("Ring full.\n");
#endif // WITH_WARN
sizeTemp = capacity;
else
}
// if first set peek tempHead
if (sizeTemp == 1)
// update from temp variables
size[i] = sizeTemp;
}

__device__ static __forceinline__ void
float* data, const int i,
const int capacity, const int N,
const float val)
{
*sizeTemp = size[i];
// wrap tail if needed
int x = *headTemp + *sizeTemp;
x &= capacity - 1;
data[(x * N) + i] = val;
// increase size
(*sizeTemp)++;
if (*sizeTemp >= capacity)
{
#ifdef WITH_WARN
printf("Ring full.\n");
#endif // WITH_WARN
*sizeTemp = capacity;
}
// if first set peek tempHead
if (*sizeTemp == 1)
}

__device__ static __forceinline__ void
float* data, const int i, const int capacity,
const int N)
{
// use temporary variables
int sizeTemp = size[i];
// if empty
if (sizeTemp == 0)
return;
// update size and peek head
sizeTemp--;
if (sizeTemp < 0)
sizeTemp = 0;
if (sizeTemp > 0) // if else cheaper than trying to do in one
else
// update from temporary variables
size[i] = sizeTemp;
}

__device__ static __forceinline__ void
const int i, const int capacity, const int N)
{
// if empty
if (*sizeTemp == 0)
return;
// update size and peek head
(*sizeTemp)--;
if (*sizeTemp < 0)
*sizeTemp = 0;
if (*sizeTemp > 0) // if else cheaper than trying to do in one
else
}

__device__ static __forceinline__ void
int* size, const int i)
{
size[i] = *sizeTemp;
}

__device__ static __forceinline__ float
{
}

__device__ static __forceinline__ int
Size_ring_GPU(int* size, const int i)
{
return size[i];
}

__device__ static __forceinline__ float
Iterate_ring_GPU(int* head, float* data, const int i,
const int x, const int capacity, const int N)
{
// Wrap
int temp = head[i] + x;
temp &= capacity - 1;
// Return pointer to it
return data[(temp * N) + i];
}

__device__ static __forceinline__ float
const int i, const int x,
const int capacity, const int N)
{
// Wrap i
int temp = headTemp + x;
temp &= capacity - 1;
// Return pointer to it
return data[(temp * N) + i];
}


Here, on the device, I included additional versions of functions: Loading and Loaded. These are used to reduce redundant memory access to head and size as follows:

int headTemp = 0, sizeTemp = 0;
data, i, capacity, N, 1.234f);
capacity, N);


and

int headTemp = head[i];
int y = 0;
float temp = 0.f;
for (y = 0; y < Size_ring_GPU(size, i); y++)


Now, throughout the device code implementation i can be thought of as blockIdx.x * blockDimx.x + threadIdx.x, and with this in mind and the need for memory coalescing for performance reasons, this should hopefully explain the indexing to data_h and data_d (which I keep similar on the host to facilitate copying between host and device). Moreover, as N is a multiple of the warp size, if head and size are identical for all i then memory access should be coalesced and fast. However, during execution head and size for each i will not be identical, leading to memory access being fragmented and less coalesced.

This leads me to my question(s):

• Is it possible to extend my implementation to mitigate this effect? (For example, when size == 0 I reset head = 0 such that, with low activity, buffers will realign towards data[i]. Perhaps I should implement some sort of defrag, and run it periodically on the buffer?)
• Are there any other modifications to be made to increase performance (device code mainly), stability etc.? (General comments will be great as well.)
• Why do you inline/force inline? Typically this decision is best left to the compiler. Have you done profiling comparisons against inlined/non-lined code? Mar 1 '16 at 18:27
• @Reinderien good idea, I haven't actually profiled that difference, I'll give it a go Mar 1 '16 at 18:49

Unfortunately I can't answer your initial question, as the analysis would take more time than I can afford; but the second question I can give you some bullet-points on:

• From my experience the CUDA compiler is not as smart as the mainstream C/C++ compilers, and there's a lot of things that would be optimized out in more advanced compilers that aren't in CUDA, for example, ternary use vs if/else blocks. Likewise, combining statements may have an effect, or not.
• You've got a lot of branching there, the less you can do, the more synchronized your threads will be and the faster block processing will complete. Using the ternary operator may produce faster code, depending on the CUDA version.
• I forget which subsets of the C++ language CUDA supports (it's not quite C or C++, more like C+) but use prefix increment if you can.
• In other compilers such as GCC, combining ++ and if statements can sometimes have small performance benefits, where you'd think they'd be automatically optimized. I would expect CUDA to be no different, but experimentation might be required.
• Remember the const rules: read from right to left, so "const int * const i" means "i is a constant pointer to a constant int". Again, I forget what level of const use CUDA supports. The more const, the more assumptions the compiler can make, generally speaking - which may or may not lead to better opportunity for optimization.
• As Reinderien noted, forcing inline is best left to the compiler, with the exception of very small functions. Even then, using 'inline' might be better practice. -

Bearing all that in mind, see below - none of the code logic should've changed. There are some additional suggestions I've included in the code. Apologies if I've expected any particular feature which is not actually available in CUDA (it's been a while since I used it):

__device__ static void
float* const data, const int i, const int capacity,
const int N, const float val)
{
// use temp variables
int sizeTemp = size[i];

// wrap tail if needed and set value
data[(((headTemp + sizeTemp) ? (capacity - 1)) * N) + i] = val;

// increase size
if (++sizeTemp >= capacity)
{
#ifdef WITH_WARN
printf("Ring full.\n");
#endif // WITH_WARN
sizeTemp = capacity;

}

// if first set peek tempHead
if (sizeTemp == 1)

// update from temp variables
size[i] = sizeTemp;
}

// NOTE: What is the logic of including headTemp and sizeTemp in the parameters?
// Why not use temporary values within the function as with the above function then copy those values back? This will be faster.
// If you need their values afterwards you can always obtain it from head, size and i.
__device__ static void
int* const size, float* const peek_head,
float* const data, const int i,
const int capacity, const int N,
const float val)
{
*sizeTemp = size[i];

// wrap tail if needed
data[(((*headTemp + *sizeTemp) & (capacity - 1)) * N) + i] = val;

// increase size

// NOTE: If you can find a function to reduce *sizetemp to a MAX value of capacity in CUDA, *use that* instead of creating branches!
if (++(*sizeTemp) >= capacity)
{
#ifdef WITH_WARN
printf("Ring full.\n");
#endif // WITH_WARN
*sizeTemp = capacity;

}
// if first set peek tempHead
if (*sizeTemp == 1)
}

__device__ static void
float* const data, const int i, const int capacity,
const int N)
{
// use temporary variables
int sizeTemp = size[i];

// if empty
if (sizeTemp == 0)
return;

// NOTE: an alternative you may want to test and see whether it yields better performance:

// update size and peek head

if (--sizeTemp < 0)
sizeTemp = 0;

peek_head[i] = (sizeTemp > 0) ? data[(headTemp * N) + i] :  -1.f;

// update from temporary variables
size[i] = sizeTemp;
}

__device__ static void
float* const peek_head, float* const data,
const int i, const int capacity, const int N)
{
// if empty
if (*sizeTemp == 0)
return;

// update size and peek head
if (--(*sizeTemp) < 0)
*sizeTemp = 0;

peek_head[i] = (*sizeTemp > 0) ? data[(*headTemp * N) + i] : -1.f;
}

__device__ static inline void
int* const size, const int i)
{
size[i] = *sizeTemp;
}

__device__ static inline float
Peek_ring_GPU(float* const peek_head, const int i)
{
}

__device__ static inline int
Size_ring_GPU(int* const size, const int i)
{
return size[i];
}

__device__ static inline float
Iterate_ring_GPU(int* const head, float* const data, const int i,
const int x, const int capacity, const int N)
{
// Wrap and return pointer to it:
return data[(((head[i] + x) & (capacity - 1)) * N) + i];
}

__device__ static inline float