Skip to content

GitLab

"driver/driver.hip.cpp" did not exist on "d2a488ddeca43e1b6bf6f1c3ceb4abf067c49962"
Unverified Commit 97e6d514 authored by Chao Liu's avatar Chao Liu Committed by GitHub
Browse files

Merge pull request #4 from ROCmSoftwarePlatform/separate_online_compile 

Separate online compile
parents f6edda61 7b1ec41e
Hide whitespace changes
Inline Side-by-side
Showing with 62 additions and 54 deletions
external/half/include/half.hpp
View file @ 97e6d514
......@@ -2404,8 +2404,7 @@ unsigned int gamma(unsigned int arg)
0.0114684895434781459556 }; double t = arg + 4.65, s = p[0]; for(unsigned int i=0; i<5; ++i)
s += p[i+1] / (arg+i);
return std::log(s) + (arg-0.5)*std::log(t) - t;
*/ static const f31 pi(0xC90FDAA2, 1),
lbe(0xB8AA3B29, 0);
*/ static const f31 pi(0xC90FDAA2, 1), lbe(0xB8AA3B29, 0);
unsigned int abs = arg & 0x7FFF, sign = arg & 0x8000;
bool bsign = sign != 0;
f31 z(abs), x = sign ? (z + f31(0x80000000, 0)) : z, t = x + f31(0x94CCCCCD, 2),
......
host/CMakeLists.txt
View file @ 97e6d514
add_subdirectory(host_tensor)
add_subdirectory(online_compilation)
add_subdirectory(online_compile)
add_subdirectory(driver_offline)
add_subdirectory(driver_online)
host/driver_offline/CMakeLists.txt
View file @ 97e6d514
include_directories(BEFORE
include
${PROJECT_SOURCE_DIR}/host/host_tensor/include
${PROJECT_SOURCE_DIR}/host/solver/include
${PROJECT_SOURCE_DIR}/composable_kernel/include
${PROJECT_SOURCE_DIR}/composable_kernel/include/utility
${PROJECT_SOURCE_DIR}/composable_kernel/include/tensor_description
......
host/driver_online/CMakeLists.txt
View file @ 97e6d514
include_directories(BEFORE
include
${PROJECT_BINARY_DIR}/host/online_compilation/include
${PROJECT_SOURCE_DIR}/host/online_compilation/include
${PROJECT_BINARY_DIR}/host/online_compile/include
${PROJECT_SOURCE_DIR}/host/online_compile/include
${PROJECT_SOURCE_DIR}/host/host_tensor/include
${PROJECT_SOURCE_DIR}/host/solver/include
${PROJECT_SOURCE_DIR}/composable_kernel/include
${PROJECT_SOURCE_DIR}/composable_kernel/include/utility
${PROJECT_SOURCE_DIR}/composable_kernel/include/tensor_description
......@@ -18,4 +19,4 @@ set(CONV_FWD_DRIVER_ONLINE_SOURCE conv_fwd_driver_online.cpp)
add_executable(conv_fwd_driver_online ${CONV_FWD_DRIVER_ONLINE_SOURCE})
target_link_libraries(conv_fwd_driver_online PRIVATE host_tensor)
target_link_libraries(conv_fwd_driver_online PRIVATE online_compilation)
target_link_libraries(conv_fwd_driver_online PRIVATE online_compile)
host/driver_online/conv_fwd_driver_online.cpp
View file @ 97e6d514
......@@ -39,11 +39,11 @@ int main(int argc, char* argv[])
using size_t = std::size_t;
hipStream_t stream;
olCompile::Handle* handle;
online_compile::Handle* handle;
MY_HIP_CHECK(hipStreamCreate(&stream));
handle = new olCompile::Handle(stream);
handle = new online_compile::Handle(stream);
constexpr auto I0 = Number<0>{};
constexpr auto I1 = Number<1>{};
......
host/driver_online/include/online_device_dynamic_convolution_forward_implicit_gemm_v4r4_dlops_nchw_kcyx_nkhw.hpp
View file @ 97e6d514
......@@ -216,7 +216,7 @@ template <typename TInWei,
typename InLeftPads,
typename InRightPads>
void online_device_dynamic_convolution_forward_implicit_gemm_v4r4_dlops_nchw_kcyx_nkhw(
olCompile::Handle* handle,
online_compile::Handle* handle,
const InLengths& in_n_c_hi_wi_lengths,
const WeiLengths& wei_k_c_y_x_lengths,
const OutLengths& out_n_k_ho_wo_lengths,
......
host/driver_online/include/online_device_dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nchw_kcyx_nkhw.hpp
View file @ 97e6d514
......@@ -212,7 +212,7 @@ template <typename TInWei,
typename InLeftPads,
typename InRightPads>
void online_device_dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nchw_kcyx_nkhw(
olCompile::Handle* handle,
online_compile::Handle* handle,
const InLengths& in_n_c_hi_wi_lengths,
const WeiLengths& wei_k_c_y_x_lengths,
const OutLengths& out_n_k_ho_wo_lengths,
......
host/driver_online/include/online_device_dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk.hpp
View file @ 97e6d514
......@@ -213,7 +213,7 @@ template <typename TInWei,
typename InLeftPads,
typename InRightPads>
void online_device_dynamic_convolution_forward_implicit_gemm_v4r4_xdlops_nhwc_kyxc_nhwk(
olCompile::Handle* handle,
online_compile::Handle* handle,
const InLengths& in_n_hi_wi_c_lengths,
const WeiLengths& wei_k_y_x_c_lengths,
const OutLengths& out_n_ho_wo_k_lengths,
......
host/driver_online/include/online_device_dynamic_convolution_forward_implicit_gemm_v6r1_dlops_nchw_kcyx_nkhw.hpp
View file @ 97e6d514
......@@ -20,7 +20,7 @@ template <typename TInWei,
typename InLeftPads,
typename InRightPads>
void online_device_dynamic_convolution_forward_implicit_gemm_v6r1_dlops_nchw_kcyx_nkhw(
olCompile::Handle* handle,
online_compile::Handle* handle,
const InLengths& in_n_c_hi_wi_lengths,
const WeiLengths& wei_k_c_y_x_lengths,
const OutLengths& out_n_k_ho_wo_lengths,
......@@ -100,13 +100,13 @@ void online_device_dynamic_convolution_forward_implicit_gemm_v6r1_dlops_nchw_kcy
"dynamic_convolution_forward_implicit_gemm_v6r1_dlops_nchw_kcyx_nkhw.cpp";
std::string algo_name = "implicit_gemm_conv_fwd_v6r1_dlops_nchw";
std::string compile_param_string = " -std=c++17 " + compile_param.GetCompileParameterString();
std::string compile_param_string = get_ck_hip_online_compile_common_flag() + compile_param.GetCompileParameterString();
std::string network_config = compile_param_string;
std::vector<float> kernel1_times;
std::vector<float> kernel2_times;
for(index_t i = 0; i < nrepeat; ++i)
for(index_t i = 0; i < nrepeat + 1; ++i)
{
KernelTimer timer1, timer2;
std::string kernel_name;
......@@ -164,11 +164,11 @@ void online_device_dynamic_convolution_forward_implicit_gemm_v6r1_dlops_nchw_kcy
auto ave_time1 =
std::accumulate(
std::next(kernel1_times.begin()), kernel1_times.end(), 0., std::plus<float>{}) /
(nrepeat - 1);
nrepeat;
auto ave_time2 =
std::accumulate(
std::next(kernel2_times.begin()), kernel2_times.end(), 0., std::plus<float>{}) /
(nrepeat - 1);
nrepeat;
float perf = (float)(conv_problem_desc.CalculateFlop()) /
(std::size_t(1000) * 1000 * 1000) / (ave_time1 + ave_time2);
......
host/driver_online/include/online_driver_common.hpp
View file @ 97e6d514
......@@ -3,6 +3,13 @@
namespace ck_driver {
inline auto get_ck_hip_online_compile_common_flag()
{
std::string param = " -std=c++17";
return param;
}
// greatest common divisor, aka highest common factor
inline int gcd(int x, int y)
{
......
host/online_compilation/CMakeLists.txt → host/online_compile/CMakeLists.txt
View file @ 97e6d514
......@@ -67,10 +67,10 @@ else()
set(OLC_DEBUG 0)
endif()
configure_file("${PROJECT_SOURCE_DIR}/host/online_compilation/include/config.h.in" "${PROJECT_BINARY_DIR}/host/online_compilation/include/config.h")
configure_file("${PROJECT_SOURCE_DIR}/host/online_compile/include/config.h.in" "${PROJECT_BINARY_DIR}/host/online_compile/include/config.h")
include_directories(BEFORE
${PROJECT_BINARY_DIR}/host/online_compilation/include
${PROJECT_BINARY_DIR}/host/online_compile/include
)
message(STATUS "Hip compiler flags: ${HIP_COMPILER_FLAGS}")
......@@ -97,7 +97,7 @@ set(ONLINE_COMPILATION_SOURCE
)
include_directories(BEFORE
${PROJECT_BINARY_DIR}/host/online_compilation/include
${PROJECT_BINARY_DIR}/host/online_compile/include
include
)
......@@ -152,17 +152,17 @@ add_custom_command(
)
## the library target
add_library(online_compilation SHARED ${ONLINE_COMPILATION_SOURCE})
add_library(online_compile SHARED ${ONLINE_COMPILATION_SOURCE})
target_include_directories(online_compilation PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/online_compilation/include/)
target_include_directories(online_compilation PRIVATE ${PROJECT_BINARY_DIR})
target_include_directories(online_compilation PRIVATE ${PROJECT_SOURCE_DIR}/external/half/include/)
target_include_directories(online_compile PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/online_compile/include/)
target_include_directories(online_compile PRIVATE ${PROJECT_BINARY_DIR})
target_include_directories(online_compile PRIVATE ${PROJECT_SOURCE_DIR}/external/half/include/)
target_link_libraries(online_compilation PRIVATE hip::device)
target_link_libraries(online_compilation INTERFACE hip::host)
target_link_libraries(online_compilation PRIVATE Boost::filesystem)
target_link_libraries(online_compile PRIVATE hip::device)
target_link_libraries(online_compile INTERFACE hip::host)
target_link_libraries(online_compile PRIVATE Boost::filesystem)
target_compile_features(online_compilation PUBLIC)
set_target_properties(online_compilation PROPERTIES POSITION_INDEPENDENT_CODE ON)
target_compile_features(online_compile PUBLIC)
set_target_properties(online_compile PROPERTIES POSITION_INDEPENDENT_CODE ON)
install(TARGETS online_compilation LIBRARY DESTINATION lib)
install(TARGETS online_compile LIBRARY DESTINATION lib)
host/online_compilation/hip_utility/binary_cache.cpp → host/online_compile/hip_utility/binary_cache.cpp
View file @ 97e6d514
......@@ -35,7 +35,7 @@
#include <fstream>
#include <iostream>
namespace olCompile {
namespace online_compile {
OLC_DECLARE_ENV_VAR(OLC_DISABLE_CACHE)
OLC_DECLARE_ENV_VAR(HOME)
......@@ -62,14 +62,14 @@ boost::filesystem::path GetCachePath()
return user_path;
}
static bool IsCacheDisabled() { return olCompile::IsEnabled(OLC_DISABLE_CACHE{}); }
static bool IsCacheDisabled() { return online_compile::IsEnabled(OLC_DISABLE_CACHE{}); }
boost::filesystem::path
GetCacheFile(const std::string& device, const std::string& name, const std::string& args)
{
// std::string filename = (is_kernel_str ? olCompile::md5(name) : name) + ".o";
// std::string filename = (is_kernel_str ? online_compile::md5(name) : name) + ".o";
std::string filename = name + ".o";
return GetCachePath() / olCompile::md5(device + ":" + args) / filename;
return GetCachePath() / online_compile::md5(device + ":" + args) / filename;
}
boost::filesystem::path LoadBinary(const TargetProperties& target,
......@@ -77,7 +77,7 @@ boost::filesystem::path LoadBinary(const TargetProperties& target,
const std::string& name,
const std::string& args)
{
if(olCompile::IsCacheDisabled())
if(online_compile::IsCacheDisabled())
return {};
(void)num_cu;
......@@ -97,7 +97,7 @@ void SaveBinary(const boost::filesystem::path& binary_path,
const std::string& name,
const std::string& args)
{
if(olCompile::IsCacheDisabled())
if(online_compile::IsCacheDisabled())
{
boost::filesystem::remove(binary_path);
}
......@@ -109,4 +109,4 @@ void SaveBinary(const boost::filesystem::path& binary_path,
}
}
} // namespace olCompile
} // namespace online_compile
host/online_compilation/hip_utility/exec_utils.cpp → host/online_compile/hip_utility/exec_utils.cpp
View file @ 97e6d514
......@@ -38,7 +38,7 @@
#include <sys/wait.h>
#endif // __linux__
namespace olCompile {
namespace online_compile {
namespace exec {
int Run(const std::string& p, std::istream* in, std::ostream* out)
......@@ -53,7 +53,7 @@ int Run(const std::string& p, std::istream* in, std::ostream* out)
OLC_MANAGE_PTR(FILE*, pclose) pipe{popen(p.c_str(), file_mode)};
if(!pipe)
throw std::runtime_error("olCompile::exec::Run(): popen(" + p + ", " + file_mode +
throw std::runtime_error("online_compile::exec::Run(): popen(" + p + ", " + file_mode +
") failed");
if(redirect_stdin || redirect_stdout)
......@@ -74,7 +74,7 @@ int Run(const std::string& p, std::istream* in, std::ostream* out)
buffer[in->gcount()] = 0;
if(fputs(buffer.data(), pipe.get()) == EOF)
throw std::runtime_error("olCompile::exec::Run(): fputs() failed");
throw std::runtime_error("online_compile::exec::Run(): fputs() failed");
}
}
}
......@@ -90,4 +90,4 @@ int Run(const std::string& p, std::istream* in, std::ostream* out)
}
} // namespace exec
} // namespace olCompile
} // namespace online_compile
host/online_compilation/hip_utility/handlehip.cpp → host/online_compile/hip_utility/handlehip.cpp
View file @ 97e6d514
......@@ -50,7 +50,7 @@
OLC_DECLARE_ENV_VAR(OLC_DEVICE_CU)
namespace olCompile {
namespace online_compile {
std::size_t GetAvailableMemory()
{
......@@ -182,24 +182,24 @@ KernelInvoke Handle::Run(Kernel k) const { return k.Invoke(this->GetStream()); }
Program Handle::LoadProgram(const std::string& program_name, std::string params) const
{
if((!olCompile::EndsWith(program_name, ".mlir-cpp")) &&
(!olCompile::EndsWith(program_name, ".mlir")))
if((!online_compile::EndsWith(program_name, ".mlir-cpp")) &&
(!online_compile::EndsWith(program_name, ".mlir")))
{
params += " -mcpu=" + this->GetTargetProperties().Name();
}
auto hsaco = olCompile::LoadBinary(
auto hsaco = online_compile::LoadBinary(
this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
if(hsaco.empty())
{
auto p = HIPOCProgram{program_name, params, this->GetTargetProperties()};
auto path = olCompile::GetCachePath() / boost::filesystem::unique_path();
auto path = online_compile::GetCachePath() / boost::filesystem::unique_path();
if(p.IsCodeObjectInMemory())
olCompile::WriteFile(p.GetCodeObjectBlob(), path);
online_compile::WriteFile(p.GetCodeObjectBlob(), path);
else
boost::filesystem::copy_file(p.GetCodeObjectPathname(), path);
olCompile::SaveBinary(path, this->GetTargetProperties(), program_name, params);
online_compile::SaveBinary(path, this->GetTargetProperties(), program_name, params);
return p;
}
......@@ -245,7 +245,7 @@ std::size_t Handle::GetGlobalMemorySize() const
std::size_t Handle::GetMaxComputeUnits() const
{
int result;
const char* const num_cu = olCompile::GetStringEnv(OLC_DEVICE_CU{});
const char* const num_cu = online_compile::GetStringEnv(OLC_DEVICE_CU{});
if(num_cu != nullptr && strlen(num_cu) > 0)
{
return boost::lexical_cast<std::size_t>(num_cu);
......@@ -282,4 +282,4 @@ std::ostream& Handle::Print(std::ostream& os) const
return os;
}
} // namespace olCompile
} // namespace online_compile
host/online_compilation/hip_utility/hip_build_utils.cpp → host/online_compile/hip_utility/hip_build_utils.cpp
View file @ 97e6d514
......@@ -45,7 +45,7 @@ OLC_DECLARE_ENV_VAR(OLC_DEBUG_HIP_DUMP)
#define OLC_HIP_COMPILER "/opt/rocm/llvm/bin/clang++"
namespace olCompile {
namespace online_compile {
bool IsHccCompiler()
{
......@@ -155,12 +155,12 @@ static boost::filesystem::path HipBuildImpl(boost::optional<TmpDir>& tmp_dir,
params += " -mllvm -amdgpu-function-calls=false";
}
if(olCompile::IsEnabled(OLC_DEBUG_HIP_VERBOSE{}))
if(online_compile::IsEnabled(OLC_DEBUG_HIP_VERBOSE{}))
{
params += " -v";
}
if(olCompile::IsEnabled(OLC_DEBUG_HIP_DUMP{}))
if(online_compile::IsEnabled(OLC_DEBUG_HIP_DUMP{}))
{
if(IsHccCompiler())
{
......@@ -247,7 +247,7 @@ static external_tool_version_t HipCompilerVersionImpl()
break;
std::stringstream out;
if(olCompile::exec::Run(path + " --version", nullptr, &out) != 0)
if(online_compile::exec::Run(path + " --version", nullptr, &out) != 0)
break;
std::string line;
......@@ -343,4 +343,4 @@ bool operator<=(const external_tool_version_t& lhs, const external_tool_version_
return !(lhs > rhs);
}
} // namespace olCompile
} // namespace online_compile
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment