update GGL to log launched kernel information
Change-Id: Ied0aa6055673c687071b4a579aecd17f0f3f09ce
[ROCm/hip commit: 6d4af1ab1f]
This commit is contained in:
@@ -21,6 +21,7 @@ THE SOFTWARE.
|
|||||||
*/
|
*/
|
||||||
|
|
||||||
#pragma once
|
#pragma once
|
||||||
|
|
||||||
#if GENERIC_GRID_LAUNCH == 1
|
#if GENERIC_GRID_LAUNCH == 1
|
||||||
|
|
||||||
#include "concepts.hpp"
|
#include "concepts.hpp"
|
||||||
@@ -71,7 +72,7 @@ namespace hip_impl
|
|||||||
|
|
||||||
template<FunctionalProcedure F, typename... Ts>
|
template<FunctionalProcedure F, typename... Ts>
|
||||||
using is_new_grid_launch_t = typename std::conditional<
|
using is_new_grid_launch_t = typename std::conditional<
|
||||||
std::is_callable<F(Ts...)>{},
|
is_callable<F(Ts...)>{},
|
||||||
New_grid_launch_tag,
|
New_grid_launch_tag,
|
||||||
Old_grid_launch_tag>::type;
|
Old_grid_launch_tag>::type;
|
||||||
}
|
}
|
||||||
@@ -118,6 +119,7 @@ namespace hip_impl
|
|||||||
// TODO: these are workarounds, they should be removed.
|
// TODO: these are workarounds, they should be removed.
|
||||||
|
|
||||||
hc::accelerator_view lock_stream_hip_(hipStream_t&, void*&);
|
hc::accelerator_view lock_stream_hip_(hipStream_t&, void*&);
|
||||||
|
void print_prelaunch_trace_(const char*, dim3, dim3, int, hipStream_t);
|
||||||
void unlock_stream_hip_(
|
void unlock_stream_hip_(
|
||||||
hipStream_t, void*, const char*, hc::accelerator_view*);
|
hipStream_t, void*, const char*, hc::accelerator_view*);
|
||||||
|
|
||||||
@@ -137,7 +139,13 @@ namespace hip_impl
|
|||||||
void* lck_stream = nullptr;
|
void* lck_stream = nullptr;
|
||||||
auto acc_v = lock_stream_hip_(stream, lck_stream);
|
auto acc_v = lock_stream_hip_(stream, lck_stream);
|
||||||
auto stream_guard = make_RAII_guard(
|
auto stream_guard = make_RAII_guard(
|
||||||
[](){ /* perhaps use a slimmed down ihipPrintKernelLaunch here */ },
|
std::bind(
|
||||||
|
print_prelaunch_trace_,
|
||||||
|
kernel_name,
|
||||||
|
num_blocks,
|
||||||
|
dim_blocks,
|
||||||
|
group_mem_bytes,
|
||||||
|
stream),
|
||||||
std::bind(
|
std::bind(
|
||||||
unlock_stream_hip_, stream, lck_stream, kernel_name, &acc_v));
|
unlock_stream_hip_, stream, lck_stream, kernel_name, &acc_v));
|
||||||
|
|
||||||
@@ -841,16 +849,15 @@ namespace hip_impl
|
|||||||
group_mem_bytes,\
|
group_mem_bytes,\
|
||||||
stream,\
|
stream,\
|
||||||
...)\
|
...)\
|
||||||
do {\
|
do {\
|
||||||
hipLaunchKernelGGL(\
|
hipLaunchKernelGGL(\
|
||||||
kernel_name,\
|
kernel_name,\
|
||||||
num_blocks,\
|
num_blocks,\
|
||||||
dim_blocks,\
|
dim_blocks,\
|
||||||
group_mem_bytes,\
|
group_mem_bytes,\
|
||||||
stream,\
|
stream,\
|
||||||
hipLaunchParm{},\
|
hipLaunchParm{},\
|
||||||
##__VA_ARGS__);\
|
##__VA_ARGS__);\
|
||||||
} while(0)
|
} while(0)
|
||||||
|
|
||||||
}
|
}
|
||||||
#endif //GENERIC_GRID_LAUNCH
|
#endif //GENERIC_GRID_LAUNCH
|
||||||
|
|||||||
@@ -21,6 +21,7 @@ THE SOFTWARE.
|
|||||||
*/
|
*/
|
||||||
|
|
||||||
#pragma once
|
#pragma once
|
||||||
|
#include "concepts.hpp"
|
||||||
|
|
||||||
#include <type_traits> // For std::conditional, std::decay, std::enable_if,
|
#include <type_traits> // For std::conditional, std::decay, std::enable_if,
|
||||||
// std::false_type, std result_of and std::true_type.
|
// std::false_type, std result_of and std::true_type.
|
||||||
@@ -29,9 +30,6 @@ THE SOFTWARE.
|
|||||||
namespace std
|
namespace std
|
||||||
{ // TODO: these should be removed as soon as possible.
|
{ // TODO: these should be removed as soon as possible.
|
||||||
#if (__cplusplus < 201406L)
|
#if (__cplusplus < 201406L)
|
||||||
template<typename...>
|
|
||||||
using void_t = void;
|
|
||||||
|
|
||||||
#if (__cplusplus < 201402L)
|
#if (__cplusplus < 201402L)
|
||||||
template<bool cond, typename T = void>
|
template<bool cond, typename T = void>
|
||||||
using enable_if_t = typename enable_if<cond, T>::type;
|
using enable_if_t = typename enable_if<cond, T>::type;
|
||||||
@@ -43,88 +41,80 @@ namespace std
|
|||||||
using result_of_t = typename result_of<F(Ts...)>::type;
|
using result_of_t = typename result_of<F(Ts...)>::type;
|
||||||
template<typename T>
|
template<typename T>
|
||||||
using remove_reference_t = typename remove_reference<T>::type;
|
using remove_reference_t = typename remove_reference<T>::type;
|
||||||
template<
|
|
||||||
FunctionalProcedure F,
|
|
||||||
unsigned int n = 0u,
|
|
||||||
typename = void>
|
|
||||||
struct is_callable_impl : is_callable_impl<F, n + 1u> {};
|
|
||||||
|
|
||||||
// Pointer to member function, call through non-pointer.
|
|
||||||
template<FunctionalProcedure F, typename C, typename... Ts>
|
|
||||||
struct is_callable_impl<
|
|
||||||
F(C, Ts...),
|
|
||||||
0u,
|
|
||||||
void_t<decltype((declval<C>().*declval<F>())(declval<Ts>()...))>
|
|
||||||
> : true_type {
|
|
||||||
};
|
|
||||||
|
|
||||||
// Pointer to member function, call through pointer.
|
|
||||||
template<FunctionalProcedure F, typename C, typename... Ts>
|
|
||||||
struct is_callable_impl<
|
|
||||||
F(C, Ts...),
|
|
||||||
1u,
|
|
||||||
void_t<decltype(((*declval<C>()).*declval<F>())(declval<Ts>()...))>
|
|
||||||
> : std::true_type {
|
|
||||||
};
|
|
||||||
|
|
||||||
// Pointer to member data, call through non-pointer, no args.
|
|
||||||
template<FunctionalProcedure F, typename C>
|
|
||||||
struct is_callable_impl<
|
|
||||||
F(C),
|
|
||||||
2u,
|
|
||||||
void_t<decltype(declval<C>().*declval<F>())>
|
|
||||||
> : true_type {
|
|
||||||
};
|
|
||||||
|
|
||||||
// Pointer to member data, call through pointer, no args.
|
|
||||||
template<FunctionalProcedure F, typename C>
|
|
||||||
struct is_callable_impl<
|
|
||||||
F(C),
|
|
||||||
3u,
|
|
||||||
void_t<decltype(*declval<C>().*declval<F>())>
|
|
||||||
> : true_type {
|
|
||||||
};
|
|
||||||
|
|
||||||
// General call, n args.
|
|
||||||
template<FunctionalProcedure F, typename... Ts>
|
|
||||||
struct is_callable_impl<
|
|
||||||
F(Ts...),
|
|
||||||
4u,
|
|
||||||
void_t<decltype(declval<F>()(declval<Ts>()...))>
|
|
||||||
> : true_type {
|
|
||||||
};
|
|
||||||
|
|
||||||
// Not callable.
|
|
||||||
template<FunctionalProcedure F>
|
|
||||||
struct is_callable_impl<F, 5u> : false_type {};
|
|
||||||
|
|
||||||
template<typename Call>
|
|
||||||
struct is_callable : is_callable_impl<Call> {};
|
|
||||||
#else
|
|
||||||
template<typename, typename = void>
|
|
||||||
struct is_callable_impl : false_type {};
|
|
||||||
|
|
||||||
template<FunctionalProcedure F, typename... Ts>
|
|
||||||
struct is_callable_impl<
|
|
||||||
F(Ts...),
|
|
||||||
void_t<result_of_t<F(Ts...)>>> : true_type {};
|
|
||||||
|
|
||||||
template<typename F>
|
|
||||||
struct is_callable : is_callable_impl<F> {};
|
|
||||||
#endif
|
#endif
|
||||||
template<typename...>
|
|
||||||
struct disjunction : false_type {};
|
|
||||||
template<typename B1>
|
|
||||||
struct disjunction<B1> : B1 {};
|
|
||||||
template<typename B1, typename... Bs>
|
|
||||||
struct disjunction<B1, Bs...>
|
|
||||||
: conditional_t<B1{} == true, B1, disjunction<Bs...>>
|
|
||||||
{};
|
|
||||||
#endif
|
#endif
|
||||||
}
|
}
|
||||||
|
|
||||||
namespace hip_impl // Only for documentation, macros ignore namespaces.
|
namespace hip_impl
|
||||||
{
|
{
|
||||||
|
template<typename...>
|
||||||
|
using void_t_ = void;
|
||||||
|
|
||||||
|
#if (__cplusplus < 201402L)
|
||||||
|
template<
|
||||||
|
FunctionalProcedure F,
|
||||||
|
unsigned int n = 0u,
|
||||||
|
typename = void>
|
||||||
|
struct is_callable_impl : is_callable_impl<F, n + 1u> {};
|
||||||
|
|
||||||
|
// Pointer to member function, call through non-pointer.
|
||||||
|
template<FunctionalProcedure F, typename C, typename... Ts>
|
||||||
|
struct is_callable_impl<
|
||||||
|
F(C, Ts...),
|
||||||
|
0u,
|
||||||
|
void_t_<decltype((std::declval<C>().*std::declval<F>())(
|
||||||
|
std::declval<Ts>()...))>
|
||||||
|
> : std::true_type {};
|
||||||
|
|
||||||
|
// Pointer to member function, call through pointer.
|
||||||
|
template<FunctionalProcedure F, typename C, typename... Ts>
|
||||||
|
struct is_callable_impl<
|
||||||
|
F(C, Ts...),
|
||||||
|
1u,
|
||||||
|
void_t_<decltype(((*std::declval<C>()).*std::declval<F>())(
|
||||||
|
std::declval<Ts>()...))>
|
||||||
|
> : std::true_type {};
|
||||||
|
|
||||||
|
// Pointer to member data, call through non-pointer, no args.
|
||||||
|
template<FunctionalProcedure F, typename C>
|
||||||
|
struct is_callable_impl<
|
||||||
|
F(C),
|
||||||
|
2u,
|
||||||
|
void_t_<decltype(std::declval<C>().*std::declval<F>())>
|
||||||
|
> : std::true_type {};
|
||||||
|
|
||||||
|
// Pointer to member data, call through pointer, no args.
|
||||||
|
template<FunctionalProcedure F, typename C>
|
||||||
|
struct is_callable_impl<
|
||||||
|
F(C),
|
||||||
|
3u,
|
||||||
|
void_t_<decltype(*std::declval<C>().*std::declval<F>())>
|
||||||
|
> : std::true_type {};
|
||||||
|
|
||||||
|
// General call, n args.
|
||||||
|
template<FunctionalProcedure F, typename... Ts>
|
||||||
|
struct is_callable_impl<
|
||||||
|
F(Ts...),
|
||||||
|
4u,
|
||||||
|
void_t_<decltype(std::declval<F>()(std::declval<Ts>()...))>
|
||||||
|
> : std::true_type {};
|
||||||
|
|
||||||
|
// Not callable.
|
||||||
|
template<FunctionalProcedure F>
|
||||||
|
struct is_callable_impl<F, 5u> : std::false_type {};
|
||||||
|
|
||||||
|
template<typename Call>
|
||||||
|
struct is_callable : is_callable_impl<Call> {};
|
||||||
|
#else
|
||||||
|
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 {};
|
||||||
|
#endif
|
||||||
|
|
||||||
#define count_macro_args_impl_hip_(\
|
#define count_macro_args_impl_hip_(\
|
||||||
_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15,\
|
_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15,\
|
||||||
_16, _17, _18, _19, _20, _21, _22, _23, _24, _25, _26, _27, _28, _29,\
|
_16, _17, _18, _19, _20, _21, _22, _23, _24, _25, _26, _27, _28, _29,\
|
||||||
|
|||||||
@@ -27,6 +27,9 @@ THE SOFTWARE.
|
|||||||
#include "hc.hpp"
|
#include "hc.hpp"
|
||||||
#include "trace_helper.h"
|
#include "trace_helper.h"
|
||||||
|
|
||||||
|
#include <iostream>
|
||||||
|
#include <sstream>
|
||||||
|
|
||||||
namespace hip_impl
|
namespace hip_impl
|
||||||
{
|
{
|
||||||
hc::accelerator_view lock_stream_hip_(
|
hc::accelerator_view lock_stream_hip_(
|
||||||
@@ -42,6 +45,39 @@ namespace hip_impl
|
|||||||
return (*static_cast<L*>(locked_stream))->_av;
|
return (*static_cast<L*>(locked_stream))->_av;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
void print_prelaunch_trace_(
|
||||||
|
const char* kernel_name,
|
||||||
|
dim3 num_blocks,
|
||||||
|
dim3 dim_blocks,
|
||||||
|
int group_mem_bytes,
|
||||||
|
hipStream_t stream)
|
||||||
|
{
|
||||||
|
if ((HIP_TRACE_API & (1 << TRACE_CMD)) ||
|
||||||
|
HIP_PROFILE_API ||
|
||||||
|
(COMPILE_HIP_DB && HIP_TRACE_API)) {
|
||||||
|
std::stringstream os;
|
||||||
|
os << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum()
|
||||||
|
<< " hipLaunchKernel '" << kernel_name << "'"
|
||||||
|
<< " gridDim:" << num_blocks
|
||||||
|
<< " groupDim:" << dim_blocks
|
||||||
|
<< " sharedMem:+" << group_mem_bytes
|
||||||
|
<< " " << *stream;
|
||||||
|
|
||||||
|
if (HIP_PROFILE_API == 0x1) {
|
||||||
|
std::string shortAtpString("hipLaunchKernel:");
|
||||||
|
shortAtpString += kernel_name;
|
||||||
|
MARKER_BEGIN(shortAtpString.c_str(), "HIP");
|
||||||
|
} else if (HIP_PROFILE_API == 0x2) {
|
||||||
|
MARKER_BEGIN(os.str().c_str(), "HIP");
|
||||||
|
}
|
||||||
|
|
||||||
|
if (COMPILE_HIP_DB && HIP_TRACE_API) {
|
||||||
|
std::cerr << API_COLOR << os.str() << API_COLOR_END
|
||||||
|
<< std::endl;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
void unlock_stream_hip_(
|
void unlock_stream_hip_(
|
||||||
hipStream_t stream,
|
hipStream_t stream,
|
||||||
void* locked_stream,
|
void* locked_stream,
|
||||||
|
|||||||
Verwijs in nieuw issue
Block a user