Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_native_vector_types

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>() {

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>
 
 extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int);

include/hip/hcc_detail/hip_atomic.h

@@ -0,0 +1,269 @@
+#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*>(address)};
+    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*>(address)};
+    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)
+{
+    __device__
+    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)
+{
+    __device__
+    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.

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);
 }

include/hip/hcc_detail/hip_runtime.h

@@ -103,6 +103,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>
@@ -194,82 +195,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);

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);

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

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);
 }