Merge 'master' into 'amd-master'

Change-Id: I640c6dfa3523f006ad0c2e05b325017609192484

CMakeLists.txt

@@ -271,6 +271,21 @@ if(HIP_PLATFORM STREQUAL "hcc")
         target_link_libraries(hiprtc PUBLIC stdc++fs)
     endif()
 
+
+    if(HIP_PLATFORM STREQUAL "hcc")
+        find_package(amd_comgr REQUIRED CONFIG
+            PATHS
+                /opt/rocm/
+            PATH_SUFFIXES
+                cmake/amd_comgr
+                lib/cmake/amd_comgr
+        )
+        MESSAGE(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
+    endif()
+
+    target_link_libraries(hip_hcc PRIVATE amd_comgr)
+    target_link_libraries(hip_hcc_static PRIVATE amd_comgr)
+
     string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
     foreach(TARGET hip_hcc hip_hcc_static)
         target_include_directories(${TARGET} SYSTEM INTERFACE $<INSTALL_INTERFACE:$<INSTALL_PREFIX>/include>;${HSA_PATH}/include)

bin/hipcc

@@ -108,7 +108,7 @@ $HIP_RUNTIME= $hipConfig{'HIP_RUNTIME'};
 # If using VDI runtime, need to find HIP_VDI_HOME
 if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
     my $hipcc_dir = dirname($0);
-    if (-e "$hipcc_dir/.hipVersion") {
+    if (-e "$hipcc_dir/../lib/bitcode") {
         $HIP_VDI_HOME = abs_path($hipcc_dir . "/..");
     } else {
         $HIP_VDI_HOME = "/opt/rocm/hip";

docs/markdown/hip_kernel_language.md

@@ -793,10 +793,13 @@ hipcc now supports compiling C++/HIP kernels to binary code objects.
 The user can specify the target for which the binary can be generated. HIP/HCC does not yet support fat binaries so only a single target may be specified.
 The file format for binary is `.co` which means Code Object. The following command builds the code object using `hipcc`.
 
-`hipcc --genco --target-isa=[TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]`
-```[TARGET GPU] = gfx803/gfx701
+`hipcc --genco --targets [TARGET GPU] [INPUT FILE] -o [OUTPUT FILE]`
+
+```
+[TARGET GPU] = gfx900 gfx803 gfx701
 [INPUT FILE] = Name of the file containing kernels
-[OUTPUT FILE] = Name of the generated code object file```
+[OUTPUT FILE] = Name of the generated code object file
+```
 
 Note that one important fact to remember when using binary code objects is that the number of arguments to the kernel are different on HCC and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel.
 

include/hip/hcc_detail/device_functions.h

@@ -1007,7 +1007,7 @@ void __syncthreads()
    SIZE         15:11   Range: 1..32
  */
 
-#define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG
+#define GETREG_IMMED(SZ,OFF,REG) (((SZ) << 11) | ((OFF) << 6) | (REG))
 
 /*
   __smid returns the wave's assigned Compute Unit and Shader Engine.

include/hip/hcc_detail/functional_grid_launch.hpp

@@ -22,7 +22,6 @@ THE SOFTWARE.
 
 #pragma once
 
-#include "code_object_bundle.hpp"
 #include "concepts.hpp"
 #include "helpers.hpp"
 #include "program_state.hpp"
@@ -57,10 +56,10 @@ template <
     std::size_t n,
     typename... Ts,
     typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
-inline std::vector<std::uint8_t> make_kernarg(
+inline hip_impl::kernarg make_kernarg(
     const std::tuple<Ts...>&,
     const kernargs_size_align&,
-    std::vector<std::uint8_t> kernarg) {
+    hip_impl::kernarg kernarg) {
     return kernarg;
 }
 
@@ -68,10 +67,10 @@ template <
     std::size_t n,
     typename... Ts,
     typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
-inline std::vector<std::uint8_t> make_kernarg(
+inline hip_impl::kernarg make_kernarg(
     const std::tuple<Ts...>& formals,
     const kernargs_size_align& size_align,
-    std::vector<std::uint8_t> kernarg) {
+    hip_impl::kernarg kernarg) {
     using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
 
     static_assert(
@@ -96,7 +95,7 @@ inline std::vector<std::uint8_t> make_kernarg(
 }
 
 template <typename... Formals, typename... Actuals>
-inline std::vector<std::uint8_t> make_kernarg(
+inline hip_impl::kernarg make_kernarg(
     void (*kernel)(Formals...), std::tuple<Actuals...> actuals) {
     static_assert(sizeof...(Formals) == sizeof...(Actuals),
         "The count of formal arguments must match the count of actuals.");
@@ -104,7 +103,7 @@ inline std::vector<std::uint8_t> make_kernarg(
     if (sizeof...(Formals) == 0) return {};
 
     std::tuple<Formals...> to_formals{std::move(actuals)};
-    std::vector<std::uint8_t> kernarg;
+    hip_impl::kernarg kernarg;
     kernarg.reserve(sizeof(to_formals));
 
     auto& ps = hip_impl::get_program_state();

include/hip/hcc_detail/helpers.hpp

@@ -84,12 +84,21 @@ struct is_callable_impl<F(Ts...), 4u, void_t_<decltype(std::declval<F>()(std::de
 // Not callable.
 template <FunctionalProcedure F>
 struct is_callable_impl<F, 5u> : std::false_type {};
-#else
+#elif  (__cplusplus < 201703L)
 template <typename, typename = void>
 struct is_callable_impl : std::false_type {};
 
 template <FunctionalProcedure F, typename... Ts>
 struct is_callable_impl<F(Ts...), void_t_<std::result_of_t<F(Ts...)> > > : std::true_type {};
+#else
+
+// C++17
+
+template <typename, typename = void>
+struct is_callable_impl : std::false_type {};
+
+template <FunctionalProcedure F, typename... Ts>
+struct is_callable_impl<F(Ts...), void_t_<std::invoke_result<F(Ts...)> > > : std::true_type {};
 #endif
 template <typename Call>
 struct is_callable : is_callable_impl<Call> {};

include/hip/hcc_detail/hip_complex.h

@@ -120,51 +120,102 @@ THE SOFTWARE.
         ret.y = lhs.y * rhs;                                                                       \
         return ret;                                                                                \
     }
-#define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T)                                      \
-    explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {}                               \
-    __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {}
 
 #endif
 
-struct hipFloatComplex {
-#ifdef __cplusplus
-   public:
-   typedef float value_type;
-    __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {}
-    explicit __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {}
-    __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {}
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned int)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed int)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, double)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned long long)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed long long)
-#endif
-    float x, y;
-} __attribute__((aligned(8)));
+typedef float2 hipFloatComplex;
+
+__device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
+
+__device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
+
+__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
+    hipFloatComplex z;
+    z.x = a;
+    z.y = b;
+    return z;
+}
+
+__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
+    hipFloatComplex ret;
+    ret.x = z.x;
+    ret.y = -z.y;
+    return ret;
+}
+
+__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
+    return z.x * z.x + z.y * z.y;
+}
+
+__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
+    return make_hipFloatComplex(p.x + q.x, p.y + q.y);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
+    return make_hipFloatComplex(p.x - q.x, p.y - q.y);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
+    return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
+}
+
+__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
+    float sqabs = hipCsqabsf(q);
+    hipFloatComplex ret;
+    ret.x = (p.x * q.x + p.y * q.y) / sqabs;
+    ret.y = (p.y * q.x - p.x * q.y) / sqabs;
+    return ret;
+}
+
+__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
+
+
+typedef double2 hipDoubleComplex;
+
+__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
+
+__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
+
+__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
+    hipDoubleComplex z;
+    z.x = a;
+    z.y = b;
+    return z;
+}
+
+__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
+    hipDoubleComplex ret;
+    ret.x = z.x;
+    ret.y = z.y;
+    return ret;
+}
+
+__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
+    return z.x * z.x + z.y * z.y;
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
+    return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
+    return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
+    return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
+}
+
+__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
+    double sqabs = hipCsqabs(q);
+    hipDoubleComplex ret;
+    ret.x = (p.x * q.x + p.y * q.y) / sqabs;
+    ret.y = (p.y * q.x - p.x * q.y) / sqabs;
+    return ret;
+}
+
+__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); }
 
-struct hipDoubleComplex {
-#ifdef __cplusplus
-   public:
-   typedef double value_type;
-    __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {}
-    explicit __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {}
-    __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {}
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned int)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed int)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, float)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned long long)
-    MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed long long)
-#endif
-    double x, y;
-} __attribute__((aligned(16)));
 
 #if __cplusplus
 
@@ -214,93 +265,6 @@ COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
 
 #endif
 
-__device__ __host__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
-
-__device__ __host__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
-
-__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
-    hipFloatComplex z;
-    z.x = a;
-    z.y = b;
-    return z;
-}
-
-__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
-    hipFloatComplex ret;
-    ret.x = z.x;
-    ret.y = -z.y;
-    return ret;
-}
-
-__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z) {
-    return z.x * z.x + z.y * z.y;
-}
-
-__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
-    return make_hipFloatComplex(p.x + q.x, p.y + q.y);
-}
-
-__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
-    return make_hipFloatComplex(p.x - q.x, p.y - q.y);
-}
-
-__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
-    return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
-}
-
-__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
-    float sqabs = hipCsqabsf(q);
-    hipFloatComplex ret;
-    ret.x = (p.x * q.x + p.y * q.y) / sqabs;
-    ret.y = (p.y * q.x - p.x * q.y) / sqabs;
-    return ret;
-}
-
-__device__ __host__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
-
-__device__ __host__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
-
-__device__ __host__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
-
-__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
-    hipDoubleComplex z;
-    z.x = a;
-    z.y = b;
-    return z;
-}
-
-__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
-    hipDoubleComplex ret;
-    ret.x = z.x;
-    ret.y = z.y;
-    return ret;
-}
-
-__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z) {
-    return z.x * z.x + z.y * z.y;
-}
-
-__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
-    return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
-}
-
-__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
-    return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
-}
-
-__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
-    return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
-}
-
-__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
-    double sqabs = hipCsqabs(q);
-    hipDoubleComplex ret;
-    ret.x = (p.x * q.x + p.y * q.y) / sqabs;
-    ret.y = (p.y * q.x - p.x * q.y) / sqabs;
-    return ret;
-}
-
-__device__ __host__ static inline double hipCabs(hipDoubleComplex z) { return sqrtf(hipCsqabs(z)); }
 
 typedef hipFloatComplex hipComplex;
 

include/hip/hcc_detail/hip_runtime_api.h

@@ -2599,180 +2599,41 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con
 
 hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
 
-struct Agent_global {
-
-    Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
-    Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) 
-      : name(nullptr), address(address), byte_cnt(byte_cnt) {
-      if (name)
-        this->name = strdup(name);
-    }
-
-    Agent_global& operator=(Agent_global&& t) {
-      if (this == &t) return *this;
-
-      if (name) free(name);
-      name = t.name;
-      address = t.address;
-      byte_cnt = t.byte_cnt;
-
-      t.name = nullptr;
-      t.address = nullptr;
-      t.byte_cnt = 0;
-
-      return *this;
-    }
-
-    Agent_global(Agent_global&& t) 
-      : name(nullptr), address(nullptr), byte_cnt(0) {
-      *this = std::move(t);
-    }
-
-    // not needed, delete them to prevent bugs
-    Agent_global(const Agent_global&) = delete;
-    Agent_global& operator=(Agent_global& t) = delete;
-
-    ~Agent_global() { if (name) free(name); }
-    
-    char* name;
-    hipDeviceptr_t address;
-    uint32_t byte_cnt;
-};
-
 #if !__HIP_VDI__
 #if defined(__cplusplus)
 } // extern "C"
 #endif
 
 namespace hip_impl {
-hsa_executable_t executable_for(hipModule_t);
-const char* hash_for(hipModule_t);
+    class agent_globals_impl;
+    class agent_globals {
+        public:
+            agent_globals();
+            ~agent_globals();
+            agent_globals(const agent_globals&) = delete;
 
-template<typename ForwardIterator>
-std::pair<hipDeviceptr_t, std::size_t> read_global_description(
-    ForwardIterator f, ForwardIterator l, const char* name) {
-    const auto it = std::find_if(f, l, [=](const Agent_global& x) {
-        return strcmp(x.name, name) == 0;
-    });
+            hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                    hipModule_t hmod, const char* name);
+            hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                    const char* name);
+        private:
+            agent_globals_impl* impl;
+    };
 
-    return it == l ?
-        std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
-}
-
-std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
-                                             hsa_executable_t executable);
-hsa_agent_t this_agent();
-
-
-class agent_globals_impl {
-private:
-    std::pair<
-        std::mutex,
-        std::unordered_map<
-            std::string, std::vector<Agent_global>>> globals_from_module;
-
-    std::unordered_map<
-        hsa_agent_t,
-        std::pair<
-            std::once_flag,
-            std::vector<Agent_global>>> globals_from_process;
-
-public:
-
-    hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
-            hipModule_t hmod, const char* name) {
-        // the key of the map would the hash of code object associated with the
-        // hipModule_t instance
-        std::string key(hash_for(hmod));
-
-        if (globals_from_module.second.count(key) == 0) {
-            std::lock_guard<std::mutex> lck{globals_from_module.first};
-
-            if (globals_from_module.second.count(key) == 0) {
-                globals_from_module.second.emplace(
-                        key, read_agent_globals(this_agent(), executable_for(hmod)));
-            }
-        }
-
-        const auto it0 = globals_from_module.second.find(key);
-        if (it0 == globals_from_module.second.cend()) {
-            hip_throw(
-                    std::runtime_error{"agent_globals data structure corrupted."});
-        }
-
-        std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
-                it0->second.cend(), name);
-
-        return *dptr ? hipSuccess : hipErrorNotFound;
+    inline
+    __attribute__((visibility("hidden")))
+    agent_globals& get_agent_globals() {
+        static agent_globals ag;
+        return ag;
     }
 
+    extern "C"
+    inline
+    __attribute__((visibility("hidden")))
     hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
-            const char* name) {
-
-        auto agent = this_agent();
-
-        std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
-            std::vector<Agent_global> tmp0;
-            for (auto&& executable : hip_impl::get_program_state().executables(aa)) {
-                auto tmp1 = read_agent_globals(aa, executable);
-                tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
-                            make_move_iterator(tmp1.end()));
-            }
-            globals_from_process[aa].second = move(move(tmp0));
-        }, agent);
-
-        const auto it = globals_from_process.find(agent);
-
-        if (it == globals_from_process.cend()) return hipErrorNotInitialized;
-
-        std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
-                it->second.second.cend(), name);
-
-        return *dptr ? hipSuccess : hipErrorNotFound;
+        const char* name) {
+        return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
     }
-  
-};
-
-class agent_globals {
-public:
-    agent_globals() : impl(new agent_globals_impl()) { 
-        if (!impl) 
-            hip_throw(
-                std::runtime_error{"Error when constructing agent global data structures."});
-    }
-    ~agent_globals() { delete impl; }
-
-    hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
-                                             hipModule_t hmod, const char* name) {
-        return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
-    }
-
-    hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
-                                              const char* name) {
-        return impl->read_agent_global_from_process(dptr, bytes, name);
-    }
-
-private:
-    agent_globals_impl* impl;
-};
-
-inline
-__attribute__((visibility("hidden")))
-agent_globals& get_agent_globals() {
-    static agent_globals ag;
-    return ag;
-}
-
-
-extern "C"
-inline
-__attribute__((visibility("hidden")))
-hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
-                                          const char* name) {
-    return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
-}
-
-
 } // Namespace hip_impl.
 
 #if defined(__cplusplus)
@@ -2906,6 +2767,21 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
 hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
    int* numBlocks, const void* f, int  blockSize, size_t dynamicSMemSize, unsigned int flags);
 
+/**
+ * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched
+ * on respective streams before enqueuing any other work on the specified streams from any other threads
+ *
+ *
+ * @param [in] hipLaunchParams          List of launch parameters, one per device.
+ * @param [in] numDevices               Size of the launchParamsList array.
+ * @param [in] flags                    Flags to control launch behavior.
+ *
+ * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
+ */
+hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
+                                              int  numDevices, unsigned int  flags);
+
+
 
 // doxygen end Version Management
 /**
@@ -3262,6 +3138,12 @@ inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchP
     return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags);
 }
 
+template <class T>
+inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
+                                                     unsigned int  numDevices, unsigned int  flags = 0) {
+    return hipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags);
+}
+
 
 /*
  * @brief Unbinds the textuer bound to @p tex

include/hip/hcc_detail/hip_texture_types.h

@@ -45,10 +45,15 @@ THE SOFTWARE.
  *                                                                              *
  *                                                                              *
  *******************************************************************************/
+#if __HIP__
+#define __HIP_TEXTURE_ATTRIB __attribute__((device_builtin_texture_type))
+#else
+#define __HIP_TEXTURE_ATTRIB
+#endif
 
 template <class T, int texType = hipTextureType1D,
           enum hipTextureReadMode mode = hipReadModeElementType>
-struct texture : public textureReference {
+struct __HIP_TEXTURE_ATTRIB texture : public textureReference {
     texture(int norm = 0, enum hipTextureFilterMode fMode = hipFilterModePoint,
             enum hipTextureAddressMode aMode = hipAddressModeClamp) {
         normalized = norm;

include/hip/hcc_detail/hip_vector_types.h

@@ -114,7 +114,7 @@ THE SOFTWARE.
             typename U,
             typename std::enable_if<
                 std::is_convertible<U, T>{}>::type* = nullptr>
-        inline __host__ __device__
+        explicit inline __host__ __device__
         HIP_vector_type(U x) noexcept
         {
             for (auto i = 0u; i != rank; ++i) data[i] = x;
@@ -286,14 +286,14 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator+(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} += y;
+        return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
     inline __host__ __device__
     HIP_vector_type<T, n> operator+(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
-        return y + x;
+        return HIP_vector_type<T, n>{x} += y;
     }
 
     template<typename T, unsigned int n>
@@ -308,7 +308,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator-(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} -= y;
+        return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
     inline __host__ __device__
@@ -330,14 +330,14 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator*(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} *= y;
+        return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
     inline __host__ __device__
     HIP_vector_type<T, n> operator*(
         U x, const HIP_vector_type<T, n>& y) noexcept
     {
-        return y * x;
+        return HIP_vector_type<T, n>{x} *= y;
     }
 
     template<typename T, unsigned int n>
@@ -352,7 +352,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator/(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} /= y;
+        return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
     }
     template<typename T, unsigned int n, typename U>
     inline __host__ __device__
@@ -423,7 +423,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator%(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} %= y;
+        return HIP_vector_type<T, n>{x} %= HIP_vector_type<T, n>{y};
     }
     template<
         typename T,
@@ -456,7 +456,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator^(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} ^= y;
+        return HIP_vector_type<T, n>{x} ^= HIP_vector_type<T, n>{y};
     }
     template<
         typename T,
@@ -489,7 +489,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator|(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} |= y;
+        return HIP_vector_type<T, n>{x} |= HIP_vector_type<T, n>{y};
     }
     template<
         typename T,
@@ -522,7 +522,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator&(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} &= y;
+        return HIP_vector_type<T, n>{x} &= HIP_vector_type<T, n>{y};
     }
     template<
         typename T,
@@ -555,7 +555,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator>>(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} >>= y;
+        return HIP_vector_type<T, n>{x} >>= HIP_vector_type<T, n>{y};
     }
     template<
         typename T,
@@ -588,7 +588,7 @@ THE SOFTWARE.
     HIP_vector_type<T, n> operator<<(
         const HIP_vector_type<T, n>& x, U y) noexcept
     {
-        return HIP_vector_type<T, n>{x} <<= y;
+        return HIP_vector_type<T, n>{x} <<= HIP_vector_type<T, n>{y};
     }
     template<
         typename T,
@@ -668,20 +668,20 @@ __MAKE_VECTOR_TYPE__(double, double);
 
 #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \
     static inline __device__ __host__ \
-    type make_##type(comp x) { type r = {x}; return r; }
+    type make_##type(comp x) { type r{x}; return r; }
 
 #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \
     static inline __device__ __host__ \
-    type make_##type(comp x, comp y) { type r = {x, y}; return r; }
+    type make_##type(comp x, comp y) { type r{x, y}; return r; }
 
 #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \
     static inline __device__ __host__ \
-    type make_##type(comp x, comp y, comp z) { type r = {x, y, z}; return r; }
+    type make_##type(comp x, comp y, comp z) { type r{x, y, z}; return r; }
 
 #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \
     static inline __device__ __host__ \
     type make_##type(comp x, comp y, comp z, comp w) { \
-        type r = {x, y, z, w}; \
+        type r{x, y, z, w}; \
         return r; \
     }
 

include/hip/hcc_detail/program_state.hpp

@@ -30,41 +30,25 @@ THE SOFTWARE.
 #include <cstddef>
 #include <cstdint>
 #include <cstdlib>
-#include <stdexcept>
-#include <string>
-#include <unordered_map>
-#include <vector>
 
 struct ihipModuleSymbol_t;
 using hipFunction_t = ihipModuleSymbol_t*;
 
-namespace std {
-template<>
-struct hash<hsa_agent_t> {
-    size_t operator()(hsa_agent_t x) const {
-        return hash<decltype(x.handle)>{}(x.handle);
-    }
-};
-
-template<>
-struct hash<hsa_isa_t> {
-    size_t operator()(hsa_isa_t x) const {
-        return hash<decltype(x.handle)>{}(x.handle);
-    }
-};
-}  // namespace std
-
-inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) {
-    return x.handle == y.handle;
-}
-inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
-    return x.handle == y.handle;
-}
-
 namespace hip_impl {
 
-[[noreturn]]
-void hip_throw(const std::exception&);
+struct kernarg_impl;
+class kernarg {
+public:
+    kernarg();
+    kernarg(kernarg&&);
+    ~kernarg();
+    std::uint8_t* data();
+    std::size_t   size();
+    void reserve(std::size_t);
+    void resize(std::size_t);
+private:
+    kernarg_impl* impl;
+};
 
 class kernargs_size_align;
 class program_state_impl;
@@ -72,6 +56,7 @@ class program_state {
 public:
     program_state();
     ~program_state();
+    program_state(const program_state&) = delete;
 
     hipFunction_t kernel_descriptor(std::uintptr_t,
                                     hsa_agent_t);
@@ -83,12 +68,8 @@ public:
 
     void* global_addr_by_name(const char* name);
 
-    // to fix later
-    const std::vector<hsa_executable_t>& executables(hsa_agent_t agent);
-
-    program_state(const program_state&) = delete;
-
 private:
+    friend class agent_globals_impl;
     program_state_impl* impl;
 };
 

src/hip_hcc.cpp

@@ -2492,13 +2492,4 @@ namespace hip_impl {
         #endif
     }
 
-    std::mutex executables_cache_mutex;
-
-    std::vector<hsa_executable_t>& executables_cache(
-            std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
-        static std::unordered_map<std::string,
-            std::unordered_map<hsa_isa_t,
-                std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
-        return cache[elf][isa][agent];
-    }
 } // Namespace hip_impl.

src/hip_module.cpp

@@ -55,6 +55,22 @@ THE SOFTWARE.
 using namespace ELFIO;
 using namespace std;
 
+// For HIP implicit kernargs.
+static const size_t HIP_IMPLICIT_KERNARG_SIZE = 48;
+static const size_t HIP_IMPLICIT_KERNARG_ALIGNMENT = 8;
+
+struct amd_kernel_code_v3_t {
+  uint32_t group_segment_fixed_size;
+  uint32_t private_segment_fixed_size;
+  uint8_t reserved0[8];
+  int64_t kernel_code_entry_byte_offset;
+  uint8_t reserved1[24];
+  uint32_t compute_pgm_rsrc1;
+  uint32_t compute_pgm_rsrc2;
+  uint16_t kernel_code_properties;
+  uint8_t reserved2[6];
+};
+
 // calculate MD5 checksum
 inline std::string checksum(size_t size, const char *source) {
     // FNV-1a hashing, 64-bit version
@@ -146,33 +162,28 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         ihipDevice_t* currentDevice = ihipGetDevice(deviceId);
         hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
 
-        void* config[5] = {0};
-        size_t kernArgSize;
-
-        std::vector<char> tmp{};
+        std::vector<char> kernargs{};
         if (kernelParams) {
             if (extra) return hipErrorInvalidValue;
 
             for (auto&& x : f->_kernarg_layout) {
                 const auto p{static_cast<const char*>(*kernelParams)};
 
-                tmp.insert(
-                    tmp.cend(),
+                kernargs.insert(
+                    kernargs.cend(),
                     round_up_to_next_multiple_nonnegative(
-                        tmp.size(), x.second) - tmp.size(),
+                        kernargs.size(), x.second) - kernargs.size(),
                     '\0');
-                tmp.insert(tmp.cend(), p, p + x.first);
+                kernargs.insert(kernargs.cend(), p, p + x.first);
 
                 ++kernelParams;
             }
-            config[1] = static_cast<void*>(tmp.data());
-
-            kernArgSize = tmp.size();
         } else if (extra) {
-            memcpy(config, extra, sizeof(size_t) * 5);
-            if (config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER &&
-                config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END) {
-                kernArgSize = *(size_t*)(config[3]);
+            if (extra[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER &&
+                extra[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && extra[4] == HIP_LAUNCH_PARAM_END) {
+                auto args = (char*)extra[1];
+                size_t argSize = *(size_t*)(extra[3]);
+                kernargs.insert(kernargs.end(), args, args+argSize);
             } else {
                 return hipErrorNotInitialized;
             }
@@ -181,6 +192,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
             return hipErrorInvalidValue;
         }
 
+        // Insert 48-bytes at the end for implicit kernel arguments and fill with value zero.
+        size_t padSize = (~kernargs.size() + 1) & (HIP_IMPLICIT_KERNARG_ALIGNMENT - 1);
+        kernargs.insert(kernargs.end(), padSize + HIP_IMPLICIT_KERNARG_SIZE, 0);
 
         /*
           Kernel argument preparation.
@@ -189,7 +203,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         lp.dynamic_group_mem_bytes =
             sharedMemBytes;  // TODO - this should be part of preLaunchKernel.
         hStream = ihipPreLaunchKernel(
-            hStream, dim3(globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ),
+            hStream, dim3(globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ),
             dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str());
 
 
@@ -206,10 +220,20 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         aql.grid_size_x = globalWorkSizeX;
         aql.grid_size_y = globalWorkSizeY;
         aql.grid_size_z = globalWorkSizeZ;
-        aql.group_segment_size =
-            f->_header->workgroup_group_segment_byte_size + sharedMemBytes;
-        aql.private_segment_size =
-            f->_header->workitem_private_segment_byte_size;
+        bool is_code_object_v3 = f->_name.find(".kd") != std::string::npos;
+        if (is_code_object_v3) {
+            const auto* header =
+                reinterpret_cast<const amd_kernel_code_v3_t*>(f->_header);
+            aql.group_segment_size =
+                header->group_segment_fixed_size + sharedMemBytes;
+            aql.private_segment_size =
+                header->private_segment_fixed_size;
+        } else {
+            aql.group_segment_size =
+                f->_header->workgroup_group_segment_byte_size + sharedMemBytes;
+            aql.private_segment_size =
+                f->_header->workitem_private_segment_byte_size;
+        }
         aql.kernel_object = f->_object;
         aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
         aql.header =
@@ -230,7 +254,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
 
         hc::completion_future cf;
 
-        lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize,
+        lp.av->dispatch_hsa_kernel(&aql, kernargs.data(), kernargs.size(),
                                    (startEvent || stopEvent) ? &cf : nullptr
 #if (__hcc_workweek__ > 17312)
                                    ,
@@ -289,16 +313,6 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
 }
 
-hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
-                              hipModule_t hmod, const char* name) {
-    HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
-    if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
-
-    if (!name) return hipErrorNotInitialized;
-
-    return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
-}
-
 namespace hip_impl {
     hsa_executable_t executable_for(hipModule_t hmod) {
         return hmod->executable;
@@ -323,10 +337,171 @@ namespace hip_impl {
 
         return currentDevice->_hsaAgent;
     }
+
+    struct Agent_global {
+        Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
+        Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) 
+            : name(nullptr), address(address), byte_cnt(byte_cnt) {
+                if (name)
+                    this->name = strdup(name);
+        }
+
+        Agent_global& operator=(Agent_global&& t) {
+            if (this == &t) return *this;
+
+            if (name) free(name);
+            name = t.name;
+            address = t.address;
+            byte_cnt = t.byte_cnt;
+
+            t.name = nullptr;
+            t.address = nullptr;
+            t.byte_cnt = 0;
+
+            return *this;
+        }
+
+        Agent_global(Agent_global&& t) 
+            : name(nullptr), address(nullptr), byte_cnt(0) {
+                *this = std::move(t);
+        }
+
+        // not needed, delete them to prevent bugs
+        Agent_global(const Agent_global&) = delete;
+        Agent_global& operator=(Agent_global& t) = delete;
+
+        ~Agent_global() { if (name) free(name); }
+
+        char* name;
+        hipDeviceptr_t address;
+        uint32_t byte_cnt;
+    };
+
+    template<typename ForwardIterator>
+    std::pair<hipDeviceptr_t, std::size_t> read_global_description(
+        ForwardIterator f, ForwardIterator l, const char* name) {
+        const auto it = std::find_if(f, l, [=](const Agent_global& x) {
+            return strcmp(x.name, name) == 0;
+        });
+
+        return it == l ?
+            std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
+    }
+
+    std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
+                                            hsa_executable_t executable);
+    class agent_globals_impl {
+        private:
+            std::pair<
+                std::mutex,
+                std::unordered_map<
+                    std::string, std::vector<Agent_global>>> globals_from_module;
+
+            std::unordered_map<
+                hsa_agent_t,
+                std::pair<
+                    std::once_flag,
+                    std::vector<Agent_global>>> globals_from_process;
+
+        public:
+
+            hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                    hipModule_t hmod, const char* name) {
+                // the key of the map would the hash of code object associated with the
+                // hipModule_t instance
+                std::string key(hash_for(hmod));
+
+                if (globals_from_module.second.count(key) == 0) {
+                    std::lock_guard<std::mutex> lck{globals_from_module.first};
+
+                    if (globals_from_module.second.count(key) == 0) {
+                        globals_from_module.second.emplace(
+                                key, read_agent_globals(this_agent(), executable_for(hmod)));
+                    }
+                }
+
+                const auto it0 = globals_from_module.second.find(key);
+                if (it0 == globals_from_module.second.cend()) {
+                    hip_throw(
+                            std::runtime_error{"agent_globals data structure corrupted."});
+                }
+
+                std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
+                        it0->second.cend(), name);
+                // HACK for SWDEV-173477
+                //
+                // For code objects with global symbols of length 0, ROCR runtime's fix
+                // may not be working correctly. Therefore the
+                // result from read_agent_globals() can't be trusted entirely.
+                //
+                // As a workaround to tame applications which depend on the existence of
+                // global symbols with length 0, always return hipSuccess here.
+                //
+                // This behavior shall be reverted once ROCR runtime has been fixed to
+                // address SWDEV-173477 and SWDEV-190701
+
+                //return *dptr ? hipSuccess : hipErrorNotFound;
+                return hipSuccess;                
+            }
+
+            hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                    const char* name) {
+
+                auto agent = this_agent();
+
+                std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
+                    std::vector<Agent_global> tmp0;
+                    for (auto&& executable : hip_impl::get_program_state().impl->get_executables(aa)) {
+                        auto tmp1 = read_agent_globals(aa, executable);
+                        tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
+                            make_move_iterator(tmp1.end()));
+                    }
+                    globals_from_process[aa].second = move(move(tmp0));
+                }, agent);
+
+                const auto it = globals_from_process.find(agent);
+
+                if (it == globals_from_process.cend()) return hipErrorNotInitialized;
+
+                std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
+                        it->second.second.cend(), name);
+
+                return *dptr ? hipSuccess : hipErrorNotFound;
+            }
+
+    };
+
+    agent_globals::agent_globals() : impl(new agent_globals_impl()) { 
+        if (!impl) 
+            hip_throw(
+                std::runtime_error{"Error when constructing agent global data structures."});
+    }
+    agent_globals::~agent_globals() { delete impl; }
+
+    hipError_t agent_globals::read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                                             hipModule_t hmod, const char* name) {
+        return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
+    }
+
+    hipError_t agent_globals::read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                                              const char* name) {
+        return impl->read_agent_global_from_process(dptr, bytes, name);
+    }
+
 } // Namespace hip_impl.
 
+hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
+                              hipModule_t hmod, const char* name) {
+    HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
+    if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
+
+    if (!name) return hipErrorNotInitialized;
+
+    return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
+}
+
 namespace {
-inline void track(const Agent_global& x, hsa_agent_t agent) {
+inline void track(const hip_impl::Agent_global& x, hsa_agent_t agent) {
     tprintf(DB_MEM, "  add variable '%s' with ptr=%p size=%u to tracker\n", x.name,
             x.address, x.byte_cnt);
 
@@ -347,7 +522,7 @@ inline void track(const Agent_global& x, hsa_agent_t agent) {
 
 }
 
-template <typename Container = vector<Agent_global>>
+template <typename Container = vector<hip_impl::Agent_global>>
 inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t agent,
                                                 hsa_executable_symbol_t x, void* out) {
     using namespace hip_impl;
@@ -358,7 +533,7 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag
     hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t);
 
     if (t == HSA_SYMBOL_KIND_VARIABLE) {
-        Agent_global tmp(name(x).c_str(), address(x), size(x));
+        hip_impl::Agent_global tmp(name(x).c_str(), address(x), size(x));
         static_cast<Container*>(out)->push_back(std::move(tmp));
 
         track(static_cast<Container*>(out)->back(),agent);
@@ -462,6 +637,12 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch
 
     auto kernel = find_kernel_by_name(hmod->executable, name, agent);
 
+    if (kernel.handle == 0u) {
+        std::string name_str(name);
+        name_str.append(".kd");
+        kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent);
+    }
+
     if (kernel.handle == 0u) return hipErrorNotFound;
 
     // TODO: refactor the whole ihipThisThat, which is a mess and yields the
@@ -486,7 +667,11 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod,
 }
 
 namespace {
-hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header) {
+const amd_kernel_code_v3_t *header_v3(const ihipModuleSymbol_t& kd) {
+  return reinterpret_cast<const amd_kernel_code_v3_t*>(kd._header);
+}
+
+hipFuncAttributes make_function_attributes(const ihipModuleSymbol_t& kd) {
     hipFuncAttributes r{};
 
     hipDeviceProp_t prop{};
@@ -495,16 +680,31 @@ hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header) {
     //       available per CU, therefore we hardcode it to 64 KiRegisters.
     prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024;
 
-    r.localSizeBytes = header.workitem_private_segment_byte_size;
-    r.sharedSizeBytes = header.workgroup_group_segment_byte_size;
+    bool is_code_object_v3 = kd._name.find(".kd") != std::string::npos;
+    if (is_code_object_v3) {
+        r.localSizeBytes = header_v3(kd)->private_segment_fixed_size;
+        r.sharedSizeBytes = header_v3(kd)->group_segment_fixed_size;
+    } else {
+        r.localSizeBytes = kd._header->workitem_private_segment_byte_size;
+        r.sharedSizeBytes = kd._header->workgroup_group_segment_byte_size;
+    }
     r.maxDynamicSharedSizeBytes = prop.sharedMemPerBlock - r.sharedSizeBytes;
-    r.numRegs = header.workitem_vgpr_count;
+    if (is_code_object_v3) {
+        r.numRegs = ((header_v3(kd)->compute_pgm_rsrc1 & 0x3F) + 1) << 2;
+    } else {
+        r.numRegs = kd._header->workitem_vgpr_count;
+    }
     r.maxThreadsPerBlock = r.numRegs ?
         std::min(prop.maxThreadsPerBlock, prop.regsPerBlock / r.numRegs) :
         prop.maxThreadsPerBlock;
-    r.binaryVersion =
-        header.amd_machine_version_major * 10 +
-        header.amd_machine_version_minor;
+    if (is_code_object_v3) {
+        r.binaryVersion = 0; // FIXME: should it be the ISA version or code
+                             //        object format version?
+    } else {
+        r.binaryVersion =
+            kd._header->amd_machine_version_major * 10 +
+            kd._header->amd_machine_version_minor;
+    }
     r.ptxVersion = prop.major * 10 + prop.minor; // HIP currently presents itself as PTX 3.0.
 
     return r;
@@ -520,11 +720,10 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
 
     auto agent = this_agent();
     auto kd = get_program_state().kernel_descriptor(reinterpret_cast<uintptr_t>(func), agent);
-    const auto header = kd->_header;
 
-    if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."};
+    if (!kd->_header) throw runtime_error{"Ill-formed Kernel_descriptor."};
 
-    *attr = make_function_attributes(*header);
+    *attr = make_function_attributes(*kd);
 
     return hipSuccess;
 }
@@ -555,11 +754,9 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
     (*module)->executable = get_program_state().load_executable(
                                             content.data(), content.size(), (*module)->executable,
                                             this_agent());
-    istringstream elf{content};
-    ELFIO::elfio reader;
-    if (reader.load(elf)) {
-        program_state_impl::read_kernarg_metadata(reader, (*module)->kernargs);
-    }
+
+    std::vector<char> blob(content.cbegin(), content.cend());
+    program_state_impl::read_kernarg_metadata(blob, (*module)->kernargs);
 
     // compute the hash of the code object
     (*module)->hash = checksum(content.length(), content.data());

src/program_state.cpp

@@ -1,4 +1,6 @@
 #include "../include/hip/hcc_detail/program_state.hpp"
+// contains implementation of program_state_impl
+#include "program_state.inl"
 
 #include <hsa/hsa.h>
 
@@ -7,57 +9,86 @@
 #include <unordered_map>
 #include <vector>
 
-// contains implementation of program_state_impl
-#include "program_state.inl"
-
 namespace hip_impl {
+    
+    kernarg::kernarg() : impl(new kernarg_impl) { 
+    }
 
-std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
-    return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
-}
+    kernarg::kernarg(kernarg&& k) : impl(k.impl) {
+        k.impl = nullptr;
+    }
 
-std::size_t kernargs_size_align::alignment(std::size_t n) const{
-    return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
-}
+    kernarg::~kernarg() {
+        if (impl)
+            delete(impl);
+    }
 
-program_state::program_state() :
-    impl(new program_state_impl) {
-    if (!impl) hip_throw(std::runtime_error {
-                         "Unknown error when constructing program state."});
-}
+    std::uint8_t* kernarg::data() {
+        return impl->v.data();
+    }
 
-program_state::~program_state() {
-    delete(impl);
-}
+    std::size_t kernarg::size() {
+        return impl->v.size();
+    }
 
-void* program_state::global_addr_by_name(const char* name) {
-    const auto it = impl->get_globals().find(name);
-    if (it == impl->get_globals().end())
-      return nullptr;
-    else
-      return it->second;
-}
+    void kernarg::reserve(std::size_t c) {
+        impl->v.reserve(c);
+    }
 
-hsa_executable_t program_state::load_executable(const char* data,
-                                                const size_t data_size,
-                                                hsa_executable_t executable,
-                                                hsa_agent_t agent) {
-    return impl->load_executable(data, data_size, executable, agent);
-}
+    void kernarg::resize(std::size_t c) {
+        impl->v.resize(c);
+    }
 
-const std::vector<hsa_executable_t>& program_state::executables(hsa_agent_t agent) {
-    return impl->get_executables(agent);
-}
+    std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
+        return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
+    }
 
-hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
-                                                          hsa_agent_t agent) {
-    auto& kd = impl->kernel_descriptor(function_address, agent);
-    return kd;
-}
+    std::size_t kernargs_size_align::alignment(std::size_t n) const{
+        return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
+    }
 
-kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
-  kernargs_size_align t;
-  t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
-  return t;
-}
+    program_state::program_state() : impl(new program_state_impl) {
+        if (!impl) hip_throw(std::runtime_error {
+            "Unknown error when constructing program state."});
+    }
+
+    program_state::~program_state() {
+        delete(impl);
+    }
+
+    void* program_state::global_addr_by_name(const char* name) {
+        const auto it = impl->get_globals().find(name);
+        if (it == impl->get_globals().end())
+            return nullptr;
+        else
+            return it->second;
+    }
+
+    hsa_executable_t program_state::load_executable(const char* data,
+        const size_t data_size,
+        hsa_executable_t executable,
+        hsa_agent_t agent) {
+        return impl->load_executable(data, data_size, executable, agent);
+    }
+
+    hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
+        hsa_agent_t agent) {
+        auto& kd = impl->kernel_descriptor(function_address, agent);
+        return kd;
+    }
+
+    kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
+        kernargs_size_align t;
+        t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
+        return t;
+    }
+
+    std::mutex executables_cache_mutex;
+    std::vector<hsa_executable_t>& executables_cache(
+        std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
+        static std::unordered_map<std::string,
+            std::unordered_map<hsa_isa_t,
+                std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
+        return cache[elf][isa][agent];
+    }
 };

src/program_state.inl

@@ -17,6 +17,7 @@
 #include <hsa/hsa.h>
 #include <hsa/hsa_ext_amd.h>
 #include <hsa/hsa_ven_amd_loader.h>
+#include <amd_comgr.h>
 
 #include <link.h>
 
@@ -26,12 +27,36 @@
 #include <cstdio>
 #include <memory>
 #include <mutex>
+#include <stdexcept>
 #include <string>
 #include <sstream>
 #include <unordered_map>
 #include <utility>
 #include <vector>
 
+namespace std {
+template<>
+struct hash<hsa_agent_t> {
+    size_t operator()(hsa_agent_t x) const {
+        return hash<decltype(x.handle)>{}(x.handle);
+    }
+};
+
+template<>
+struct hash<hsa_isa_t> {
+    size_t operator()(hsa_isa_t x) const {
+        return hash<decltype(x.handle)>{}(x.handle);
+    }
+};
+}  // namespace std
+
+inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) {
+    return x.handle == y.handle;
+}
+inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
+    return x.handle == y.handle;
+}
+
 namespace hip_impl {
 
 [[noreturn]]
@@ -540,9 +565,13 @@ public:
 
         std::call_once(functions[agent].first, [this](hsa_agent_t aa) {
             for (auto&& function : get_function_names()) {
-                const auto it = get_kernels(aa).find(function.second);
+                auto it = get_kernels(aa).find(function.second);
 
-                if (it == get_kernels(aa).cend()) continue;
+                if (it == get_kernels(aa).cend()) {
+                    it = get_kernels(aa).find(function.second + ".kd");
+                    if (it == get_kernels(aa).cend())
+                        continue;
+                }
 
                 for (auto&& kernel_symbol : it->second) {
                     functions[aa].second.emplace(
@@ -556,92 +585,172 @@ public:
     }
 
     static
-    std::size_t parse_args(
-            const std::string& metadata,
-            std::size_t f,
-            std::size_t l,
+    std::string metadata_to_string(const amd_comgr_metadata_node_t& md) {
+        std::string str;
+        size_t size;
+
+        if (amd_comgr_get_metadata_string(md, &size, NULL)
+            == AMD_COMGR_STATUS_SUCCESS) {
+            str.resize(size - 1);
+            amd_comgr_get_metadata_string(md, &size, &str[0]);
+        }
+        return str;
+    }
+
+    static
+    void parse_args(
+            const amd_comgr_metadata_node_t& args_md,
+            bool is_code_object_v3,
             std::vector<std::pair<std::size_t, std::size_t>>& size_align) {
-        if (f == l) return f;
-        if (!size_align.empty()) return l;
+        size_t arg_count = 0;
+        if (amd_comgr_get_metadata_list_size(args_md, &arg_count)
+            != AMD_COMGR_STATUS_SUCCESS)
+            return;
 
-        do {
-            static constexpr size_t size_sz{5};
-            f = metadata.find("Size:", f) + size_sz;
+        for (size_t i = 0; i < arg_count; ++i) {
+            amd_comgr_metadata_node_t arg_md;
 
-            if (l <= f) return f;
+            if (amd_comgr_index_list_metadata(args_md, i, &arg_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
 
-            auto size = std::strtoul(&metadata[f], nullptr, 10);
+            amd_comgr_metadata_node_t arg_size_md;
+            if (amd_comgr_metadata_lookup(arg_md,
+                                          is_code_object_v3 ? ".size" : "Size",
+                                          &arg_size_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
 
-            static constexpr size_t align_sz{6};
-            f = metadata.find("Align:", f) + align_sz;
+            size_t arg_size = std::stoul(metadata_to_string(arg_size_md));
 
-            char* l{};
-            auto align = std::strtoul(&metadata[f], &l, 10);
+            if (amd_comgr_destroy_metadata(arg_size_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
 
-            f += (l - &metadata[f]) + 1;
+            size_t arg_align;
 
-            size_align.emplace_back(size, align);
-        } while (true);
+            if (is_code_object_v3) {
+                amd_comgr_metadata_node_t arg_offset_md;
+                if (amd_comgr_metadata_lookup(arg_md, ".offset", &arg_offset_md)
+                    != AMD_COMGR_STATUS_SUCCESS)
+                    return;
+
+                size_t arg_offset
+                    = std::stoul(metadata_to_string(arg_offset_md));
+
+                if (amd_comgr_destroy_metadata(arg_offset_md)
+                    != AMD_COMGR_STATUS_SUCCESS)
+                    return;
+
+                arg_align = 1;
+                while (arg_offset && (arg_offset & 1) == 0) {
+                    arg_offset >>= 1;
+                    arg_align <<= 1;
+                }
+            } else {
+                amd_comgr_metadata_node_t arg_align_md;
+                if (amd_comgr_metadata_lookup(arg_md, "Align", &arg_align_md)
+                    != AMD_COMGR_STATUS_SUCCESS)
+                    return;
+
+                arg_align = std::stoul(metadata_to_string(arg_align_md));
+
+                if (amd_comgr_destroy_metadata(arg_align_md)
+                    != AMD_COMGR_STATUS_SUCCESS)
+                    return;
+            }
+
+            size_align.emplace_back(arg_size, arg_align);
+
+            if (amd_comgr_destroy_metadata(arg_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
+        }
     }
 
     static
     void read_kernarg_metadata(
-            ELFIO::elfio& reader,
+            const std::vector<char>& blob,
             std::unordered_map<
             std::string,
             std::vector<std::pair<std::size_t, std::size_t>>>& kernargs) {
-        // TODO: this is inefficient.
-        auto it = find_section_if(reader, [](const ELFIO::section* x) {
-                return x->get_type() == SHT_NOTE;
-                });
+        amd_comgr_data_t dataIn;
+        amd_comgr_status_t status;
 
-        if (!it) return;
+        if (amd_comgr_create_data(AMD_COMGR_DATA_KIND_RELOCATABLE, &dataIn)
+            != AMD_COMGR_STATUS_SUCCESS)
+            return;
 
-        const ELFIO::note_section_accessor acc{reader, it};
-        for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) {
-            ELFIO::Elf_Word type{};
-            std::string name{};
-            void* desc{};
-            ELFIO::Elf_Word desc_size{};
+        if (amd_comgr_set_data(dataIn, blob.size(), blob.data())
+            != AMD_COMGR_STATUS_SUCCESS)
+            return;
 
-            acc.get_note(i, type, name, desc, desc_size);
+        amd_comgr_metadata_node_t metadata;
+        if (amd_comgr_get_data_metadata(dataIn, &metadata)
+            != AMD_COMGR_STATUS_SUCCESS)
+            return;
 
-            if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
-
-            std::string tmp{
-                static_cast<char*>(desc), static_cast<char*>(desc) + desc_size};
-
-            auto dx = tmp.find("Kernels:");
-
-            if (dx == std::string::npos) continue;
-
-            static constexpr decltype(tmp.size()) kernels_sz{8};
-            dx += kernels_sz;
-
-            do {
-                dx = tmp.find("Name:", dx);
-
-                if (dx == std::string::npos) break;
-
-                static constexpr decltype(tmp.size()) name_sz{5};
-                dx = tmp.find_first_not_of(" '", dx + name_sz);
-
-                auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx);
-                dx += fn.size();
-
-                auto dx1 = tmp.find("CodeProps", dx);
-                dx = tmp.find("Args:", dx);
-
-                if (dx1 < dx) {
-                    dx = dx1;
-                    continue;
-                }
-                if (dx == std::string::npos) break;
-
-                static constexpr decltype(tmp.size()) args_sz{5};
-                dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]);
-            } while (true);
+        bool is_code_object_v3 = false;
+        amd_comgr_metadata_node_t kernels_md;
+        if (amd_comgr_metadata_lookup(metadata, "Kernels", &kernels_md)
+            != AMD_COMGR_STATUS_SUCCESS) {
+            if (amd_comgr_metadata_lookup(metadata,
+                                          "amdhsa.kernels",
+                                          &kernels_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
+            is_code_object_v3 = true;
         }
+
+        size_t kernel_count = 0;
+        if (amd_comgr_get_metadata_list_size(kernels_md, &kernel_count)
+            != AMD_COMGR_STATUS_SUCCESS)
+            return;
+
+        for (size_t i = 0; i < kernel_count; i++) {
+            amd_comgr_metadata_node_t kernel_md;
+
+            if (amd_comgr_index_list_metadata(kernels_md, i, &kernel_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                continue;
+
+            amd_comgr_metadata_node_t name_md;
+            if (amd_comgr_metadata_lookup(kernel_md,
+                                          is_code_object_v3 ? ".name" : "Name",
+                                          &name_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                continue;
+
+            std::string kernel_name_str = metadata_to_string(name_md);
+
+            if (amd_comgr_destroy_metadata(name_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                continue;
+
+            if (is_code_object_v3)
+                kernel_name_str.append(".kd");
+
+
+            amd_comgr_metadata_node_t args_md;
+            if (amd_comgr_metadata_lookup(kernel_md,
+                                          is_code_object_v3 ? ".args" : "Args",
+                                          &args_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                continue;
+
+            parse_args(args_md, is_code_object_v3, kernargs[kernel_name_str]);
+
+            if (amd_comgr_destroy_metadata(args_md) != AMD_COMGR_STATUS_SUCCESS
+                || amd_comgr_destroy_metadata(kernel_md)
+                   != AMD_COMGR_STATUS_SUCCESS)
+                continue;
+        }
+
+        if (amd_comgr_destroy_metadata(kernels_md) != AMD_COMGR_STATUS_SUCCESS
+            || amd_comgr_destroy_metadata(metadata) != AMD_COMGR_STATUS_SUCCESS)
+            return;
+
+        amd_comgr_release_data(dataIn);
     }
 
     const std::unordered_map<std::string, 
@@ -651,13 +760,7 @@ public:
             for (auto&& name_and_isa_blobs : get_code_object_blobs()) {
                 for (auto&& isa_blobs : name_and_isa_blobs.second) {
                     for (auto&& blob : isa_blobs.second) {
-                        std::stringstream tmp{std::string{blob.cbegin(), blob.cend()}};
-
-                        ELFIO::elfio reader;
-
-                        if (!reader.load(tmp)) continue;
-
-                        read_kernarg_metadata(reader, kernargs.second);
+                        read_kernarg_metadata(blob, kernargs.second);
                     }
                 }
             }
@@ -711,12 +814,20 @@ public:
 
         auto it1 = get_kernargs().find(it->second);
         if (it1 == get_kernargs().end()) {
-            hip_throw(std::runtime_error{
-                      "Missing metadata for __global__ function: " + it->second});
+            it1 = get_kernargs().find(it->second + ".kd");
+            if (it1 == get_kernargs().end()) {
+                hip_throw(std::runtime_error{
+                          "Missing metadata for __global__ function: " + it->second});
+            }
         }
 
         return it1->second;
     }
 };  // class program_state_impl
 
+struct kernarg_impl {
+    std::vector<std::uint8_t> v;
+};
+
+
 };

tests/src/hiprtc/hiprtcGetLoweredName.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
- * TEST: %t
+ * TEST: %t EXCLUDE_HIP_PLATFORM hcc
  * HIT_END
  */
 #include <test_common.h>

tests/src/hiprtc/hiprtcGetTypeName.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
- * TEST: %t
+ * TEST: %t EXCLUDE_HIP_PLATFORM hcc
  * HIT_END
  */
 

tests/src/hiprtc/saxpy.cpp

@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp LINK_OPTIONS hiprtc EXCLUDE_HIP_PLATFORM nvcc
- * TEST: %t
+ * TEST: %t EXCLUDE_HIP_PLATFORM hcc
  * HIT_END
  */