Choose whether or not to use functional grid_launch based on the version of HCC used to compile.
[ROCm/clr commit: c6ab6f292b]
Этот коммит содержится в:
Alex Voicu
@@ -1,3 +1,25 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hsa/hsa.h>
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -0,0 +1,159 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "code_object_bundle.hpp"
|
||||
#include "concepts.hpp"
|
||||
#include "helpers.hpp"
|
||||
#include "program_state.hpp"
|
||||
|
||||
#include "hc.hpp"
|
||||
#include "hip/hip_hcc.h"
|
||||
#include "hip_runtime.h"
|
||||
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <functional>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
#include <stdexcept>
|
||||
#include <string>
|
||||
#include <tuple>
|
||||
#include <type_traits>
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<std::is_integral<T>{}>::type* = nullptr>
|
||||
inline
|
||||
T round_up_to_next_multiple_nonnegative(T x, T y)
|
||||
{
|
||||
T tmp = x + y - 1;
|
||||
return tmp - tmp % y;
|
||||
}
|
||||
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg()
|
||||
{
|
||||
return {};
|
||||
}
|
||||
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(std::vector<std::uint8_t> kernarg)
|
||||
{
|
||||
return kernarg;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(std::vector<uint8_t> kernarg, T x)
|
||||
{
|
||||
kernarg.resize(
|
||||
round_up_to_next_multiple_nonnegative(kernarg.size(), alignof(T)) +
|
||||
sizeof(T));
|
||||
|
||||
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::move(x)};
|
||||
|
||||
return kernarg;
|
||||
}
|
||||
|
||||
template<typename T, typename... Ts>
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(
|
||||
std::vector<std::uint8_t> kernarg, T x, Ts... xs)
|
||||
{
|
||||
return make_kernarg(
|
||||
make_kernarg(std::move(kernarg), std::move(x)), std::move(xs)...);
|
||||
}
|
||||
|
||||
template<typename... Ts>
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(Ts... xs)
|
||||
{
|
||||
std::vector<std::uint8_t> kernarg;
|
||||
kernarg.reserve(sizeof(std::tuple<Ts...>));
|
||||
|
||||
return make_kernarg(std::move(kernarg), std::move(xs)...);
|
||||
}
|
||||
|
||||
void hipLaunchKernelGGLImpl(
|
||||
std::uintptr_t function_address,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t sharedMemBytes,
|
||||
hipStream_t stream,
|
||||
void** kernarg);
|
||||
} // Namespace hip_impl.
|
||||
|
||||
template<typename... Args, typename F = void (*)(Args...)>
|
||||
inline
|
||||
void hipLaunchKernelGGL(
|
||||
F kernel,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t sharedMemBytes,
|
||||
hipStream_t stream,
|
||||
Args... args)
|
||||
{
|
||||
auto kernarg = hip_impl::make_kernarg(std::move(args)...);
|
||||
std::size_t kernarg_size = kernarg.size();
|
||||
|
||||
void* config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, kernarg.data(),
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &kernarg_size,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
hip_impl::hipLaunchKernelGGLImpl(
|
||||
reinterpret_cast<std::uintptr_t>(kernel),
|
||||
numBlocks,
|
||||
dimBlocks,
|
||||
sharedMemBytes,
|
||||
stream,
|
||||
&config[0]);
|
||||
}
|
||||
|
||||
template<typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
|
||||
inline
|
||||
void hipLaunchKernel(
|
||||
F kernel,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t groupMemBytes,
|
||||
hipStream_t stream,
|
||||
Args... args)
|
||||
{
|
||||
hipLaunchKernelGGL(
|
||||
kernel,
|
||||
numBlocks,
|
||||
dimBlocks,
|
||||
groupMemBytes,
|
||||
stream,
|
||||
hipLaunchParm{},
|
||||
std::move(args)...);
|
||||
}
|
||||
|
||||
@@ -20,143 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#if GENERIC_GRID_LAUNCH == 1
|
||||
|
||||
#include "code_object_bundle.hpp"
|
||||
#include "concepts.hpp"
|
||||
#include "helpers.hpp"
|
||||
#include "program_state.hpp"
|
||||
|
||||
#include "hc.hpp"
|
||||
#include "hip/hip_hcc.h"
|
||||
#include "hip_runtime.h"
|
||||
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <functional>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
#include <stdexcept>
|
||||
#include <string>
|
||||
#include <tuple>
|
||||
#include <type_traits>
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<std::is_integral<T>{}>::type* = nullptr>
|
||||
inline
|
||||
T round_up_to_next_multiple_nonnegative(T x, T y)
|
||||
{
|
||||
T tmp = x + y - 1;
|
||||
return tmp - tmp % y;
|
||||
}
|
||||
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg()
|
||||
{
|
||||
return {};
|
||||
}
|
||||
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(std::vector<std::uint8_t> kernarg)
|
||||
{
|
||||
return kernarg;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(std::vector<uint8_t> kernarg, T x)
|
||||
{
|
||||
kernarg.resize(
|
||||
round_up_to_next_multiple_nonnegative(kernarg.size(), alignof(T)) +
|
||||
sizeof(T));
|
||||
|
||||
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::move(x)};
|
||||
|
||||
return kernarg;
|
||||
}
|
||||
|
||||
template<typename T, typename... Ts>
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(
|
||||
std::vector<std::uint8_t> kernarg, T x, Ts... xs)
|
||||
{
|
||||
return make_kernarg(
|
||||
make_kernarg(std::move(kernarg), std::move(x)), std::move(xs)...);
|
||||
}
|
||||
|
||||
template<typename... Ts>
|
||||
inline
|
||||
std::vector<std::uint8_t> make_kernarg(Ts... xs)
|
||||
{
|
||||
std::vector<std::uint8_t> kernarg;
|
||||
kernarg.reserve(sizeof(std::tuple<Ts...>));
|
||||
|
||||
return make_kernarg(std::move(kernarg), std::move(xs)...);
|
||||
}
|
||||
|
||||
void hipLaunchKernelGGLImpl(
|
||||
std::uintptr_t function_address,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t sharedMemBytes,
|
||||
hipStream_t stream,
|
||||
void** kernarg);
|
||||
} // Namespace hip_impl.
|
||||
|
||||
template<typename... Args, typename F = void (*)(Args...)>
|
||||
inline
|
||||
void hipLaunchKernelGGL(
|
||||
F kernel,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t sharedMemBytes,
|
||||
hipStream_t stream,
|
||||
Args... args)
|
||||
{
|
||||
auto kernarg = hip_impl::make_kernarg(std::move(args)...);
|
||||
std::size_t kernarg_size = kernarg.size();
|
||||
|
||||
void* config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, kernarg.data(),
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &kernarg_size,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
hip_impl::hipLaunchKernelGGLImpl(
|
||||
reinterpret_cast<std::uintptr_t>(kernel),
|
||||
numBlocks,
|
||||
dimBlocks,
|
||||
sharedMemBytes,
|
||||
stream,
|
||||
&config[0]);
|
||||
}
|
||||
|
||||
template<typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
|
||||
inline
|
||||
void hipLaunchKernel(
|
||||
F kernel,
|
||||
const dim3& numBlocks,
|
||||
const dim3& dimBlocks,
|
||||
std::uint32_t groupMemBytes,
|
||||
hipStream_t stream,
|
||||
Args... args)
|
||||
{
|
||||
hipLaunchKernelGGL(
|
||||
kernel,
|
||||
numBlocks,
|
||||
dimBlocks,
|
||||
groupMemBytes,
|
||||
stream,
|
||||
hipLaunchParm{},
|
||||
std::move(args)...);
|
||||
}
|
||||
|
||||
#endif //GENERIC_GRID_LAUNCH
|
||||
#if __hcc_workweek__ >= 17481
|
||||
#define FUNCTIONAL_GRID_LAUNCH
|
||||
#include "functional_grid_launch.hpp"
|
||||
#else
|
||||
#include "macro_based_grid_launch.hpp"
|
||||
#endif
|
||||
#endif //GENERIC_GRID_LAUNCH
|
||||
Разница между файлами не показана из-за своего большого размера
Загрузить разницу
@@ -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);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -20,119 +20,8 @@ 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);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
#if defined(FUNCTIONAL_GRID_LAUNCH)
|
||||
#include "functional_grid_launch.inl"
|
||||
#else
|
||||
#include "macro_based_grid_launch.inl"
|
||||
#endif
|
||||
@@ -37,6 +37,7 @@ THE SOFTWARE.
|
||||
|
||||
#include "elfio/elfio.hpp"
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hcc_detail/program_state.hpp"
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
|
||||
@@ -0,0 +1,99 @@
|
||||
/*
|
||||
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"
|
||||
|
||||
// Internal header, do not percolate upwards.
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "hc.hpp"
|
||||
#include "trace_helper.h"
|
||||
|
||||
#include <iostream>
|
||||
#include <sstream>
|
||||
|
||||
namespace hip_impl
|
||||
{
|
||||
hc::accelerator_view lock_stream_hip_(
|
||||
hipStream_t& stream, void*& locked_stream)
|
||||
{ // This allocated but does not take ownership of locked_stream. If it is
|
||||
// not deleted elsewhere it will leak.
|
||||
using L = decltype(stream->lockopen_preKernelCommand());
|
||||
|
||||
HIP_INIT();
|
||||
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
locked_stream = new L{stream->lockopen_preKernelCommand()};
|
||||
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_KCMD)) ||
|
||||
HIP_PROFILE_API ||
|
||||
(COMPILE_HIP_DB && (HIP_TRACE_API & (1<<TRACE_ALL)))) {
|
||||
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::string fullStr;
|
||||
recordApiTrace(&fullStr, os.str());
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void unlock_stream_hip_(
|
||||
hipStream_t stream,
|
||||
void* locked_stream,
|
||||
const char* kernel_name,
|
||||
hc::accelerator_view* acc_v)
|
||||
{ // Precondition: acc_v is the accelerator_view associated with stream
|
||||
// which is guarded by locked_stream;
|
||||
// locked_stream is deletable.
|
||||
using L = decltype(stream->lockopen_preKernelCommand());
|
||||
|
||||
stream->lockclose_postKernelCommand(kernel_name, acc_v);
|
||||
|
||||
delete static_cast<L*>(locked_stream);
|
||||
locked_stream = nullptr;
|
||||
if(HIP_PROFILE_API) {
|
||||
MARKER_END();
|
||||
}
|
||||
}
|
||||
}
|
||||
Ссылка в новой задаче
Block a user