Choose whether or not to use functional grid_launch based on the version of HCC used to compile.

src/functional_grid_launch.inl

@@ -0,0 +1,138 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hcc_detail/grid_launch_GGL.hpp"
+#include "hip/hcc_detail/program_state.hpp"
+
+#include "hip/hip_runtime_api.h"
+
+// Internal header, do not percolate upwards.
+#include "hip_hcc_internal.h"
+#include "hc.hpp"
+#include "trace_helper.h"
+
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <stdexcept>
+
+#include <iostream>
+
+using namespace hc;
+using namespace std;
+
+namespace hip_impl
+{
+    namespace
+    {
+        inline
+        string name(uintptr_t function_address)
+        {
+            const auto it = function_names().find(function_address);
+
+            if (it == function_names().cend())  {
+                throw runtime_error{
+                    "Invalid function passed to hipLaunchKernelGGL."};
+            }
+
+            return it->second;
+        }
+
+        inline
+        string name(hsa_agent_t agent)
+        {
+            char n[64] = {};
+            hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, n);
+
+            return string{n};
+        }
+
+        inline
+        hsa_agent_t target_agent(hipStream_t stream)
+        {
+            if (stream) {
+                return *static_cast<hsa_agent_t*>(
+                    stream->locked_getAv()->get_hsa_agent());
+            }
+            else if (
+                ihipGetTlsDefaultCtx() && ihipGetTlsDefaultCtx()->getDevice()) {
+                return ihipGetDevice(
+                    ihipGetTlsDefaultCtx()->getDevice()->_deviceId)->_hsaAgent;
+            }
+            else {
+                return *static_cast<hsa_agent_t*>(
+                    accelerator{}.get_default_view().get_hsa_agent());
+            }
+        }
+    }
+
+    void hipLaunchKernelGGLImpl(
+        uintptr_t function_address,
+        const dim3& numBlocks,
+        const dim3& dimBlocks,
+        uint32_t sharedMemBytes,
+        hipStream_t stream,
+        void** kernarg)
+    {
+        const auto it0 = functions().find(function_address);
+
+        if (it0 == functions().cend()) {
+            throw runtime_error{
+                "No device code available for function: " +
+                name(function_address)
+            };
+        }
+
+        auto agent = target_agent(stream);
+
+        const auto it1 = find_if(
+            it0->second.cbegin(),
+            it0->second.cend(),
+            [=](const pair<hsa_agent_t, Kernel_descriptor>& x) {
+                return x.first.handle == agent.handle;
+            });
+
+        if (it1 == it0->second.cend()) {
+            throw runtime_error{
+                "No code available for function: " + name(function_address) +
+                ", for agent: " + name(agent)
+            };
+        }
+
+        for (auto&& agent_kernel : it0->second) {
+            if (agent.handle == agent_kernel.first.handle) {
+                hipModuleLaunchKernel(
+                    agent_kernel.second,
+                    numBlocks.x,
+                    numBlocks.y,
+                    numBlocks.z,
+                    dimBlocks.x,
+                    dimBlocks.y,
+                    dimBlocks.z,
+                    sharedMemBytes,
+                    stream,
+                    nullptr,
+                    kernarg);
+            }
+        }
+    }
+}