Diff 211759

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu

Show First 20 Lines • Show All 56 Lines • ▼ Show 20 Lines	if (numberOfActiveOMPThreads > 1) {
WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE);		WARPSIZE * ((numberOfActiveOMPThreads + WARPSIZE - 1) / WARPSIZE);

PRINT(LD_SYNC,		PRINT(LD_SYNC,
"call kmpc_barrier with %d omp threads, sync parameter %d\n",		"call kmpc_barrier with %d omp threads, sync parameter %d\n",
(int)numberOfActiveOMPThreads, (int)threads);		(int)numberOfActiveOMPThreads, (int)threads);
// Barrier #1 is for synchronization among active threads.		// Barrier #1 is for synchronization among active threads.
named_sync(L1_BARRIER, threads);		named_sync(L1_BARRIER, threads);
}		}
		} else {
		// Still need to flush the memory per the standard.
		__kmpc_flush(loc_ref);
} // numberOfActiveOMPThreads > 1		} // numberOfActiveOMPThreads > 1
PRINT0(LD_SYNC, "completed kmpc_barrier\n");		PRINT0(LD_SYNC, "completed kmpc_barrier\n");
}		}
}		}

// Emit a simple barrier call in SPMD mode. Assumes the caller is in an L0		// Emit a simple barrier call in SPMD mode. Assumes the caller is in an L0
// parallel region and that all worker threads participate.		// parallel region and that all worker threads participate.
EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) {		EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) {
▲ Show 20 Lines • Show All 71 Lines • Show Last 20 Lines

				// RUN: %compile-run-and-check

				#include <omp.h>
				#include <stdio.h>

				int main(int argc, char *argv[]) {
				int data, out, flag = 0;
				#pragma omp target teams num_teams(2) map(tofrom \
				: out) map(to \
				: data, flag) \
				thread_limit(1)
				#pragma omp parallel num_threads(1)
				{
				if (omp_get_team_num() == 0) {
				/* Write to the data buffer that will be read by thread in team 1 */
				data = 42;
				/* Flush data to thread in team 1 */
				#pragma omp barrier
				/* Set flag to release thread in team 1 */
				#pragma omp atomic write
				flag = 1;
				} else if (omp_get_team_num() == 1) {
				/* Loop until we see the update to the flag */
				int val;
				do {
				#pragma omp atomic read
				val = flag;
				} while (val < 1);
				out = data;
				#pragma omp barrier
				}
				}
				// CHECK: out=42.
				/* Value of out will be 42 */
				printf("out=%d.\n", out);
				return !(out == 42);
				}