Skip to content

Instantly share code, notes, and snippets.

@jepio
Created September 29, 2015 08:02
Show Gist options
  • Save jepio/2c796c124d6229747d0b to your computer and use it in GitHub Desktop.
Save jepio/2c796c124d6229747d0b to your computer and use it in GitHub Desktop.
Replacing 3 syncs with one predicated sync.
__global__ void three_syncs(int *data, int size, int thresh)
{
__shared__ int flag;
int tid = threaIdx.x;
do {
__syncthreads();
if (!tid)
flag = 0;
__syncthreads();
if (tid < size) {
data[tid]--;
// ...stuff
if (data[tid] > thresh)
flag = 1;
}
__syncthreads();
} while (flag);
}
__global__ void one_sync_or(int *data, int size, int thresh)
{
int flag;
int tid = threadIdx.x;
do {
flag = 0;
if (tid < size) {
data[tid]--;
// ..stuff
if (data[tid] > thresh)
flag = 1
}
flag = __syncthreads_or(flag);
} while (flag);
};
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment