EXSWCPHIPT-42 - Add HIP RTC support to the test framework (#2719)

* EXSWCPHIPT-42 - Add HIP RTC support to the test framework
* Removed ifdef from hipTestContext class

tests/catch/include/hip_test_common.hh

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2021 - 2022 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
@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 #pragma once
 #include "hip_test_context.hh"
+#include <hip_test_rtc.hh>
 #include <catch.hpp>
 #include <stdlib.h>
 
@@ -171,8 +172,60 @@ static inline void HIP_SKIP_TEST(char const* const reason) noexcept {
   // ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test.
   std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl;
 }
+
+/**
+ * @brief Helper template that returns the expected arguments of a kernel.
+ *
+ * @return constexpr std::tuple<FArgs...> the expected arguments of the kernel.
+ */
+template <typename... FArgs> std::tuple<FArgs...> getExpectedArgs(void(FArgs...)){};
+
+/**
+ * @brief Asserts that the types of the arguments of a function match exactly with the types in the
+ * function signature.
+ * This is necessary because HIP RTC does not do implicit casting of the kernel
+ * parameters.
+ * In order to get the kernel function signature, this function should only called when
+ * RTC is disabled.
+ *
+ * @tparam F the kernel function
+ * @tparam Args the parameters that will be passed to the kernel.
+ */
+template <typename F, typename... Args> void validateArguments(F f, Args...) {
+  using expectedArgsTuple = decltype(getExpectedArgs(f));
+  static_assert(std::is_same<expectedArgsTuple, std::tuple<Args...>>::value,
+                "Kernel arguments types must match exactly!");
 }
 
+/**
+ * @brief Launch a kernel using either HIP or HIP RTC.
+ *
+ * @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a
+ * template).
+ * @tparam K The kernel type. Expects a function or template when RTC is disabled. Expects a
+ * function pointer instead when RTC is enabled.
+ * @tparam Dim Can be either dim3 or int.
+ * @tparam Args A list of kernel arguments to be forwarded.
+ * @param kernel The kernel to be launched (defined in kernels.hh)
+ * @param numBlocks
+ * @param numThreads
+ * @param memPerBlock
+ * @param stream
+ * @param packedArgs A list of kernel arguments to be forwarded.
+ */
+template <typename... Typenames, typename K, typename Dim, typename... Args>
+void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock,
+                  hipStream_t stream, Args&&... packedArgs) {
+#ifndef RTC_TESTING
+    validateArguments(kernel, packedArgs...);
+    kernel<<<numBlocks, numThreads, memPerBlock, stream>>>(std::forward<Args>(packedArgs)...);
+#else
+  launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
+                                std::forward<Args>(packedArgs)...);
+#endif
+}
+}  // namespace HipTest
+
 
 // This must be called in the beginning of image test app's main() to indicate whether image
 // is supported.