From 513ba4eb8ab3080361ef9759d0b9f6808c3df1be Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Sat, 2 Jun 2018 12:27:17 +0100
Subject: [PATCH 01/11] Switch the atomic implementation to use Clang 
 builtins.

---
 .../include/hip/hcc_detail/device_functions.h |   3 +-
 hipamd/include/hip/hcc_detail/hip_atomic.h    | 265 +++++++++++++++
 hipamd/include/hip/hcc_detail/hip_runtime.h   |  77 +----
 hipamd/include/hip/hip_runtime.h              |   2 +-
 hipamd/src/device_util.cpp                    | 123 -------
 .../src/deviceLib/hipSimpleAtomicsTest.cpp    | 316 ++++++++++--------
 6 files changed, 447 insertions(+), 339 deletions(-)
 create mode 100644 hipamd/include/hip/hcc_detail/hip_atomic.h

diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h
index 28d874b27a..20a365ebbe 100644
--- a/hipamd/include/hip/hcc_detail/device_functions.h
+++ b/hipamd/include/hip/hcc_detail/device_functions.h
@@ -23,7 +23,8 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
 
-#include <hip/hip_runtime.h>
+#include "host_defines.h"
+
 #include <hip/hip_vector_types.h>
 
 
diff --git a/hipamd/include/hip/hcc_detail/hip_atomic.h b/hipamd/include/hip/hcc_detail/hip_atomic.h
new file mode 100644
index 0000000000..4af1794ba0
--- /dev/null
+++ b/hipamd/include/hip/hcc_detail/hip_atomic.h
@@ -0,0 +1,265 @@
+#pragma once
+
+#include "device_functions.h"
+
+__device__
+inline
+int atomicCAS(int* address, int compare, int val)
+{
+    __atomic_compare_exchange_n(
+        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+
+    return compare;
+}
+__device__
+inline
+unsigned int atomicCAS(
+    unsigned int* address, unsigned int compare, unsigned int val)
+{
+    __atomic_compare_exchange_n(
+        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+
+    return compare;
+}
+__device__
+inline
+unsigned long long atomicCAS(
+    unsigned long long* address,
+    unsigned long long compare,
+    unsigned long long val)
+{
+    __atomic_compare_exchange_n(
+        address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+
+    return compare;
+}
+
+__device__
+inline
+int atomicAdd(int* address, int val)
+{
+    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicAdd(unsigned int* address, unsigned int val)
+{
+    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicAdd(
+    unsigned long long* address, unsigned long long val)
+{
+    return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+float atomicAdd(float* address, float val)
+{
+    unsigned int* uaddr{reinterpret_cast<unsigned int*>(uaddr)};
+    unsigned int old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+    unsigned int r;
+
+    do {
+        r = old;
+        old = atomicCAS(uaddr, r, __float_as_uint(val + __uint_as_float(r)));
+    } while (r != old);
+
+    return __uint_as_float(r);
+}
+__device__
+inline
+double atomicAdd(double* address, double val)
+{
+    unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(uaddr)};
+    unsigned long long old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+    unsigned long long r;
+
+    do {
+        r = old;
+        old = atomicCAS(
+            uaddr, r, __double_as_longlong(val + __longlong_as_double(r)));
+    } while (r != old);
+
+    return __longlong_as_double(r);
+}
+
+__device__
+inline
+int atomicSub(int* address, int val)
+{
+    return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicSub(unsigned int* address, unsigned int val)
+{
+    return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED);
+}
+
+__device__
+inline
+int atomicExch(int* address, int val)
+{
+    return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicExch(unsigned int* address, unsigned int val)
+{
+    return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicExch(unsigned long long* address, unsigned long long val)
+{
+    return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+float atomicExch(float* address, float val)
+{
+    return __uint_as_float(__atomic_exchange_n(
+        reinterpret_cast<unsigned int*>(address),
+        __float_as_uint(val),
+        __ATOMIC_RELAXED));
+}
+
+__device__
+inline
+int atomicMin(int* address, int val)
+{
+    return __sync_fetch_and_min(address, val);
+}
+__device__
+inline
+unsigned int atomicMin(unsigned int* address, unsigned int val)
+{
+    return __sync_fetch_and_umin(address, val);
+}
+__device__
+inline
+unsigned long long atomicMin(
+    unsigned long long* address, unsigned long long val)
+{
+    unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+    while (val < tmp) { tmp = atomicCAS(address, tmp, val); }
+
+    return tmp;
+}
+
+__device__
+inline
+int atomicMax(int* address, int val)
+{
+    return __sync_fetch_and_max(address, val);
+}
+__device__
+inline
+unsigned int atomicMax(unsigned int* address, unsigned int val)
+{
+    return __sync_fetch_and_umax(address, val);
+}
+__device__
+inline
+unsigned long long atomicMax(
+    unsigned long long* address, unsigned long long val)
+{
+    unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+    while (tmp < val) { tmp = atomicCAS(address, tmp, val); }
+
+    return tmp;
+}
+
+__device__
+inline
+unsigned int atomicInc(unsigned int* address, unsigned int val)
+{
+    extern unsigned int __builtin_amdgcn_atomic_inc(
+        unsigned int*,
+        unsigned int,
+        unsigned int,
+        unsigned int,
+        bool) __asm("llvm.amdgcn.atomic.inc.i32.p0i32");
+
+    return __builtin_amdgcn_atomic_inc(
+        address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
+}
+
+__device__
+inline
+unsigned int atomicDec(unsigned int* address, unsigned int val)
+{
+    extern unsigned int __builtin_amdgcn_atomic_dec(
+        unsigned int*,
+        unsigned int,
+        unsigned int,
+        unsigned int,
+        bool) __asm("llvm.amdgcn.atomic.dec.i32.p0i32");
+
+    return __builtin_amdgcn_atomic_dec(
+        address, val, __ATOMIC_RELAXED, 1 /* Device scope */, false);
+}
+
+__device__
+inline
+int atomicAnd(int* address, int val)
+{
+    return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicAnd(unsigned int* address, unsigned int val)
+{
+    return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicAnd(
+    unsigned long long* address, unsigned long long val)
+{
+    return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
+}
+
+__device__
+inline
+int atomicOr(int* address, int val)
+{
+    return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicOr(unsigned int* address, unsigned int val)
+{
+    return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicOr(
+    unsigned long long* address, unsigned long long val)
+{
+    return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
+}
+
+__device__
+inline
+int atomicXor(int* address, int val)
+{
+    return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicXor(unsigned int* address, unsigned int val)
+{
+    return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicXor(
+    unsigned long long* address, unsigned long long val)
+{
+    return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
+}
+
+// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block.
\ No newline at end of file
diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index 92f06e9174..c77b4b85bd 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -88,6 +88,7 @@ extern int HIP_TRACE_API;
 #ifdef __cplusplus
 #include <hip/hcc_detail/hip_ldg.h>
 #endif
+#include <hip/hcc_detail/hip_atomic.h>
 #include <hip/hcc_detail/host_defines.h>
 #include <hip/hcc_detail/math_functions.h>
 #include <hip/hcc_detail/device_functions.h>
@@ -175,82 +176,6 @@ __device__ clock_t clock();
 // abort
 __device__ void abort();
 
-// atomicAdd()
-__device__ int atomicAdd(int* address, int val);
-__device__ unsigned int atomicAdd(unsigned int* address, unsigned int val);
-
-__device__ unsigned long long int atomicAdd(unsigned long long int* address,
-                                            unsigned long long int val);
-
-__device__ float atomicAdd(float* address, float val);
-
-
-// atomicSub()
-__device__ int atomicSub(int* address, int val);
-
-__device__ unsigned int atomicSub(unsigned int* address, unsigned int val);
-
-
-// atomicExch()
-__device__ int atomicExch(int* address, int val);
-
-__device__ unsigned int atomicExch(unsigned int* address, unsigned int val);
-
-__device__ unsigned long long int atomicExch(unsigned long long int* address,
-                                             unsigned long long int val);
-
-__device__ float atomicExch(float* address, float val);
-
-
-// atomicMin()
-__device__ int atomicMin(int* address, int val);
-__device__ unsigned int atomicMin(unsigned int* address, unsigned int val);
-__device__ unsigned long long int atomicMin(unsigned long long int* address,
-                                            unsigned long long int val);
-
-
-// atomicMax()
-__device__ int atomicMax(int* address, int val);
-__device__ unsigned int atomicMax(unsigned int* address, unsigned int val);
-__device__ unsigned long long int atomicMax(unsigned long long int* address,
-                                            unsigned long long int val);
-
-
-// atomicCAS()
-__device__ int atomicCAS(int* address, int compare, int val);
-__device__ unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val);
-__device__ unsigned long long int atomicCAS(unsigned long long int* address,
-                                            unsigned long long int compare,
-                                            unsigned long long int val);
-
-
-// atomicAnd()
-__device__ int atomicAnd(int* address, int val);
-__device__ unsigned int atomicAnd(unsigned int* address, unsigned int val);
-__device__ unsigned long long int atomicAnd(unsigned long long int* address,
-                                            unsigned long long int val);
-
-
-// atomicOr()
-__device__ int atomicOr(int* address, int val);
-__device__ unsigned int atomicOr(unsigned int* address, unsigned int val);
-__device__ unsigned long long int atomicOr(unsigned long long int* address,
-                                           unsigned long long int val);
-
-
-// atomicXor()
-__device__ int atomicXor(int* address, int val);
-__device__ unsigned int atomicXor(unsigned int* address, unsigned int val);
-__device__ unsigned long long int atomicXor(unsigned long long int* address,
-                                            unsigned long long int val);
-
-// atomicInc()
-__device__ unsigned int atomicInc(unsigned int* address, unsigned int val);
-
-
-// atomicDec()
-__device__ unsigned int atomicDec(unsigned int* address, unsigned int val);
-
 // warp vote function __all __any __ballot
 __device__ int __all(int input);
 __device__ int __any(int input);
diff --git a/hipamd/include/hip/hip_runtime.h b/hipamd/include/hip/hip_runtime.h
index 157fc88a43..937ba61ecf 100644
--- a/hipamd/include/hip/hip_runtime.h
+++ b/hipamd/include/hip/hip_runtime.h
@@ -64,4 +64,4 @@ THE SOFTWARE.
 #include <hip/hip_runtime_api.h>
 #include <hip/hip_vector_types.h>
 
-#endif
+#endif
\ No newline at end of file
diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp
index 5ce014b2b9..11f992a510 100644
--- a/hipamd/src/device_util.cpp
+++ b/hipamd/src/device_util.cpp
@@ -761,129 +761,6 @@ __device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); };
 // abort
 __device__ void abort() { return hc::abort(); }
 
-// atomicAdd()
-__device__ int atomicAdd(int* address, int val) { return hc::atomic_fetch_add(address, val); }
-__device__ unsigned int atomicAdd(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_add(address, val);
-}
-__device__ unsigned long long int atomicAdd(unsigned long long int* address,
-                                            unsigned long long int val) {
-    return (long long int)hc::atomic_fetch_add((uint64_t*)address, (uint64_t)val);
-}
-__device__ float atomicAdd(float* address, float val) { return hc::atomic_fetch_add(address, val); }
-
-// atomicSub()
-__device__ int atomicSub(int* address, int val) { return hc::atomic_fetch_sub(address, val); }
-__device__ unsigned int atomicSub(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_sub(address, val);
-}
-
-// atomicExch()
-__device__ int atomicExch(int* address, int val) { return hc::atomic_exchange(address, val); }
-__device__ unsigned int atomicExch(unsigned int* address, unsigned int val) {
-    return hc::atomic_exchange(address, val);
-}
-__device__ unsigned long long int atomicExch(unsigned long long int* address,
-                                             unsigned long long int val) {
-    return (long long int)hc::atomic_exchange((uint64_t*)address, (uint64_t)val);
-}
-__device__ float atomicExch(float* address, float val) { return hc::atomic_exchange(address, val); }
-
-// atomicMin()
-__device__ int atomicMin(int* address, int val) { return hc::atomic_fetch_min(address, val); }
-__device__ unsigned int atomicMin(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_min(address, val);
-}
-__device__ unsigned long long int atomicMin(unsigned long long int* address,
-                                            unsigned long long int val) {
-    return (long long int)hc::atomic_fetch_min((uint64_t*)address, (uint64_t)val);
-}
-
-// atomicMax()
-__device__ int atomicMax(int* address, int val) { return hc::atomic_fetch_max(address, val); }
-__device__ unsigned int atomicMax(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_max(address, val);
-}
-__device__ unsigned long long int atomicMax(unsigned long long int* address,
-                                            unsigned long long int val) {
-    return (long long int)hc::atomic_fetch_max((uint64_t*)address, (uint64_t)val);
-}
-
-// atomicCAS()
-template <typename T>
-__device__ T atomicCAS_impl(T* address, T compare, T val) {
-    // the implementation assumes the atomic is lock-free and
-    // has the same size as the non-atmoic equivalent type
-    static_assert(sizeof(T) == sizeof(std::atomic<T>),
-                  "size mismatch between atomic and non-atomic types");
-
-    union {
-        T* address;
-        std::atomic<T>* atomic_address;
-    } u;
-    u.address = address;
-
-    T expected = compare;
-
-    // hcc should generate a system scope atomic CAS
-    std::atomic_compare_exchange_weak_explicit(
-        u.atomic_address, &expected, val, std::memory_order_acq_rel, std::memory_order_relaxed);
-    return expected;
-}
-
-__device__ int atomicCAS(int* address, int compare, int val) {
-    return atomicCAS_impl(address, compare, val);
-}
-__device__ unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val) {
-    return atomicCAS_impl(address, compare, val);
-}
-__device__ unsigned long long int atomicCAS(unsigned long long int* address,
-                                            unsigned long long int compare,
-                                            unsigned long long int val) {
-    return atomicCAS_impl(address, compare, val);
-}
-
-// atomicAnd()
-__device__ int atomicAnd(int* address, int val) { return hc::atomic_fetch_and(address, val); }
-__device__ unsigned int atomicAnd(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_and(address, val);
-}
-__device__ unsigned long long int atomicAnd(unsigned long long int* address,
-                                            unsigned long long int val) {
-    return (long long int)hc::atomic_fetch_and((uint64_t*)address, (uint64_t)val);
-}
-
-// atomicOr()
-__device__ int atomicOr(int* address, int val) { return hc::atomic_fetch_or(address, val); }
-__device__ unsigned int atomicOr(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_or(address, val);
-}
-__device__ unsigned long long int atomicOr(unsigned long long int* address,
-                                           unsigned long long int val) {
-    return (long long int)hc::atomic_fetch_or((uint64_t*)address, (uint64_t)val);
-}
-
-// atomicXor()
-__device__ int atomicXor(int* address, int val) { return hc::atomic_fetch_xor(address, val); }
-__device__ unsigned int atomicXor(unsigned int* address, unsigned int val) {
-    return hc::atomic_fetch_xor(address, val);
-}
-__device__ unsigned long long int atomicXor(unsigned long long int* address,
-                                            unsigned long long int val) {
-    return (long long int)hc::atomic_fetch_xor((uint64_t*)address, (uint64_t)val);
-}
-
-// atomicInc
-__device__ unsigned int atomicInc(unsigned int* address, unsigned int val) {
-    return hc::__atomic_wrapinc(address, val);
-}
-
-// atomicDec
-__device__ unsigned int atomicDec(unsigned int* address, unsigned int val) {
-    return hc::__atomic_wrapdec(address, val);
-}
-
-
 // warp vote function __all __any __ballot
 __device__ int __all(int input) { return hc::__all(input); }
 
diff --git a/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
index ce1d0372ff..129d7c1aa8 100644
--- a/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
+++ b/hipamd/tests/src/deviceLib/hipSimpleAtomicsTest.cpp
@@ -23,134 +23,37 @@ THE SOFTWARE.
  * HIT_END
  */
 
-// includes, system
-#include <stdlib.h>
-#include <stdio.h>
-#include <string.h>
-#include <math.h>
-
-
 // Includes HIP Runtime
 #include "hip/hip_runtime.h"
 #include <test_common.h>
 
+// includes, system
+#include <algorithm>
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+#include <math.h>
+#include <type_traits>
+
 #define EXIT_WAIVED 2
 
 const char* sampleName = "hipSimpleAtomicsTest";
 
+using namespace std;
+
 ////////////////////////////////////////////////////////////////////////////////
 // Auto-Verification Code
 bool testResult = true;
 
 ////////////////////////////////////////////////////////////////////////////////
-// Declaration, forward
-void runTest(int argc, char** argv);
 
+bool computeGoldBitwise(...) {
+    return true;
+}
 
-#define min(a, b) (a) < (b) ? (a) : (b)
-#define max(a, b) (a) > (b) ? (a) : (b)
-
-int computeGold(int* gpuData, const int len) {
-    int val = 0;
-
-    for (int i = 0; i < len; ++i) {
-        val += 10;
-    }
-
-    if (val != gpuData[0]) {
-        printf("atomicAdd failed\n");
-        return false;
-    }
-
-    val = 0;
-
-    for (int i = 0; i < len; ++i) {
-        val -= 10;
-    }
-
-    if (val != gpuData[1]) {
-        printf("atomicSub failed\n");
-        return false;
-    }
-
-    bool found = false;
-
-    for (int i = 0; i < len; ++i) {
-        // third element should be a member of [0, len)
-        if (i == gpuData[2]) {
-            found = true;
-            break;
-        }
-    }
-
-    if (!found) {
-        printf("atomicExch failed\n");
-        return false;
-    }
-
-    val = -(1 << 8);
-
-    for (int i = 0; i < len; ++i) {
-        // fourth element should be len-1
-        val = max(val, i);
-    }
-
-    if (val != gpuData[3]) {
-        printf("atomicMax failed\n");
-        return false;
-    }
-
-    val = 1 << 8;
-
-    for (int i = 0; i < len; ++i) {
-        val = min(val, i);
-    }
-
-    if (val != gpuData[4]) {
-        printf("atomicMin failed\n");
-        return false;
-    }
-
-    int limit = 17;
-    val = 0;
-
-    for (int i = 0; i < len; ++i) {
-        val = (val >= limit) ? 0 : val + 1;
-    }
-
-    if (val != gpuData[5]) {
-        printf("atomicInc failed\n");
-        return false;
-    }
-
-    limit = 137;
-    val = 0;
-
-    for (int i = 0; i < len; ++i) {
-        val = ((val == 0) || (val > limit)) ? limit : val - 1;
-    }
-
-    if (val != gpuData[6]) {
-        printf("atomicDec failed\n");
-        return false;
-    }
-
-    found = false;
-
-    for (int i = 0; i < len; ++i) {
-        // eighth element should be a member of [0, len)
-        if (i == gpuData[7]) {
-            found = true;
-            break;
-        }
-    }
-
-    if (!found) {
-        printf("atomicCAS failed\n");
-        return false;
-    }
-
-    val = 0xff;
+template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
+bool computeGoldBitwise(T* gpuData, int len) {
+    T val = 0xff;
 
     for (int i = 0; i < len; ++i) {
         // 9th element should be 1
@@ -189,22 +92,142 @@ int computeGold(int* gpuData, const int len) {
     return true;
 }
 
-__global__ void testKernel(hipLaunchParm lp, int* g_odata) {
+template<typename T>
+bool computeGold(T* gpuData, int len) {
+    T val = 0;
+
+    for (int i = 0; i < len; ++i) {
+        val += 10;
+    }
+
+    if (val != gpuData[0]) {
+        printf("atomicAdd failed\n");
+        return false;
+    }
+
+    val = 0;
+
+    for (int i = 0; i < len; ++i) {
+        val -= 10;
+    }
+
+    if (val != gpuData[1]) {
+        printf("atomicSub failed\n");
+        return false;
+    }
+
+    bool found = false;
+
+    for (T i = 0; i < len; ++i) {
+        // third element should be a member of [0, len)
+        if (i == gpuData[2]) {
+            found = true;
+            break;
+        }
+    }
+
+    if (!found) {
+        printf("atomicExch failed\n");
+        return false;
+    }
+
+    val = -(1 << 8);
+
+    for (T i = 0; i < len; ++i) {
+        // fourth element should be len-1
+        val = max(val, i);
+    }
+
+    if (val != gpuData[3]) {
+        printf("atomicMax failed\n");
+        return false;
+    }
+
+    val = 1 << 8;
+
+    for (T i = 0; i < len; ++i) {
+        val = min(val, i);
+    }
+
+    if (val != gpuData[4]) {
+        printf("atomicMin failed\n");
+        return false;
+    }
+
+    int limit = 17;
+    val = 0;
+
+    for (int i = 0; i < len; ++i) {
+        val = (val >= limit) ? 0 : val + 1;
+    }
+
+    if (val != gpuData[5]) {
+        printf("atomicInc failed\n");
+        return false;
+    }
+
+    limit = 137;
+    val = 0;
+
+    for (int i = 0; i < len; ++i) {
+        val = ((val == 0) || (val > limit)) ? limit : val - 1;
+    }
+
+    if (val != gpuData[6]) {
+        printf("atomicDec failed\n");
+        return false;
+    }
+
+    found = false;
+
+    for (T i = 0; i < len; ++i) {
+        // eighth element should be a member of [0, len)
+        if (i == gpuData[7]) {
+            found = true;
+            break;
+        }
+    }
+    if (!found) {
+        printf("atomicCAS failed\n");
+        return false;
+    }
+
+    return computeGoldBitwise(gpuData, len);
+}
+
+__device__
+void testKernelExch(...) {}
+
+template<typename T, typename enable_if<!is_same<T, double>{}>::type* = nullptr>
+__device__
+void testKernelExch(T* g_odata) {
     // access thread id
-    const unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
-
-    // Test various atomic instructions
-
-    // Arithmetic atomic instructions
-
-    // Atomic addition
-    atomicAdd(&g_odata[0], 10);
-
-    // Atomic subtraction (final should be 0)
-    atomicSub(&g_odata[1], 10);
+    const T tid = blockDim.x * blockIdx.x + threadIdx.x;
 
     // Atomic exchange
     atomicExch(&g_odata[2], tid);
+}
+
+__device__
+void testKernelSub(...) {}
+
+template<
+    typename T, 
+    typename enable_if<
+        is_same<T, int>{} || is_same<T, unsigned int>{}>::type* = nullptr>
+void testKernelSub(T* g_odata) {
+    // Atomic subtraction (final should be 0)
+    atomicSub(&g_odata[1], 10);
+}
+
+__device__
+void testKernelIntegral(...) {}
+
+template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
+__device__
+void testKernelIntegral(T* g_odata) {
+    // access thread id
+    const T tid = blockDim.x * blockIdx.x + threadIdx.x;
 
     // Atomic maximum
     atomicMax(&g_odata[3], tid);
@@ -231,20 +254,21 @@ __global__ void testKernel(hipLaunchParm lp, int* g_odata) {
 
     // Atomic XOR
     atomicXor(&g_odata[10], tid);
+
+    testKernelSub(g_odata);
 }
 
+template<typename T>
+__global__ void testKernel(T* g_odata) {
+    // Atomic addition
+    atomicAdd(&g_odata[0], 10);
 
-int main(int argc, char** argv) {
-    printf("%s starting...\n", sampleName);
-
-    runTest(argc, argv);
-
-    hipDeviceReset();
-    printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!");
-    exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
+    testKernelIntegral(g_odata);
+    testKernelExch(g_odata);
 }
 
-void runTest(int argc, char** argv) {
+template<typename T>
+void runTest() {
     hipDeviceProp_t deviceProp;
     deviceProp.major = 0;
     deviceProp.minor = 0;
@@ -262,10 +286,10 @@ void runTest(int argc, char** argv) {
     unsigned int numThreads = 256;
     unsigned int numBlocks = 64;
     unsigned int numData = 11;
-    unsigned int memSize = sizeof(int) * numData;
+    unsigned int memSize = sizeof(T) * numData;
 
     // allocate mem for the result on host side
-    int* hOData = (int*)malloc(memSize);
+    T* hOData = (T*)malloc(memSize);
 
     // initialize the memory
     for (unsigned int i = 0; i < numData; i++) hOData[i] = 0;
@@ -274,13 +298,14 @@ void runTest(int argc, char** argv) {
     hOData[8] = hOData[10] = 0xff;
 
     // allocate device memory for result
-    int* dOData;
+    T* dOData;
     hipMalloc((void**)&dOData, memSize);
     // copy host memory to device to initialize to zero
     hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice);
 
     // execute the kernel
-    hipLaunchKernel(testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData);
+    hipLaunchKernelGGL(
+        testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData);
 
     // Copy result from device to host
     hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost);
@@ -294,3 +319,18 @@ void runTest(int argc, char** argv) {
 
     passed();
 }
+
+
+int main(int argc, char** argv) {
+    printf("%s starting...\n", sampleName);
+
+    runTest<int>();
+    runTest<unsigned int>();
+    runTest<unsigned long long>();
+    runTest<float>();
+    runTest<double>();
+
+    hipDeviceReset();
+    printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!");
+    exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
+}
\ No newline at end of file

From 13a191774147c6ae6e28023108504cce81ce353a Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Sat, 2 Jun 2018 17:46:37 +0100
Subject: [PATCH 02/11] Add missing __device__ for forward declares.

---
 hipamd/include/hip/hcc_detail/hip_atomic.h | 10 +++++++---
 1 file changed, 7 insertions(+), 3 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_atomic.h b/hipamd/include/hip/hcc_detail/hip_atomic.h
index 4af1794ba0..1304cf74a9 100644
--- a/hipamd/include/hip/hcc_detail/hip_atomic.h
+++ b/hipamd/include/hip/hcc_detail/hip_atomic.h
@@ -176,7 +176,9 @@ __device__
 inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
 {
-    extern unsigned int __builtin_amdgcn_atomic_inc(
+    __device__
+    extern 
+    unsigned int __builtin_amdgcn_atomic_inc(
         unsigned int*,
         unsigned int,
         unsigned int,
@@ -191,7 +193,9 @@ __device__
 inline
 unsigned int atomicDec(unsigned int* address, unsigned int val)
 {
-    extern unsigned int __builtin_amdgcn_atomic_dec(
+    __device__
+    extern 
+    unsigned int __builtin_amdgcn_atomic_dec(
         unsigned int*,
         unsigned int,
         unsigned int,
@@ -262,4 +266,4 @@ unsigned long long atomicXor(
     return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
 }
 
-// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block.
\ No newline at end of file
+// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block.

From 92c106bcf5a7464b7843b9df31ed9ff25fb36202 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Sun, 3 Jun 2018 03:03:55 +0100
Subject: [PATCH 03/11] Fix hideous typos.

---
 hipamd/include/hip/hcc_detail/hip_atomic.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_atomic.h b/hipamd/include/hip/hcc_detail/hip_atomic.h
index 1304cf74a9..ce4b7c9008 100644
--- a/hipamd/include/hip/hcc_detail/hip_atomic.h
+++ b/hipamd/include/hip/hcc_detail/hip_atomic.h
@@ -57,7 +57,7 @@ __device__
 inline
 float atomicAdd(float* address, float val)
 {
-    unsigned int* uaddr{reinterpret_cast<unsigned int*>(uaddr)};
+    unsigned int* uaddr{reinterpret_cast<unsigned int*>(address)};
     unsigned int old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
     unsigned int r;
 
@@ -72,7 +72,7 @@ __device__
 inline
 double atomicAdd(double* address, double val)
 {
-    unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(uaddr)};
+    unsigned long long* uaddr{reinterpret_cast<unsigned long long*>(address)};
     unsigned long long old{__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
     unsigned long long r;
 

From 62155871979b17773fbfb6b8bf94c41252321e3a Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Tue, 5 Jun 2018 08:46:57 -0400
Subject: [PATCH 04/11] Fix channel_descriptor.h about vector 3 for gcc

---
 hipamd/include/hip/hcc_detail/channel_descriptor.h | 10 ++++++++++
 1 file changed, 10 insertions(+)

diff --git a/hipamd/include/hip/hcc_detail/channel_descriptor.h b/hipamd/include/hip/hcc_detail/channel_descriptor.h
index 5fb0faeee1..de290fafc1 100644
--- a/hipamd/include/hip/hcc_detail/channel_descriptor.h
+++ b/hipamd/include/hip/hcc_detail/channel_descriptor.h
@@ -92,6 +92,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<char2>() {
     return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
 }
 
+#ifndef __GNUC__  // vector3 is the same as vector4
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<uchar3>() {
     int e = (int)sizeof(unsigned char) * 8;
@@ -103,6 +104,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<char3>() {
     int e = (int)sizeof(signed char) * 8;
     return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
 }
+#endif
 
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<uchar4>() {
@@ -152,6 +154,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<short2>() {
     return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
 }
 
+#ifndef __GNUC__
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<ushort3>() {
     int e = (int)sizeof(unsigned short) * 8;
@@ -163,6 +166,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<short3>() {
     int e = (int)sizeof(signed short) * 8;
     return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
 }
+#endif
 
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<ushort4>() {
@@ -212,6 +216,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<int2>() {
     return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
 }
 
+#ifndef __GNUC__
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<uint3>() {
     int e = (int)sizeof(unsigned int) * 8;
@@ -223,6 +228,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<int3>() {
     int e = (int)sizeof(signed int) * 8;
     return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
 }
+#endif
 
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<uint4>() {
@@ -254,11 +260,13 @@ inline hipChannelFormatDesc hipCreateChannelDesc<float2>() {
     return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindFloat);
 }
 
+#ifndef __GNUC__
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<float3>() {
     int e = (int)sizeof(float) * 8;
     return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindFloat);
 }
+#endif
 
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<float4>() {
@@ -302,6 +310,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<long2>() {
     return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
 }
 
+#ifndef __GNUC__
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<ulong3>() {
     int e = (int)sizeof(unsigned long) * 8;
@@ -313,6 +322,7 @@ inline hipChannelFormatDesc hipCreateChannelDesc<long3>() {
     int e = (int)sizeof(signed long) * 8;
     return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
 }
+#endif
 
 template <>
 inline hipChannelFormatDesc hipCreateChannelDesc<ulong4>() {

From 9566616dd2fa5d595ea8f08dc7de2c767ee6f77a Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu@amd.com>
Date: Fri, 8 Jun 2018 07:44:02 -0400
Subject: [PATCH 05/11] Add __device__ to device functions in
 hip_fp16_math_fwd.h

---
 .../hip/hcc_detail/hip_fp16_math_fwd.h        | 72 +++++++++----------
 1 file changed, 36 insertions(+), 36 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
index 81f9fe4761..fccbcfbfdc 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
@@ -28,49 +28,49 @@ THE SOFTWARE.
 
 extern "C"
 {
-    __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
-    _Float16 __ocml_cos_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
-    __attribute__((const))
+    __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
+    __device__ _Float16 __ocml_cos_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
+    __device__ __attribute__((const))
     _Float16 __ocml_fma_f16(_Float16, _Float16, _Float16);
-    __attribute__((const)) int __ocml_isinf_f16(_Float16);
-    __attribute__((const)) int __ocml_isnan_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
-    __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
-    _Float16 __ocml_sin_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
+    __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
+    __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
+    __device__ _Float16 __ocml_sin_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
 
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
 
-    __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
-    __2f16 __ocml_cos_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
-    __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
-    __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
-    inline
+    __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
+    __device__ __2f16 __ocml_cos_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
+    __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
+    __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
+    __device__ inline
     __2f16 __llvm_amdgcn_rcp_2f16(__2f16 x) // Not currently exposed by ROCDL.
     {
         return __2f16{__llvm_amdgcn_rcp_f16(x.x), __llvm_amdgcn_rcp_f16(x.y)};
     }
-    __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
-    __2f16 __ocml_sin_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
+    __device__ __2f16 __ocml_sin_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
 }
\ No newline at end of file

From 704f8c7b8f11f46adf7958122276fab6b1db183e Mon Sep 17 00:00:00 2001
From: Junjie Bai <bai@in.tum.de>
Date: Tue, 12 Jun 2018 23:53:18 -0700
Subject: [PATCH 06/11] Properly handle (empty) cmake generator expression

---
 hipamd/cmake/FindHIP.cmake | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/hipamd/cmake/FindHIP.cmake b/hipamd/cmake/FindHIP.cmake
index 6668e60332..a45fb1b44d 100644
--- a/hipamd/cmake/FindHIP.cmake
+++ b/hipamd/cmake/FindHIP.cmake
@@ -265,7 +265,7 @@ endmacro()
 set(HIP_HIPCC_INCLUDE_ARGS_USER "")
 macro(HIP_INCLUDE_DIRECTORIES)
     foreach(dir ${ARGN})
-        list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER -I${dir})
+        list(APPEND HIP_HIPCC_INCLUDE_ARGS_USER $<$<BOOL:${dir}>:-I${dir}>)
     endforeach()
 endmacro()
 
@@ -373,7 +373,7 @@ macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files
     list(REMOVE_DUPLICATES _hip_include_directories)
     if(_hip_include_directories)
         foreach(dir ${_hip_include_directories})
-            list(APPEND HIP_HIPCC_INCLUDE_ARGS -I${dir})
+            list(APPEND HIP_HIPCC_INCLUDE_ARGS $<$<BOOL:${dir}>:-I${dir}>)
         endforeach()
     endif()
 

From 04640992dcc3a5317047e2f627e92a3fb435590d Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Fri, 18 May 2018 10:14:46 -0500
Subject: [PATCH 07/11] HIP program state re-initialization logic

This commit is to support kernels dynamically loaded thru means such as
dlopen() after HIP runtime initializes.
---
 .../include/hip/hcc_detail/program_state.hpp  |   9 +-
 hipamd/src/functional_grid_launch.inl         |  16 ++-
 hipamd/src/program_state.cpp                  | 131 +++++++++++++++---
 3 files changed, 124 insertions(+), 32 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/program_state.hpp b/hipamd/include/hip/hcc_detail/program_state.hpp
index ac689fdb89..bdb87b3509 100644
--- a/hipamd/include/hip/hcc_detail/program_state.hpp
+++ b/hipamd/include/hip/hcc_detail/program_state.hpp
@@ -93,11 +93,12 @@ public:
     }
 };
 
-const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables();
+const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables(
+    bool rebuild = false);
 const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>&
-functions();
-const std::unordered_map<std::uintptr_t, std::string>& function_names();
-std::unordered_map<std::string, void*>& globals();
+functions(bool rebuild = false);
+const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
+std::unordered_map<std::string, void*>& globals(bool rebuild = false);
 
 hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
                                  hsa_agent_t agent);
diff --git a/hipamd/src/functional_grid_launch.inl b/hipamd/src/functional_grid_launch.inl
index 9ecad51476..336bb5d121 100644
--- a/hipamd/src/functional_grid_launch.inl
+++ b/hipamd/src/functional_grid_launch.inl
@@ -92,13 +92,19 @@ namespace hip_impl
         hipStream_t stream,
         void** kernarg)
     {
-        const auto it0 = functions().find(function_address);
+        auto it0 = functions().find(function_address);
 
         if (it0 == functions().cend()) {
-            throw runtime_error{
-                "No device code available for function: " +
-                name(function_address)
-            };
+            // Re-init device code maps once again to help locate kernels
+            // loaded after HIP runtime initialization via means such as
+            // dlopen().
+            it0 = functions(true).find(function_address);
+            if (it0 == functions().cend()) {
+                throw runtime_error{
+                    "No device code available for function: " +
+                    name(function_address)
+                };
+            }
         }
 
         auto agent = target_agent(stream);
diff --git a/hipamd/src/program_state.cpp b/hipamd/src/program_state.cpp
index c4478bec2f..3ad9285b8c 100644
--- a/hipamd/src/program_state.cpp
+++ b/hipamd/src/program_state.cpp
@@ -74,11 +74,15 @@ vector<string> copy_names_of_undefined_symbols(const symbol_section_accessor& se
 }
 
 const std::unordered_map<std::string, std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>&
-symbol_addresses() {
+symbol_addresses(bool rebuild = false) {
     static unordered_map<string, pair<Elf64_Addr, Elf_Xword>> r;
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
+        if (rebuild) {
+            r.clear();
+        }
+
         dl_iterate_phdr(
             [](dl_phdr_info* info, size_t, void*) {
                 static constexpr const char self[] = "/proc/self/exe";
@@ -108,7 +112,12 @@ symbol_addresses() {
                 return 0;
             },
             nullptr);
-    });
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        cons();
+    }
 
     return r;
 }
@@ -166,13 +175,18 @@ vector<char> code_object_blob_for_process() {
     return r;
 }
 
-const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs() {
+const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool rebuild = false) {
     static unordered_map<hsa_isa_t, vector<vector<char>>> r;
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
         static vector<vector<char>> blobs{code_object_blob_for_process()};
 
+        if (rebuild) {
+            blobs.clear();
+            blobs.push_back(code_object_blob_for_process());
+        }
+
         dl_iterate_phdr(
             [](dl_phdr_info* info, std::size_t, void*) {
                 elfio tmp;
@@ -194,7 +208,13 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs() {
                 }
             }
         }
-    });
+    };
+
+
+    call_once(f, cons);
+    if (rebuild) {
+        cons();
+    }
 
     return r;
 }
@@ -216,13 +236,13 @@ vector<pair<uintptr_t, string>> function_names_for(const elfio& reader, section*
     return r;
 }
 
-const vector<pair<uintptr_t, string>>& function_names_for_process() {
+const vector<pair<uintptr_t, string>>& function_names_for_process(bool rebuild = false) {
     static constexpr const char self[] = "/proc/self/exe";
 
     static vector<pair<uintptr_t, string>> r;
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
         elfio reader;
 
         if (!reader.load(self)) {
@@ -233,16 +253,26 @@ const vector<pair<uintptr_t, string>>& function_names_for_process() {
             find_section_if(reader, [](const section* x) { return x->get_type() == SHT_SYMTAB; });
 
         if (symtab) r = function_names_for(reader, symtab);
-    });
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        cons();
+    }
 
     return r;
 }
 
-const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels() {
+const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels(bool rebuild = false) {
     static unordered_map<string, vector<hsa_executable_symbol_t>> r;
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
+        if (rebuild) {
+            r.clear();
+            executables(rebuild);
+        }
+
         static const auto copy_kernels = [](hsa_executable_t, hsa_agent_t,
                                             hsa_executable_symbol_t s, void*) {
             if (type(s) == HSA_SYMBOL_KIND_KERNEL) r[name(s)].push_back(s);
@@ -256,7 +286,12 @@ const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels() {
                                                      copy_kernels, nullptr);
             }
         }
-    });
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        cons();
+    }
 
     return r;
 }
@@ -295,13 +330,18 @@ void load_code_object_and_freeze_executable(
 
 namespace hip_impl {
 const unordered_map<hsa_agent_t, vector<hsa_executable_t>>&
-executables() {  // TODO: This leaks the hsa_executable_ts, it should use RAII.
+executables(bool rebuild) {  // TODO: This leaks the hsa_executable_ts, it should use RAII.
     static unordered_map<hsa_agent_t, vector<hsa_executable_t>> r;
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
         static const auto accelerators = hc::accelerator::get_all();
 
+        if (rebuild) {
+            r.clear();
+            code_object_blobs(rebuild);
+        }
+
         for (auto&& acc : accelerators) {
             auto agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
 
@@ -335,17 +375,29 @@ executables() {  // TODO: This leaks the hsa_executable_ts, it should use RAII.
                                    },
                                    agent);
         }
-    });
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        cons();
+    }
 
     return r;
 }
 
-const unordered_map<uintptr_t, string>& function_names() {
+const unordered_map<uintptr_t, string>& function_names(bool rebuild) {
     static unordered_map<uintptr_t, string> r{function_names_for_process().cbegin(),
                                               function_names_for_process().cend()};
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
+        if (rebuild) {
+            r.clear();
+            function_names_for_process(rebuild);
+            r.insert(function_names_for_process().cbegin(),
+                     function_names_for_process().cend());
+        }
+
         dl_iterate_phdr(
             [](dl_phdr_info* info, size_t, void*) {
                 elfio tmp;
@@ -365,16 +417,30 @@ const unordered_map<uintptr_t, string>& function_names() {
                 return 0;
             },
             nullptr);
-    });
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        static mutex mtx;
+        lock_guard<mutex> lck{mtx};
+        cons();
+    }
 
     return r;
 }
 
-const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& functions() {
+const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& functions(bool rebuild) {
     static unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>> r;
     static once_flag f;
 
-    call_once(f, []() {
+    auto cons = [rebuild]() {
+        if (rebuild) {
+            r.clear();
+            function_names(rebuild);
+            kernels(rebuild);
+            globals(rebuild);
+        }
+
         for (auto&& function : function_names()) {
             const auto it = kernels().find(function.second);
 
@@ -386,15 +452,34 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
                 }
             }
         }
-    });
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        static mutex mtx;
+        lock_guard<mutex> lck{mtx};
+        cons();
+    }
 
     return r;
 }
 
-unordered_map<string, void*>& globals() {
+unordered_map<string, void*>& globals(bool rebuild) {
     static unordered_map<string, void*> r;
     static once_flag f;
-    call_once(f, []() { r.reserve(symbol_addresses().size()); });
+    auto cons =[rebuild]() {
+        if (rebuild) {
+            r.clear();
+            symbol_addresses(rebuild);
+        }
+
+        r.reserve(symbol_addresses().size());
+    };
+
+    call_once(f, cons);
+    if (rebuild) {
+        cons();
+    }
 
     return r;
 }

From 6815c7bdd71626ef5a23e78ed30682a962a83951 Mon Sep 17 00:00:00 2001
From: Nico Trost <nico.trost@amd.com>
Date: Thu, 14 Jun 2018 21:49:54 +0200
Subject: [PATCH 08/11] added missing hipCmul() to nvcc_detail/hip_complex.h

---
 hipamd/include/hip/nvcc_detail/hip_complex.h | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/hipamd/include/hip/nvcc_detail/hip_complex.h b/hipamd/include/hip/nvcc_detail/hip_complex.h
index ea2ddfdf4e..d0e45d26db 100644
--- a/hipamd/include/hip/nvcc_detail/hip_complex.h
+++ b/hipamd/include/hip/nvcc_detail/hip_complex.h
@@ -83,6 +83,10 @@ __device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, h
     return cuCsub(p, q);
 }
 
+__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
+    return cuCmul(p, q);
+}
+
 __device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
     return cuCdiv(p, q);
 }

From 007e2a4b5f4804f4bf159b766721847a51ba7788 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Fri, 15 Jun 2018 16:08:29 +0530
Subject: [PATCH 09/11] TEMP- fix memcpy2dAsync for trsm issue

---
 hipamd/src/hip_memory.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index d6c04ae98c..787e49683b 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -1680,9 +1680,12 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
             actualDest = pinnedPtr;
           }
     }
+#if 0
     if((width == dpitch) && (width == spitch)) {
             hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
-    } else {
+    } else
+#endif
+   {
         try {
             if(!isLocked){
                 for (int i = 0; i < height; ++i) 

From b883ea759d08e4bfae509f27f9fc0cfa2b0622e7 Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Fri, 15 Jun 2018 16:45:03 -0500
Subject: [PATCH 10/11] Improve performance of re-initialization logic

Keep track of shared libaries already discovered. Do not build HSA executables
for them.
---
 hipamd/src/program_state.cpp | 17 +++++++++++++----
 1 file changed, 13 insertions(+), 4 deletions(-)

diff --git a/hipamd/src/program_state.cpp b/hipamd/src/program_state.cpp
index 3ad9285b8c..3eba84f12e 100644
--- a/hipamd/src/program_state.cpp
+++ b/hipamd/src/program_state.cpp
@@ -180,21 +180,29 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool reb
     static once_flag f;
 
     auto cons = [rebuild]() {
+        // names of shared libraries who .kernel sections already loaded
+        static unordered_set<string> lib_names;
         static vector<vector<char>> blobs{code_object_blob_for_process()};
 
         if (rebuild) {
+            r.clear();
             blobs.clear();
-            blobs.push_back(code_object_blob_for_process());
         }
 
         dl_iterate_phdr(
             [](dl_phdr_info* info, std::size_t, void*) {
                 elfio tmp;
-                if (tmp.load(info->dlpi_name)) {
+                if ((lib_names.find(info->dlpi_name) == lib_names.end()) &&
+                    (tmp.load(info->dlpi_name))) {
                     const auto it = find_section_if(
                         tmp, [](const section* x) { return x->get_name() == ".kernel"; });
 
-                    if (it) blobs.emplace_back(it->get_data(), it->get_data() + it->get_size());
+                    if (it) {
+                        blobs.emplace_back(
+                            it->get_data(), it->get_data() + it->get_size());
+                        // register the shared library as already loaded
+                        lib_names.emplace(info->dlpi_name);
+                    }
                 }
                 return 0;
             },
@@ -338,7 +346,8 @@ executables(bool rebuild) {  // TODO: This leaks the hsa_executable_ts, it shoul
         static const auto accelerators = hc::accelerator::get_all();
 
         if (rebuild) {
-            r.clear();
+            // do NOT clear r so we reuse instances of hsa_executable_t
+            // created previously
             code_object_blobs(rebuild);
         }
 

From 8f521edff17afc3daec7ae3d5235096eebde9c04 Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Mon, 18 Jun 2018 16:54:18 -0500
Subject: [PATCH 11/11] Keep the map which tracks GPU kernel symbols to grow
 monotonically

---
 hipamd/src/program_state.cpp | 4 +++-
 1 file changed, 3 insertions(+), 1 deletion(-)

diff --git a/hipamd/src/program_state.cpp b/hipamd/src/program_state.cpp
index 3eba84f12e..fab71b1def 100644
--- a/hipamd/src/program_state.cpp
+++ b/hipamd/src/program_state.cpp
@@ -444,7 +444,9 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
 
     auto cons = [rebuild]() {
         if (rebuild) {
-            r.clear();
+            // do NOT clear r so we reuse instances of pair<hsa_agent_t, Kernel_descriptor>
+            // created previously
+
             function_names(rebuild);
             kernels(rebuild);
             globals(rebuild);