Created
September 29, 2015 08:02
-
-
Save jepio/2c796c124d6229747d0b to your computer and use it in GitHub Desktop.
Replacing 3 syncs with one predicated sync.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
__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