resolved issues with standard headers in device files: device_base.hpp,...

resolved issues with standard headers in device files: device_base.hpp, device_grouped_conv_fwd_multiple_abd.hpp, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

include/ck/tensor_operation/gpu/device/device_base.hpp

@@ -3,15 +3,17 @@
 
 #pragma once
 
+#ifndef CK_CODE_GEN_RTC
 #include <string>
 #include <sstream>
-
 #include "ck/stream_config.hpp"
+#endif
 
 namespace ck {
 namespace tensor_operation {
 namespace device {
 
+#ifndef CK_CODE_GEN_RTC
 struct BaseArgument
 {
     BaseArgument()                    = default;
@@ -36,13 +38,14 @@ struct BaseInvoker
 
     virtual ~BaseInvoker() {}
 };
+#endif
 
 struct BaseOperator
 {
     BaseOperator()                    = default;
     BaseOperator(const BaseOperator&) = default;
     BaseOperator& operator=(const BaseOperator&) = default;
-
+#ifndef CK_CODE_GEN_RTC
     virtual bool IsSupportedArgument(const BaseArgument*) { return false; }
     virtual std::string GetTypeString() const { return ""; }
 
@@ -66,7 +69,7 @@ struct BaseOperator
         assert(p_arg);
         p_arg->p_workspace_ = p_workspace;
     }
-
+#endif
     virtual ~BaseOperator() {}
 };
 

include/ck/tensor_operation/gpu/device/device_grouped_conv_fwd_multiple_abd.hpp

@@ -3,7 +3,9 @@
 
 #pragma once
 
+#ifndef CK_CODE_GEN_RTC
 #include <array>
+#endif
 
 #include "ck/tensor_operation/gpu/device/device_base.hpp"
 #include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp"
@@ -13,8 +15,13 @@ namespace ck {
 namespace tensor_operation {
 namespace device {
 
+#ifdef CK_CODE_GEN_RTC
+template <typename T>
+using is_tuple = decltype(ck::declval<T&>().IsTuple());
+#else
 template <typename T>
 using is_tuple = decltype(std::declval<T&>().IsTuple());
+#endif
 
 /**
  * \brief Grouped Convolution Forward
@@ -72,12 +79,18 @@ struct DeviceGroupedConvFwdMultipleABD : public BaseOperator
     static constexpr index_t NumDTensor = DsDataType::Size();
 
     static_assert(NumDTensor == DsLayout::Size(), "wrong! Inconsistent NumDTensor");
-
+#ifdef CK_CODE_GEN_RTC
+    using APointers = ck::conditional_t<isMultiA, ck::Array<const void*, NumATensor>&, const void*>;
+    using BPointers = ck::conditional_t<isMultiB, ck::Array<const void*, NumBTensor>&, const void*>;
+#else
     // If DataType is tuple, user has to pass std::array with pointers.
     using APointers =
-        std::conditional_t<isMultiA, std::array<const void*, NumATensor>&, const void*>;
+        ck::conditional_t<isMultiA, std::array<const void*, NumATensor>&, const void*>;
     using BPointers =
-        std::conditional_t<isMultiB, std::array<const void*, NumBTensor>&, const void*>;
+        ck::conditional_t<isMultiB, std::array<const void*, NumBTensor>&, const void*>;
+#endif
+
+#ifndef CK_CODE_GEN_RTC
 
     /**
      * \brief Make argument pointer for grouped conv fwd.
@@ -151,6 +164,7 @@ struct DeviceGroupedConvFwdMultipleABD : public BaseOperator
 
     virtual std::unique_ptr<BaseInvoker> MakeInvokerPointer() = 0;
 };
+#endif
 
 } // namespace device
 } // namespace tensor_operation

include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

@@ -212,9 +212,13 @@ __global__ void
 }
 
 } // namespace
-
+#ifdef CK_CODE_GEN_RTC
+template <typename T>
+using is_tuple = decltype(ck::declval<T&>().IsTuple());
+#else
 template <typename T>
 using is_tuple = decltype(std::declval<T&>().IsTuple());
+#endif
 
 //
 // @brief      Device Convolution operation.

include/ck/utility/type.hpp

@@ -116,9 +116,13 @@ constexpr T&& forward(typename remove_reference<T>::type&& t_) noexcept
 {
     return static_cast<T&&>(t_);
 }
+
+template <typename T>
+T&& declval() noexcept;
 #else
 #include <utility>
 #include <type_traits>
+using std::declval;
 using std::forward;
 using std::is_base_of;
 using std::is_class;