Adding changes for hipExtLaunchKernel for rocCLR

Change-Id: Iba52bc3bde7c37f3fb375a55ba0947e87b3cdc9b [ROCm/hip commit: 2d517fdcc6]
2020-05-08 18:18:36 +00:00
commit 7b52f0a1ea
@@ -392,10 +392,53 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri

 typedef int hipLaunchParm;

+template <std::size_t n, typename... Ts,
+          typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
+void pArgs(const std::tuple<Ts...>&, void*) {}
+
+template <std::size_t n, typename... Ts,
+          typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
+void pArgs(const std::tuple<Ts...>& formals, void** _vargs) {
+    using T = typename std::tuple_element<n, std::tuple<Ts...> >::type;
+
+    static_assert(!std::is_reference<T>{},
+                  "A __global__ function cannot have a reference as one of its "
+                  "arguments.");
+#if defined(HIP_STRICT)
+    static_assert(std::is_trivially_copyable<T>{},
+                  "Only TriviallyCopyable types can be arguments to a __global__ "
+                  "function");
+#endif
+    _vargs[n] = const_cast<void*>(reinterpret_cast<const void*>(&std::get<n>(formals)));
+    return pArgs<n + 1>(formals, _vargs);
+}
+
+template <typename... Formals, typename... Actuals>
+std::tuple<Formals...> validateArgsCountType(void (*kernel)(Formals...), std::tuple<Actuals...>(actuals)) {
+    static_assert(sizeof...(Formals) == sizeof...(Actuals), "Argument Count Mismatch");
+    std::tuple<Formals...> to_formals{std::move(actuals)};
+    return to_formals;
+}
+
+#if defined(HIP_TEMPLATE_KERNEL_LAUNCH)
+template <typename... Args, typename F = void (*)(Args...)>
+void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
+                        std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) {
+    constexpr size_t count = sizeof...(Args);
+    auto tup_ = std::tuple<Args...>{args...};
+    auto tup = validateArgsCountType(kernel, tup_);
+    void* _Args[count];
+    pArgs<0>(tup, _Args);
+
+    auto k = reinterpret_cast<void*>(kernel);
+    hipLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream);
+}
+#else
 #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...)          \
    do {                                                                                           \
        kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__);         \
    } while (0)
+#endif

 #include <hip/hip_runtime_api.h>

@@ -3378,6 +3378,11 @@ hipError_t hipLaunchKernel(const void* function_address,
                           hipStream_t stream __dparm(0));

 #if __HIP_ROCclr__ || !defined(__HCC__)
+//TODO: Move this to hip_ext.h
+hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks,
+                              void** args, size_t sharedMemBytes, hipStream_t stream,
+                              hipEvent_t startEvent, hipEvent_t stopEvent, int flags);
+
 hipError_t hipBindTexture(
    size_t* offset,
    const textureReference* tex,
@@ -23,6 +23,10 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HIP_EXT_H
 #define HIP_INCLUDE_HIP_HIP_EXT_H
 #include "hip/hip_runtime.h"
+#if defined(__cplusplus)
+#include <tuple>
+#include <type_traits>
+#endif
 #ifdef __HCC__

 // Forward declarations:
@@ -109,8 +113,29 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                    hipEvent_t stopEvent = nullptr)
                                    __attribute__((deprecated("use hipExtModuleLaunchKernel instead")));

-//#if !__HIP_ROCclr__ && defined(__cplusplus)
-#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
+#if defined(__HIP_ROCclr__) && defined(__cplusplus)
+
+extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks,
+                                         dim3 dimBlocks, void** args, size_t sharedMemBytes,
+                                         hipStream_t stream, hipEvent_t startEvent,
+                                         hipEvent_t stopEvent, int flags);
+
+template <typename... Args, typename F = void (*)(Args...)>
+inline void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
+                                  std::uint32_t sharedMemBytes, hipStream_t stream,
+                                  hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags,
+                                  Args... args) {
+    constexpr size_t count = sizeof...(Args);
+    auto tup_ = std::tuple<Args...>{args...};
+    auto tup = validateArgsCountType(kernel, tup_);
+    void* _Args[count];
+    pArgs<0>(tup, _Args);
+
+    auto k = reinterpret_cast<void*>(kernel);
+    hipExtLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream, startEvent,
+                       stopEvent, (int)flags);
+}
+#elif defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
 //kernel_descriptor and hip_impl::make_kernarg are in "grid_launch_GGL.hpp"

 namespace hip_impl {