Skip to content

Commit

Permalink
Merge pull request ROCm#152 from scchan/fix_threadfence_system_test
Browse files Browse the repository at this point in the history
add C++11 compilation flags and minor bug fixes
  • Loading branch information
bensander authored Aug 16, 2017
2 parents 1575fc2 + 4bcb337 commit 04385f1
Showing 1 changed file with 8 additions and 7 deletions.
15 changes: 8 additions & 7 deletions tests/src/deviceLib/hip_threadfence_system.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -21,27 +21,28 @@ THE SOFTWARE.
*/

/* HIT_START
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
* RUN: %t
* HIT_END
*/

#include <vector>
#include <atomic>
#include <thread>
#include <cassert>
#include <cstdio>
#include "hip/hip_runtime.h"
#include "hip/device_functions.h"
#include "test_common.h"


#define HIP_ASSERT(x) (assert((x)==hipSuccess))

__host__ void fence_system() {
std::atomic_thread_fence(std::memory_order_seq_cst);
}

__device__ void fence_system() {
__host__ __device__ void fence_system() {
#ifdef __HIP_DEVICE_COMPILE__
__threadfence_system();
#else
std::atomic_thread_fence(std::memory_order_seq_cst);
#endif
}

__host__ __device__ void round_robin(const int id, const int num_dev, const int num_iter, volatile int* data, volatile int* flag) {
Expand Down

0 comments on commit 04385f1

Please sign in to comment.