$$__device__$$ void edu_syr_pcpratts_exitMonitorMem($$__global$$ char * gc_info, char * mem, int old){ if(old == -1){ __threadfence(); atomicExch((int *) mem, -1); } }
$$__device__$$ void edu_syr_pcpratts_exitMonitorRef($$__global$$ char * gc_info, int thisref, int old){ char * mem = edu_syr_pcpratts_gc_deref(gc_info, thisref); mem += 12; if(old == -1){ __threadfence(); atomicExch((int *) mem, -1); } }
KOKKOS_FORCEINLINE_FUNCTION void memory_fence() { #if defined( KOKKOS_ATOMICS_USE_CUDA ) __threadfence(); #elif defined( KOKKOS_ATOMICS_USE_GCC ) || \ ( defined( KOKKOS_COMPILER_NVCC ) && defined( KOKKOS_ATOMICS_USE_INTEL ) ) __sync_synchronize(); #elif defined( KOKKOS_ATOMICS_USE_INTEL ) _mm_mfence(); #elif defined( KOKKOS_ATOMICS_USE_OMP31 ) #pragma omp flush #else #error "Error: memory_fence() not defined" #endif }
__device__ void org_trifort_threadfence (){ __threadfence() ; }