Exemplo n.º 1
0
$$__device__$$
void edu_syr_pcpratts_exitMonitorMem($$__global$$ char * gc_info, char * mem, int old){
  if(old == -1){   
    __threadfence(); 
    atomicExch((int *) mem, -1); 
  }
}
Exemplo n.º 2
0
$$__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); 
  }
}
Exemplo n.º 3
0
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
}
Exemplo n.º 4
0
__device__ void org_trifort_threadfence       (){ __threadfence()       ; }