From 84b36f2fc6a76d48070f7ece27e22d50aaad08fa Mon Sep 17 00:00:00 2001 From: Stella Laurenzo Date: Mon, 21 Apr 2025 20:04:26 -0700 Subject: [PATCH] WIP: Windows compatibility changes for C++ code. Rename incorrect find/replace ABC_dims back to abc. --- clients/common/blis_interface.cpp | 2 +- clients/common/hipblaslt_init_device.cpp | 18 +- clients/common/hipblaslt_parse_data.cpp | 5 +- clients/common/utility.cpp | 88 ++++++-- clients/gtest/hipblaslt_test.cpp | 16 +- clients/include/TensorDataManipulation.hpp | 4 + clients/include/datatype_interface.hpp | 2 + clients/include/hipblaslt_arguments.hpp | 2 +- clients/include/hipblaslt_init.hpp | 4 +- clients/include/hipblaslt_random.hpp | 67 ++++-- clients/include/testing_matmul.hpp | 45 ++-- library/CMakeLists.txt | 2 +- library/include/hipblaslt/hipblaslt-ext.hpp | 14 +- .../amd_detail/hipblaslt-ext-op-internal.hpp | 4 +- library/src/amd_detail/hipblaslt-ext-op.cpp | 35 +-- library/src/amd_detail/hipblaslt-ext.cpp | 6 +- .../rocblaslt/include/rocblaslt-auxiliary.h | 3 +- .../amd_detail/rocblaslt/src/CMakeLists.txt | 2 - .../rocblaslt/src/include/logging.h | 5 +- .../rocblaslt/src/rocblaslt_auxiliary.cpp | 70 ++++-- .../rocblaslt/src/rocblaslt_transform.cpp | 70 +++--- .../amd_detail/rocblaslt/src/tensile_host.cpp | 208 ++++++++---------- .../src/amd_detail/rocblaslt/src/utility.cpp | 11 + library/src/hipblaslt_ostream.cpp | 12 +- library/src/include/hipblaslt_ostream.hpp | 2 +- .../client/include/DataInitialization.hpp | 12 + .../Serialization/PlaceholderLibrary.hpp | 5 +- .../Source/lib/source/ContractionSolution.cpp | 4 +- .../Tensile/Source/lib/source/MLPNet.cpp | 1 + .../lib/source/hip/HipSolutionAdapter.cpp | 2 +- tensilelite/rocisa/rocisa/src/helper.cpp | 2 +- 31 files changed, 434 insertions(+), 289 deletions(-) diff --git a/clients/common/blis_interface.cpp b/clients/common/blis_interface.cpp index ad747b4bda..5d62cc0816 100644 --- a/clients/common/blis_interface.cpp +++ b/clients/common/blis_interface.cpp @@ -29,7 +29,7 @@ void setup_blis() { -#ifndef WIN32 +#ifndef _WIN32 bli_init(); #endif } diff --git a/clients/common/hipblaslt_init_device.cpp b/clients/common/hipblaslt_init_device.cpp index 5f4443ed98..0101828145 100644 --- a/clients/common/hipblaslt_init_device.cpp +++ b/clients/common/hipblaslt_init_device.cpp @@ -116,7 +116,7 @@ __device__ int8_t random_hpl(size_t idx) } template -void hipblaslt_init_device(ABC abc, +void hipblaslt_init_device(ABC_dims abc, hipblaslt_initialization init, bool is_nan, T* A, @@ -140,11 +140,11 @@ void hipblaslt_init_device(ABC abc, switch(init) { case hipblaslt_initialization::rand_int: - if(abc == ABC::A || abc == ABC::C) + if(abc == ABC_dims::A || abc == ABC_dims::C) fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T { return random_int(idx); }); - else if(abc == ABC::B) + else if(abc == ABC_dims::B) { stride = std::max(lda * N, stride); fill_batch(A, M, N, lda, stride, batch_count, [stride, lda](size_t idx) -> T { @@ -158,14 +158,14 @@ void hipblaslt_init_device(ABC abc, break; case hipblaslt_initialization::trig_float: stride = std::max(lda * N, stride); - if(abc == ABC::A || abc == ABC::C) + if(abc == ABC_dims::A || abc == ABC_dims::C) fill_batch(A, M, N, lda, stride, batch_count, [M, N, stride, lda](size_t idx) -> T { auto b = idx / stride; auto j = (idx - b * stride) / lda; auto i = (idx - b * stride) - j * lda; return T(sin(double(i + j * M + b * M * N))); }); - else if(abc == ABC::B) + else if(abc == ABC_dims::B) fill_batch(A, M, N, lda, stride, batch_count, [M, N, stride, lda](size_t idx) -> T { auto b = idx / stride; auto j = (idx - b * stride) / lda; @@ -179,15 +179,15 @@ void hipblaslt_init_device(ABC abc, }); break; case hipblaslt_initialization::special: - if(abc == ABC::A) + if(abc == ABC_dims::A) fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T { return T(hipblasLtHalf(65280.0)); }); - else if(abc == ABC::B) + else if(abc == ABC_dims::B) fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T { return T(hipblasLtHalf(0.0000607967376708984375)); }); - else if(abc == ABC::C) + else if(abc == ABC_dims::C) fill_batch(A, M, N, lda, stride, batch_count, [](size_t idx) -> T { return T(pseudo_random_device(idx) % 10 + 1.f); }); @@ -213,7 +213,7 @@ void hipblaslt_init_device(ABC abc, } } -void hipblaslt_init_device(ABC abc, +void hipblaslt_init_device(ABC_dims abc, hipblaslt_initialization init, bool is_nan, void* A, diff --git a/clients/common/hipblaslt_parse_data.cpp b/clients/common/hipblaslt_parse_data.cpp index 2f42d84166..34a99d7ae8 100644 --- a/clients/common/hipblaslt_parse_data.cpp +++ b/clients/common/hipblaslt_parse_data.cpp @@ -37,13 +37,16 @@ // Parse YAML data static std::string hipblaslt_parse_yaml(const std::string& yaml) { + // TODO: This function is inherently unsafe because it returns a string vs an open + // file handle which will block further colliding creates. See comments in + // hipblaslt_tempname() and under no circumstances copy this to new code. std::string tmp = hipblaslt_tempname(); auto exepath = hipblaslt_exepath(); auto cmd = exepath + "hipblaslt_gentest.py --template " + exepath + "hipblaslt_template.yaml -o " + tmp + " " + yaml; hipblaslt_cerr << cmd << std::endl; -#ifdef WIN32 +#ifdef _WIN32 int status = std::system(cmd.c_str()); if(status == -1) exit(EXIT_FAILURE); diff --git a/clients/common/utility.cpp b/clients/common/utility.cpp index 917dee322d..f5189e5b3a 100644 --- a/clients/common/utility.cpp +++ b/clients/common/utility.cpp @@ -24,16 +24,23 @@ * *******************************************************************************/ -#include "utility.hpp" +#ifdef _WIN32 +#include +// Must include windows.h before dependent headers. +#include +#endif + #include "d_vector.hpp" +#include "utility.hpp" +#include #include #include +#include +#include #include #include #include -#include - #include "Tensile/Source/client/include/Utility.hpp" #if __has_include() @@ -48,27 +55,65 @@ namespace fs = std::experimental::filesystem; /* ============================================================================================ */ // Return path of this executable -std::string hipblaslt_exepath() +static std::string get_self_path() { - std::string pathstr; - char* path = realpath("/proc/self/exe", 0); - if(path) +#ifdef _WIN32 + std::string result(MAX_PATH + 1, '\0'); + DWORD length = 0; + for(;;) { - char* p = strrchr(path, '/'); - if(p) + length = GetModuleFileNameA(nullptr, result.data(), result.size()); + if(length < result.size() - 1) { - p[1] = 0; - pathstr = path; + result.resize(length); + return result; } - free(path); + result.resize(result.size() * 2); } - return pathstr; +#else + return std::string(realpath("/proc/self/exe", 0)); +#endif +} + +std::string hipblaslt_exepath() +{ + fs::path exepath(get_self_path()); + exepath.remove_filename(); + std::string result = exepath.string(); + if(result.empty()) + result.append("/"); + return result; } /* ============================================================================================ */ // Temp directory rooted random path +// TODO: This function is inherently unsafe because it returns a string vs an open +// file handle which will block further colliding creates. On Posix, this will leak +// a file handle for the life of the process. On Windows, there is no way to ensure that +// the created file name is unique without racing. To counter this on Windows, we +// also include the process id and a process specific counter in the generated name, +// as that will at least race consistently vs based on a random number generator +// collision. This and its consumers should be rewritten and under no circumstances +// copied to new code. std::string hipblaslt_tempname() { +#ifdef _WIN32 + static std::atomic counter; + // Generate "/tmp/rocblas-XXXXXX" like file name + const std::string alphanum = "0123456789ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuv"; + int stringlength = alphanum.length() - 1; + std::string uniquestr = "hipblaslt-"; + uniquestr.append(std::to_string(GetCurrentProcessId())); + uniquestr.append("-"); + uniquestr.append(std::to_string(counter.fetch_add(1))); + uniquestr.append("-"); + + for(auto n : {0, 1, 2, 3, 4, 5}) + uniquestr += alphanum.at(rand() % stringlength); + + fs::path tmpname = fs::temp_directory_path() / uniquestr; + return tmpname.string(); +#else char tmp[] = "/tmp/hipblaslt-XXXXXX"; int fd = mkostemp(tmp, O_CLOEXEC); if(fd == -1) @@ -78,6 +123,7 @@ std::string hipblaslt_tempname() } return std::string(tmp); +#endif } /* ============================================================================================ */ @@ -145,7 +191,7 @@ double get_time_us_no_sync(void) /* ============================================================================================ */ /* device query and print out their ID and name; return number of compute-capable devices. */ -int64_t query_device_property(int device_id, hipDeviceProp_t &props) +int64_t query_device_property(int device_id, hipDeviceProp_t& props) { int device_count; hipblasStatus_t status = (hipblasStatus_t)hipGetDeviceCount(&device_count); @@ -230,6 +276,15 @@ hipblaslt_local_handle::hipblaslt_local_handle() #endif } +static void portable_setenv(const char* name, const char* value) +{ +#ifdef _WIN32 + _putenv_s(name, value); +#else + setenv(name, value, /*overwrite=*/true); +#endif +} + hipblaslt_local_handle::hipblaslt_local_handle(const Arguments& arg) : hipblaslt_local_handle() { @@ -239,7 +294,8 @@ hipblaslt_local_handle::hipblaslt_local_handle(const Arguments& arg) if(sol_selec_env) m_sol_selec_saved_status = std::string(sol_selec_env); m_sol_selec_env_set = true; - setenv("TENSILE_SOLUTION_SELECTION_METHOD", std::to_string(arg.tensile_solution_selection_method).c_str(), true); + portable_setenv("TENSILE_SOLUTION_SELECTION_METHOD", + std::to_string(arg.tensile_solution_selection_method).c_str()); } // memory guard control, with multi-threading should not change values across threads d_vector_set_pad_length(arg.pad); @@ -249,7 +305,7 @@ hipblaslt_local_handle::~hipblaslt_local_handle() { if(m_sol_selec_env_set) { - setenv("TENSILE_SOLUTION_SELECTION_METHOD", m_sol_selec_saved_status.c_str(), true); + portable_setenv("TENSILE_SOLUTION_SELECTION_METHOD", m_sol_selec_saved_status.c_str()); } hipblasLtDestroy(m_handle); } diff --git a/clients/gtest/hipblaslt_test.cpp b/clients/gtest/hipblaslt_test.cpp index 73cfa5c816..caaaf76a2a 100644 --- a/clients/gtest/hipblaslt_test.cpp +++ b/clients/gtest/hipblaslt_test.cpp @@ -30,7 +30,7 @@ #include #include #include -#ifdef WIN32 +#ifdef _WIN32 #include #define strcasecmp(A, B) _stricmp(A, B) #else @@ -135,7 +135,7 @@ static thread_local struct volatile sig_atomic_t enabled = false; // sigjmp_buf describing stack frame to go back to -#ifndef WIN32 +#ifndef _WIN32 sigjmp_buf sigjmp_buf_; #else jmp_buf sigjmp_buf_; @@ -161,7 +161,7 @@ extern "C" void hipblaslt_test_signal_handler(int sig) return; } -#ifndef WIN32 +#ifndef _WIN32 // If this is an alarm timeout, we abort if(sig == SIGALRM) { @@ -180,7 +180,7 @@ extern "C" void hipblaslt_test_signal_handler(int sig) // it is better than crashing. t_handler.signal = sig; errno = saved_errno; -#ifndef WIN32 +#ifndef _WIN32 siglongjmp(t_handler.sigjmp_buf_, true); #else longjmp(t_handler.sigjmp_buf_, true); @@ -190,7 +190,7 @@ extern "C" void hipblaslt_test_signal_handler(int sig) // Set up signal handlers void hipblaslt_test_sigaction() { -#ifndef WIN32 +#ifndef _WIN32 struct sigaction act; act.sa_flags = 0; sigfillset(&act.sa_mask); @@ -219,7 +219,7 @@ void catch_signals_and_exceptions_as_failures(std::function test, bool s // Save the current handler (to allow nested calls to this function) auto old_handler = t_handler; -#ifndef WIN32 +#ifndef _WIN32 // Set up the return point, and handle siglongjmp returning back to here if(sigsetjmp(t_handler.sigjmp_buf_, true)) { @@ -237,7 +237,7 @@ void catch_signals_and_exceptions_as_failures(std::function test, bool s #endif else { -#ifndef WIN32 +#ifndef _WIN32 // Alarm to detect deadlocks or hangs if(set_alarm) alarm(test_timeout); @@ -260,7 +260,7 @@ void catch_signals_and_exceptions_as_failures(std::function test, bool s } } -#ifndef WIN32 +#ifndef _WIN32 // Cancel the alarm if it was set if(set_alarm) alarm(0); diff --git a/clients/include/TensorDataManipulation.hpp b/clients/include/TensorDataManipulation.hpp index 227a9d5146..6c53fc5b66 100644 --- a/clients/include/TensorDataManipulation.hpp +++ b/clients/include/TensorDataManipulation.hpp @@ -29,6 +29,10 @@ #include #include #include +#ifdef _WIN32 +#include +typedef SSIZE_T ssize_t; +#endif namespace Tensor { diff --git a/clients/include/datatype_interface.hpp b/clients/include/datatype_interface.hpp index 1430647dc6..b559ab673e 100644 --- a/clients/include/datatype_interface.hpp +++ b/clients/include/datatype_interface.hpp @@ -27,6 +27,8 @@ #pragma once #include +#include + union computeTypeInterface { float f32; diff --git a/clients/include/hipblaslt_arguments.hpp b/clients/include/hipblaslt_arguments.hpp index 0b39f0aeac..9ccccec6b1 100644 --- a/clients/include/hipblaslt_arguments.hpp +++ b/clients/include/hipblaslt_arguments.hpp @@ -296,7 +296,7 @@ struct Arguments // Function to read Arguments data from stream friend std::istream& operator>>(std::istream& str, Arguments& arg); -#ifdef WIN32 +#ifdef _WIN32 // Clang specific code template friend hipblaslt_internal_ostream& operator<<(hipblaslt_internal_ostream& os, diff --git a/clients/include/hipblaslt_init.hpp b/clients/include/hipblaslt_init.hpp index ce68547661..fb13963d40 100644 --- a/clients/include/hipblaslt_init.hpp +++ b/clients/include/hipblaslt_init.hpp @@ -36,14 +36,14 @@ #include #include -enum class ABC +enum class ABC_dims { A, B, C }; -void hipblaslt_init_device(ABC abc, +void hipblaslt_init_device(ABC_dims ABC_dims, hipblaslt_initialization init, bool is_nan, void* A, diff --git a/clients/include/hipblaslt_random.hpp b/clients/include/hipblaslt_random.hpp index 05e9c28778..6b958cdc93 100644 --- a/clients/include/hipblaslt_random.hpp +++ b/clients/include/hipblaslt_random.hpp @@ -69,6 +69,27 @@ inline void hipblaslt_seedrand() /*! \brief Random number generator which generates NaN values */ class hipblaslt_nan_rng { + // The C++ standard does not support std::uniform_int_distribution on all + // data types (notably not on char), although some libc++ implementation do. + // Notably, MSVC does not. Therefore, explicitly specialize non-standard + // types to narrow from a supported type. + template + static T portable_uniform_int_distribution() + { + return std::uniform_int_distribution{}(t_hipblaslt_rng); + } + template <> + char portable_uniform_int_distribution() + { + return static_cast(std::uniform_int_distribution{}(t_hipblaslt_rng)); + } + template <> + unsigned char portable_uniform_int_distribution() + { + return static_cast( + std::uniform_int_distribution{}(t_hipblaslt_rng)); + } + // Generate random NaN values template static T random_nan_data() @@ -80,7 +101,7 @@ class hipblaslt_nan_rng T fp; } x; do - x.u = std::uniform_int_distribution{}(t_hipblaslt_rng); + x.u = portable_uniform_int_distribution(); while(!(x.u & (((UINT_T)1 << SIG) - 1))); // Reject Inf (mantissa == 0) x.u |= (((UINT_T)1 << EXP) - 1) << SIG; // Exponent = all 1's return x.fp; // NaN with random bits @@ -91,7 +112,13 @@ class hipblaslt_nan_rng template {}, int> = 0> explicit operator T() { - return std::uniform_int_distribution{}(t_hipblaslt_rng); + return portable_uniform_int_distribution(); + } + + // // Random unsigned char + explicit operator unsigned char() + { + return static_cast(std::uniform_int_distribution{}(t_hipblaslt_rng)); } // Random signed char @@ -387,31 +414,36 @@ inline std::string random_string(size_t n) /* ============================================================================================ */ /*! \brief Random number generator which generates random values in normal distribution N(0,1) */ -namespace hipblaslt_norm_dist { +namespace hipblaslt_norm_dist +{ // XORWOW state structure - struct XorwowState { + struct XorwowState + { unsigned int x[5]; unsigned int counter; }; // Device function to initialize XORWOW state - __device__ void init_xorwow(XorwowState* state, unsigned int seed) { + __device__ void init_xorwow(XorwowState* state, unsigned int seed) + { unsigned int s = seed; - for (int i = 0; i < 5; ++i) { - s = s * 69069 + (i + 1); // Scramble seed + for(int i = 0; i < 5; ++i) + { + s = s * 69069 + (i + 1); // Scramble seed state->x[i] = s; } state->counter = seed ^ 362437; } // Device function for XORWOW RNG - __device__ unsigned int xorwow_rand(XorwowState* state) { + __device__ unsigned int xorwow_rand(XorwowState* state) + { unsigned int t = state->x[4]; unsigned int s = state->x[0]; - state->x[4] = state->x[3]; - state->x[3] = state->x[2]; - state->x[2] = state->x[1]; - state->x[1] = s; + state->x[4] = state->x[3]; + state->x[3] = state->x[2]; + state->x[2] = state->x[1]; + state->x[1] = s; t ^= t >> 2; t ^= t << 1; state->x[0] = t ^ s ^ (s << 4); @@ -420,16 +452,19 @@ namespace hipblaslt_norm_dist { } // Device function for uniform distribution - __device__ float xorwow_uniform(XorwowState* state) { + __device__ float xorwow_uniform(XorwowState* state) + { return xorwow_rand(state) / 4294967296.0f; } // Device function for Box-Muller normal distribution - __device__ float box_muller_normal(XorwowState* state) { + __device__ float box_muller_normal(XorwowState* state) + { float u1 = xorwow_uniform(state); float u2 = xorwow_uniform(state); - if (u1 < 1e-10f) u1 = 1e-10f; - float r = sqrtf(-2.0f * logf(u1)); + if(u1 < 1e-10f) + u1 = 1e-10f; + float r = sqrtf(-2.0f * logf(u1)); float theta = 2.0f * 3.1415926535f * u2; return r * cosf(theta); } diff --git a/clients/include/testing_matmul.hpp b/clients/include/testing_matmul.hpp index f63856efa8..2eeeeb6995 100644 --- a/clients/include/testing_matmul.hpp +++ b/clients/include/testing_matmul.hpp @@ -29,6 +29,7 @@ #include "TensorDataManipulation.hpp" #include "allclose.hpp" #include "cblas_interface.hpp" +#include "efficiency_monitor.hpp" #include "flops.hpp" #include "hipBuffer.hpp" #include "hipblaslt_datatype2string.hpp" @@ -40,7 +41,6 @@ #include "mxDataGen.hpp" #include "near.hpp" #include "norm.hpp" -#include "efficiency_monitor.hpp" #include "unit.hpp" #include "utility.hpp" #include @@ -49,6 +49,7 @@ #include #include #include +#include #include #include @@ -1815,16 +1816,16 @@ void testing_matmul_with_bias(const Arguments& arg, else { #endif - hipblaslt_init_device(ABC::A, - arg.initialization, - alpha_isnan_type(arg, Talpha), - dA[i].buf(), - A_row[i], - A_col[i], - (arg.swizzle_a) ? A_row[i] : lda[i], - TiA, - (arg.swizzle_a) ? A_row[i] * A_col[i] : stride_a[i], - num_batches[i]); + hipblaslt_init_device(ABC_dims::A, + arg.initialization, + alpha_isnan_type(arg, Talpha), + dA[i].buf(), + A_row[i], + A_col[i], + (arg.swizzle_a) ? A_row[i] : lda[i], + TiA, + (arg.swizzle_a) ? A_row[i] * A_col[i] : stride_a[i], + num_batches[i]); #ifdef USE_ROCROLLER } if(arg.scaleB == hipblaslt_scaling_format::Block) @@ -1865,20 +1866,20 @@ void testing_matmul_with_bias(const Arguments& arg, else { #endif - hipblaslt_init_device(ABC::B, - arg.initialization, - alpha_isnan_type(arg, Talpha), - dB[i].buf(), - B_row[i], - B_col[i], - ldb[i], - TiB, - stride_b[i], - num_batches[i]); + hipblaslt_init_device(ABC_dims::B, + arg.initialization, + alpha_isnan_type(arg, Talpha), + dB[i].buf(), + B_row[i], + B_col[i], + ldb[i], + TiB, + stride_b[i], + num_batches[i]); #ifdef USE_ROCROLLER } #endif - hipblaslt_init_device(ABC::C, + hipblaslt_init_device(ABC_dims::C, arg.initialization, beta_isnan_type(arg, Talpha), dC[i].buf(), diff --git a/library/CMakeLists.txt b/library/CMakeLists.txt index 00b3eac5d3..0b0b2d79f4 100644 --- a/library/CMakeLists.txt +++ b/library/CMakeLists.txt @@ -155,7 +155,7 @@ endif() # Target link libraries if(NOT BUILD_CUDA) # Target link libraries - target_link_libraries(hipblaslt PRIVATE hip::device ${DL_LIB}) + target_link_libraries(hipblaslt PRIVATE hip::device ${CMAKE_DL_LIBS}) endif() if(HIPBLASLT_ENABLE_MARKER) diff --git a/library/include/hipblaslt/hipblaslt-ext.hpp b/library/include/hipblaslt/hipblaslt-ext.hpp index 9fa3efc258..35c92bfd13 100644 --- a/library/include/hipblaslt/hipblaslt-ext.hpp +++ b/library/include/hipblaslt/hipblaslt-ext.hpp @@ -415,7 +415,7 @@ namespace hipblaslt_ext strideE2; //!< The aux batch stride. Only works if mode is set to aux related epilogues. float act0; //!< The activation value 1. Some activations might use it. float act1; //!< The activation value 2. - int activationType; //!< The activation type. Only works if mode is set to activation related epilogues. + int activationType; //!< The activation type. Only works if mode is set to activation related epilogues. } __attribute__((packed)); /*! \ingroup types_module @@ -425,8 +425,8 @@ namespace hipblaslt_ext { public: HIPBLASLT_EXPORT virtual ~GemmInstance(){}; - HIPBLASLT_EXPORT GemmInstance(const GemmInstance& rhs) = delete; - HIPBLASLT_EXPORT GemmInstance& operator=(const GemmInstance& rhs) = delete; + GemmInstance(const GemmInstance& rhs) = delete; + GemmInstance& operator=(const GemmInstance& rhs) = delete; HIPBLASLT_EXPORT GemmInstance(GemmInstance&& rhs) noexcept; HIPBLASLT_EXPORT GemmInstance& operator=(GemmInstance&& rhs) noexcept; @@ -771,9 +771,9 @@ namespace hipblaslt_ext void* D, hipblasLtMatrixLayout_t matD); - HIPBLASLT_EXPORT Gemm(const Gemm&) = delete; + Gemm(const Gemm&) = delete; HIPBLASLT_EXPORT Gemm(Gemm&&) noexcept; - HIPBLASLT_EXPORT Gemm& operator=(const Gemm&) = delete; + Gemm& operator=(const Gemm&) = delete; HIPBLASLT_EXPORT Gemm& operator=(Gemm&&) noexcept; /*! \ingroup library_module @@ -1027,9 +1027,9 @@ namespace hipblaslt_ext hipDataType typeC, hipDataType typeD, hipblasComputeType_t typeCompute); - HIPBLASLT_EXPORT GroupedGemm(const GroupedGemm&) = delete; + GroupedGemm(const GroupedGemm&) = delete; HIPBLASLT_EXPORT GroupedGemm(GroupedGemm&&) noexcept; - HIPBLASLT_EXPORT GroupedGemm& operator=(const GroupedGemm&) = delete; + GroupedGemm& operator=(const GroupedGemm&) = delete; HIPBLASLT_EXPORT GroupedGemm& operator=(GroupedGemm&&) noexcept; /*! \ingroup library_module diff --git a/library/src/amd_detail/hipblaslt-ext-op-internal.hpp b/library/src/amd_detail/hipblaslt-ext-op-internal.hpp index 0beea09907..48b85c9793 100644 --- a/library/src/amd_detail/hipblaslt-ext-op-internal.hpp +++ b/library/src/amd_detail/hipblaslt-ext-op-internal.hpp @@ -32,8 +32,8 @@ #include #include #include +#include #include -#include #include #include #include @@ -660,7 +660,7 @@ namespace hipblaslt_ext explicit ExtOpMasterLibrary(const std::string& libPath) : libPath(libPath) { - libDir = std::string(dirname(&this->libPath[0])); + libDir = std::filesystem::path(this->libPath).parent_path().string(); load(libPath); } diff --git a/library/src/amd_detail/hipblaslt-ext-op.cpp b/library/src/amd_detail/hipblaslt-ext-op.cpp index 86eca3f030..440a1778ff 100644 --- a/library/src/amd_detail/hipblaslt-ext-op.cpp +++ b/library/src/amd_detail/hipblaslt-ext-op.cpp @@ -32,11 +32,12 @@ #include #include #include +#include #include #include #include -#include #include +#include #include #include #include @@ -162,22 +163,22 @@ namespace return libPath; } - auto soPath = rocblaslt_internal_get_so_path("hipblaslt"); - std::string libPath(dirname(&soPath[0])); + auto soPath = rocblaslt_internal_get_so_path(); + std::filesystem::path libPath(std::filesystem::path(soPath).parent_path()); - if(rocblaslt_internal_test_path(libPath + "/../Tensile/library")) - libPath += "/../Tensile/library"; - else if(rocblaslt_internal_test_path(libPath + "library")) - libPath += "/library"; - else - libPath += "/hipblaslt/library"; + auto pathIfExists = [](std::filesystem::path p) -> std::optional { + if(std::filesystem::exists(p)) + return p; + return {}; + }; - libPath += "/hipblasltExtOpLibrary.dat"; - - if(rocblaslt_internal_test_path(libPath)) - { - return libPath; - } + if(auto p + = pathIfExists(libPath / ".." / "Tensile" / "library" / "hipblasltExtOpLibrary.dat")) + return p->string(); + if(auto p = pathIfExists(libPath / "library" / "hipblasltExtOpLibrary.dat")) + return p->string(); + if(auto p = pathIfExists(libPath / "hipblaslt" / "library" / "hipblasltExtOpLibrary.dat")) + return p->string(); return DEFAULT_EXT_OP_LIBRARY_PATH; } @@ -483,9 +484,9 @@ hipblasStatus_t hipblasltAMaxWithScaleRun(const hipDataType datatype, if(datatype != HIP_R_32F || scaleDatatype != HIP_R_8F_E4M3_FNUZ && scaleDatatype != HIP_R_8F_E5M2_FNUZ #ifdef ROCM_USE_FLOAT8 - && scaleDatatype != HIP_R_8F_E4M3 && scaleDatatype != HIP_R_8F_E5M2 + && scaleDatatype != HIP_R_8F_E4M3 && scaleDatatype != HIP_R_8F_E5M2 #endif - ) + ) { return HIPBLAS_STATUS_NOT_SUPPORTED; } diff --git a/library/src/amd_detail/hipblaslt-ext.cpp b/library/src/amd_detail/hipblaslt-ext.cpp index ec286e4da7..724d359e1c 100644 --- a/library/src/amd_detail/hipblaslt-ext.cpp +++ b/library/src/amd_detail/hipblaslt-ext.cpp @@ -300,7 +300,7 @@ namespace hipblaslt_ext class GemmTuningV2::GemmTuningImpl { public: - u_int16_t splitK = 0; + uint16_t splitK = 0; int16_t wgm = 0; }; @@ -325,7 +325,7 @@ namespace hipblaslt_ext GemmTuningV2::GemmTuningV2(GemmTuningV2&& tuning) = default; GemmTuningV2& GemmTuningV2::operator=(GemmTuningV2&& tuning) = default; - void GemmTuningV2::setSplitK(u_int16_t splitK) + void GemmTuningV2::setSplitK(uint16_t splitK) { pimpl->splitK = splitK; } @@ -335,7 +335,7 @@ namespace hipblaslt_ext pimpl->wgm = wgm; } - u_int16_t GemmTuningV2::getSplitK() const + uint16_t GemmTuningV2::getSplitK() const { return pimpl->splitK; } diff --git a/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h b/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h index bf052aa9a2..a2eff8684c 100644 --- a/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h +++ b/library/src/amd_detail/rocblaslt/include/rocblaslt-auxiliary.h @@ -403,7 +403,8 @@ std::string rocblaslt_internal_get_arch_name(); // for internal use of testing existence of path bool rocblaslt_internal_test_path(const std::string&); -std::string rocblaslt_internal_get_so_path(const std::string& keyword); +// Gets the absolute path of the so/dll/exe containing this function. +std::string rocblaslt_internal_get_so_path(); void rocblaslt_log_error(const char* func, const char* var, const char* msg); #endif diff --git a/library/src/amd_detail/rocblaslt/src/CMakeLists.txt b/library/src/amd_detail/rocblaslt/src/CMakeLists.txt index 09ecb695c3..15a4e1035d 100644 --- a/library/src/amd_detail/rocblaslt/src/CMakeLists.txt +++ b/library/src/amd_detail/rocblaslt/src/CMakeLists.txt @@ -146,8 +146,6 @@ if(NOT Tensile_SKIP_BUILD) add_subdirectory(kernels) endif() -set(DL_LIB dl) - # rocBLASLt source target_sources( hipblaslt diff --git a/library/src/amd_detail/rocblaslt/src/include/logging.h b/library/src/amd_detail/rocblaslt/src/include/logging.h index ffc32cfbaa..c899d43a67 100644 --- a/library/src/amd_detail/rocblaslt/src/include/logging.h +++ b/library/src/amd_detail/rocblaslt/src/include/logging.h @@ -44,10 +44,13 @@ #include #include #include -#include #include #include +#ifndef _WIN32 +#include +#endif + /************************************************************************************ * Profile kernel arguments ************************************************************************************/ diff --git a/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp b/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp index 3c00c9dc35..db0d3d3078 100644 --- a/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp +++ b/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp @@ -24,6 +24,17 @@ * * ************************************************************************ */ +#ifdef _WIN32 +#include +#include +#include +#else +#define _GNU_SOURCE +#include +#include +#include +#endif + #include "UserDrivenTuningParser.hpp" #include "definitions.h" #include "handle.h" @@ -33,13 +44,8 @@ #include "tensile_host.hpp" #include "utility.hpp" -#ifndef WIN32 -#include -#endif - #include #include -#include #include #define TO_STR2(x) #x @@ -2119,35 +2125,55 @@ std::string rocblaslt_internal_get_arch_name() bool rocblaslt_internal_test_path(const std::string& path) { -#ifdef WIN32 +#ifdef _WIN32 return ((_access(path.c_str(), 4) != -1) || (_access(path.c_str(), 6) != -1)); #else return access(path.c_str(), R_OK) == 0; #endif } -#ifndef WIN32 -int hipblaslt_dl_iterate_phdr_callback(struct dl_phdr_info* hdr_info, size_t size, void* data) +#ifdef _WIN32 +std::string rocblaslt_internal_get_so_path() { - // uncomment to see all dependent .so files - // fprintf(stderr, "hipblaslt so file: %s\n", hdr_info->dlpi_name); - std::pair* typedData - = reinterpret_cast*>(data); - if(hdr_info->dlpi_name && strstr(hdr_info->dlpi_name, typedData->second.c_str())) + HMODULE hModule = NULL; + if(!GetModuleHandleEx(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS, + // Should be the address of code in this library. + (LPCTSTR)rocblaslt_internal_get_so_path, + &hModule)) { - typedData->first.assign(hdr_info->dlpi_name); - return 1; + throw std::runtime_error("Cannot get module for function"); } - return 0; -} -#endif -std::string rocblaslt_internal_get_so_path(const std::string& keyword) + std::string path; + path.resize(256); + for(;;) + { + auto stored_size = GetModuleFileNameA(hModule, path.data(), path.size()); + if(stored_size < path.size()) + { + // Success: size to what was stored (which does not include NUL). + path.resize(stored_size); + return path; + } + // Insufficient size. + path.resize(path.size() * 2); + } +} +#else +std::string rocblaslt_internal_get_so_path() { - std::pair result{"", keyword}; - dl_iterate_phdr(hipblaslt_dl_iterate_phdr_callback, &result); - return result.first; + Dl_info info; + if(dladdr(reinterpret_cast(&rocblaslt_internal_get_so_path), &info) == 0) + { + throw std::runtime_error("Cannot get address of module function"); + } + if(!info.dli_fname) + { + throw std::runtime_error("Containing binary does not have a file system path"); + } + return std::string(info.dli_fname); } +#endif void rocblaslt_log_error(const char* func, const char* var, const char* msg) { diff --git a/library/src/amd_detail/rocblaslt/src/rocblaslt_transform.cpp b/library/src/amd_detail/rocblaslt/src/rocblaslt_transform.cpp index be6e6f3dd0..e209d6ea91 100644 --- a/library/src/amd_detail/rocblaslt/src/rocblaslt_transform.cpp +++ b/library/src/amd_detail/rocblaslt/src/rocblaslt_transform.cpp @@ -32,37 +32,45 @@ #include #include #include -#include + +#include #include #include +#include #include #include #include namespace { - std::string transformCodeObjectPath() + std::filesystem::path transformCodeObjectPath() { +#ifdef _WIN32 + constexpr char DEFAULT_CO_PATH[] + = "C:\\opt\\rocm\\bin\\hipblaslt\\library\\hipblasltTransform.hsaco"; +#else constexpr char DEFAULT_CO_PATH[] = "/opt/rocm/lib/hipblaslt/library/hipblasltTransform.hsaco"; - auto soPath = rocblaslt_internal_get_so_path("hipblaslt"); - std::string libPath(dirname(&soPath[0])); +#endif - if(rocblaslt_internal_test_path(libPath + "/../Tensile/library")) - libPath += "/../Tensile/library"; - else if(rocblaslt_internal_test_path(libPath + "library")) - libPath += "/library"; - else - libPath += "/hipblaslt/library"; + std::string soPath = rocblaslt_internal_get_so_path(); + std::filesystem::path libPath(std::filesystem::path(soPath).parent_path()); - libPath += "/hipblasltTransform.hsaco"; + auto pathIfExists = [](std::filesystem::path p) -> std::optional { + if(std::filesystem::exists(p)) + return p; + return {}; + }; - if(rocblaslt_internal_test_path(libPath)) - { - return libPath; - } + if(auto p + = pathIfExists(libPath / ".." / "Tensile" / "library" / "hipblasltTransform.hsaco")) + return *p; + if(auto p = pathIfExists(libPath / "library" / "hipblasltTransform.hsaco")) + return *p; + if(auto p = pathIfExists(libPath / "hipblaslt" / "library" / "hipblasltTransform.hsaco")) + return *p; - return DEFAULT_CO_PATH; + return std::filesystem::path(DEFAULT_CO_PATH); } TensileLite::hip::SolutionAdapter& transformAdapter() @@ -77,7 +85,7 @@ namespace adapters.emplace_back(new TensileLite::hip::SolutionAdapter); } auto coPath = transformCodeObjectPath(); - const std::string coFolder = dirname(&coPath[0]); + const std::string coFolder = coPath.parent_path().string(); try { for(auto& adp : adapters) @@ -93,7 +101,7 @@ namespace return adapters; }(); - int deviceId{}; + int deviceId{}; HIP_CHECK_EXC(hipGetDevice(&deviceId)); return *adapter.at(deviceId); } @@ -137,10 +145,10 @@ namespace hipStream_t stream, const std::string& kernelName) { - constexpr auto TileM = RowMajC ? NumThreadsM : NumThreadsM * VectorWidth; - constexpr auto TileN = RowMajC ? NumThreadsN * VectorWidth : NumThreadsN; - const auto numWg = (m / TileM + !!(m % TileM)) * (n / TileN + !!(n % TileN)); - constexpr auto numWorkitems = NumThreadsM * NumThreadsN; + constexpr auto TileM = RowMajC ? NumThreadsM : NumThreadsM * VectorWidth; + constexpr auto TileN = RowMajC ? NumThreadsN * VectorWidth : NumThreadsN; + const auto numWg = (m / TileM + !!(m % TileM)) * (n / TileN + !!(n % TileN)); + constexpr auto numWorkitems = NumThreadsM * NumThreadsN; TensileLite::KernelArguments kArgs(false); if(scalarInDevice) @@ -192,16 +200,16 @@ namespace kArgs.appendAligned("transB", transB); } - constexpr auto NUM_WORKITEMS{NumThreadsM * NumThreadsN}; + constexpr auto NUM_WORKITEMS{NumThreadsM * NumThreadsN}; TensileLite::KernelInvocation invocation{kernelName, - "hipblasltTransform.hsaco", - false, - {NUM_WORKITEMS, 1, 1}, - {numWg, 1, batchSize}, - {numWg * NUM_WORKITEMS, 1, batchSize}, - 0, - kArgs}; - auto& adapter = transformAdapter(); + "hipblasltTransform.hsaco", + false, + {NUM_WORKITEMS, 1, 1}, + {numWg, 1, batchSize}, + {numWg * NUM_WORKITEMS, 1, batchSize}, + 0, + kArgs}; + auto& adapter = transformAdapter(); return adapter.launchKernel(invocation, stream, nullptr, nullptr); } diff --git a/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index ccccde9a81..50c42f06d4 100644 --- a/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -56,21 +56,18 @@ #include #include #include +#include #include #include #include +#include +#include #include #include +#include #include #include -#include -#include -#include -#include -#include -#include - #define HIPBLASLT_LIB_PATH "/opt/rocm/lib" #ifdef ENABLE_ROCTX @@ -262,11 +259,6 @@ RocblasltContractionProblem::RocblasltContractionProblem(hipblasOperation_t namespace { - std::string getHipblasltSoPath() - { - return rocblaslt_internal_get_so_path("libhipblaslt"); - } - static void assignAlphaBeta(rocisa::DataType type, const void* alphaPtr, const void* betaPtr, @@ -723,7 +715,7 @@ namespace : "", problem.tensor(TensileLite::ContractionProblemGemm::TENSOR::E).strides().size() ? std::to_string( - problem.tensor(TensileLite::ContractionProblemGemm::TENSOR::E).strides()[1]) + problem.tensor(TensileLite::ContractionProblemGemm::TENSOR::E).strides()[1]) : "", "--stride_a", problem.a().strides()[2], @@ -738,7 +730,7 @@ namespace : "", problem.tensor(TensileLite::ContractionProblemGemm::TENSOR::E).strides().size() ? std::to_string( - problem.tensor(TensileLite::ContractionProblemGemm::TENSOR::E).strides()[2]) + problem.tensor(TensileLite::ContractionProblemGemm::TENSOR::E).strides()[2]) : "", "--alpha", ToString(inputs.alpha), @@ -1891,7 +1883,7 @@ namespace return m_devicePropMap.at(deviceName); } #else - auto& get_device_property() const + auto& get_device_property() const { return m_deviceProp; } @@ -1901,24 +1893,23 @@ namespace return m_adapters; } - /******************************************************* - * Testpath() tests that a path exists and is readable * - *******************************************************/ - static bool TestPath(const std::string& path) - { - return rocblaslt_internal_test_path(path); - } - /********************************************************************* * Initialize adapter and library according to environment variables * * and default paths based on librocblaslt.so location and GPU * *********************************************************************/ void initialize(TensileLite::hip::SolutionAdapter& adapter, int32_t deviceId) { - std::string path; -#ifndef WIN32 - path.reserve(PATH_MAX); + bool enableYaml = false; + bool staticLib = false; + bool lazyLoad = ROCBLASLT_TENSILE_LAZY_LOAD; +#ifdef TENSILE_YAML + enableYaml = true; #endif +#ifdef HIPBLASLT_STATIC_LIB + staticLib = true; +#endif + + std::filesystem::path path; // The name of the current GPU platform std::string processor = rocblaslt_internal_get_arch_name(); @@ -1936,28 +1927,36 @@ namespace } else { - path = HIPBLASLT_LIB_PATH; - // Find the location of librocblaslt.so // Fall back on hard-coded path if static library or not found + if(staticLib) + { + path = HIPBLASLT_LIB_PATH; + } + else + { + auto hipblaslt_so_path + = std::filesystem::path(rocblaslt_internal_get_so_path()); + path = hipblaslt_so_path.parent_path(); + } -#ifndef HIPBLASLT_STATIC_LIB - auto hipblaslt_so_path = getHipblasltSoPath(); - - if(hipblaslt_so_path.size()) - path = std::string{dirname(&hipblaslt_so_path[0])}; -#endif // ifndef HIPBLASLT_STATIC_LIB + auto pathIfExists + = [](std::filesystem::path p) -> std::optional { + if(std::filesystem::exists(p)) + return p; + return {}; + }; // Find the location of the libraries - if(TestPath(path + "/../Tensile/library")) - path += "/../Tensile/library"; - else if(TestPath(path + "library")) - path += "/library"; + if(auto p = pathIfExists(path / ".." / "Tensile" / "library")) + path = *p; + else if(auto p = pathIfExists(path / "library")) + path = *p; else - path += "/hipblaslt/library"; + path = path / "hipblaslt" / "library"; - if(TestPath(path + "/" + processor)) - path += "/" + processor; + if(auto p = pathIfExists(path / processor)) + path = *p; if(get_logger_layer_mode() & rocblaslt_layer_mode_log_info) { @@ -1967,63 +1966,31 @@ namespace } } - // only load modules for the current architecture - auto dir = path + "/*" + processor + "*co"; -#if ROCBLASLT_TENSILE_LAZY_LOAD == 0 - bool no_match = false; -#ifdef WIN32 - std::replace(dir.begin(), dir.end(), '/', '\\'); - WIN32_FIND_DATAA finddata; - HANDLE hfine = FindFirstFileA(dir.c_str(), &finddata); - if(hfine != INVALID_HANDLE_VALUE) + // only load modules for the current architecture (contains the processor + // string and ends in "co"). + if(!lazyLoad) { - do + bool no_match = true; + for(const auto& entry : std::filesystem::directory_iterator(path)) { - std::string codeObjectFile = path + "\\" + finddata.cFileName; - static_cast(adapter.loadCodeObjectFile(codeObjectFile.c_str())); - } while(FindNextFileA(hfine, &finddata)); - } - else - { - no_match = true; - } - FindClose(hfine); -#else - glob_t glob_result{}; - int g = glob(dir.c_str(), GLOB_NOSORT, nullptr, &glob_result); - if(!g) - { - for(size_t i = 0; i < glob_result.gl_pathc; ++i) - static_cast(adapter.loadCodeObjectFile(glob_result.gl_pathv[i])); - } - else if(g == GLOB_NOMATCH) - { - no_match = true; - } - else - { -#if 0 - // clang-format off - static std::ostream& once = std::cerr - << "\nrocblaslt warning: glob(\"" << dir << "\", ...) returned " - << (g == GLOB_ABORTED ? "GLOB_ABORTED" - : g == GLOB_NOSPACE ? "GLOB_NOSPACE" - : "an unknown error") - << "." << std::endl; - // clang-format on -#endif - } - globfree(&glob_result); -#endif - if(no_match) - { - // static rocblaslt_internal_ostream& once - // = rocblaslt_cerr - std::cerr << "\nrocblaslt warning: No paths matched " << dir - << ". Make sure that HIPBLASLT_TENSILE_LIBPATH is set correctly." - << std::endl; + auto filename = entry.path().filename(); + if(filename.string().find(processor) != std::string::npos + && filename.extension().string() == ".co") + { + static_cast(adapter.loadCodeObjectFile(entry.path().string())); + no_match = false; + } + } + if(no_match) + { + // static rocblaslt_internal_ostream& once + // = rocblaslt_cerr + std::cerr << "\nrocblaslt warning: No paths matched " << path + << ". Make sure that HIPBLASLT_TENSILE_LIBPATH is set correctly." + << std::endl; + } } -#endif + // We initialize a local static variable with a lambda function call to // avoid race conditions when multiple threads with different device IDs try // to initialize library. This ensures that only one thread initializes @@ -2031,21 +1998,34 @@ namespace // complete. static int once = [&] { // Determine library path - std::string tensileLibPath; -#if ROCBLASLT_TENSILE_LAZY_LOAD -#ifdef TENSILE_YAML - tensileLibPath = path + "/TensileLibrary_lazy_" + processor + ".yaml"; -#else - tensileLibPath = path + "/TensileLibrary_lazy_" + processor + ".dat"; -#endif -#else -#ifdef TENSILE_YAML - tensileLibPath = path + "/TensileLibrary_" + processor + ".yaml"; -#else - tensileLibPath = path + "/TensileLibrary_" + processor + ".dat"; -#endif -#endif - if(!TestPath(tensileLibPath)) + std::filesystem::path tensileLibPath; + if(lazyLoad) + { + if(enableYaml) + { + tensileLibPath + = path / (std::string("TensileLibrary_lazy_") + processor + ".yaml"); + } + else + { + tensileLibPath + = path / (std::string("TensileLibrary_lazy_") + processor + ".dat"); + } + } + else + { + if(enableYaml) + { + tensileLibPath + = path / (std::string("TensileLibrary_") + processor + ".yaml"); + } + else + { + tensileLibPath + = path / (std::string("TensileLibrary_") + processor + ".dat"); + } + } + if(!std::filesystem::exists(tensileLibPath)) { std::cerr << "\nrocblaslt error: Cannot read " << tensileLibPath << ": " << strerror(errno) << std::endl; @@ -2076,7 +2056,7 @@ namespace // Load library auto lib = TensileLite::LoadLibraryFilePreload( - tensileLibPath, std::vector{}); + tensileLibPath.string(), std::vector{}); #else // Get device prop hipDeviceProp_t prop; @@ -2085,7 +2065,7 @@ namespace // Load library auto lib = TensileLite::LoadLibraryFile( - tensileLibPath); + tensileLibPath.string()); #endif if(!lib) std::cerr << "\nrocblaslt error: Could not load " << tensileLibPath @@ -2095,12 +2075,12 @@ namespace using MSL = TensileLite::MasterSolutionLibrary; m_library = std::dynamic_pointer_cast(lib); - m_tensileLibPath = tensileLibPath; + m_tensileLibPath = tensileLibPath.string(); } return 0; }(); - static_cast(adapter.initializeLazyLoading(processor, path)); + static_cast(adapter.initializeLazyLoading(processor, path.string())); if(!m_library && once != 0) { diff --git a/library/src/amd_detail/rocblaslt/src/utility.cpp b/library/src/amd_detail/rocblaslt/src/utility.cpp index 76013ead2b..b233d8b1ad 100644 --- a/library/src/amd_detail/rocblaslt/src/utility.cpp +++ b/library/src/amd_detail/rocblaslt/src/utility.cpp @@ -25,7 +25,18 @@ *******************************************************************************/ #include "utility.hpp" #include +#include + +#ifdef _WIN32 +#include +#else #include +#endif + +#include +#include +#include + std::ostream* get_logger_os() { LoggerSingleton& s = LoggerSingleton::getInstance(); diff --git a/library/src/hipblaslt_ostream.cpp b/library/src/hipblaslt_ostream.cpp index 549772e82f..a4cf1b81b3 100644 --- a/library/src/hipblaslt_ostream.cpp +++ b/library/src/hipblaslt_ostream.cpp @@ -30,7 +30,7 @@ static void hipblaslt_abort_once [[noreturn]] (); #include #include #include -#ifdef WIN32 +#ifdef _WIN32 #include #include #include @@ -52,7 +52,7 @@ static void hipblaslt_abort_once [[noreturn]] (); // Abort function which is called only once by hipblaslt_abort static void hipblaslt_abort_once() { -#ifndef WIN32 +#ifndef _WIN32 // Make sure the alarm and abort actions are default signal(SIGALRM, SIG_DFL); signal(SIGABRT, SIG_DFL); @@ -108,7 +108,7 @@ std::shared_ptr hipblaslt_internal_ostream:: && std::is_same{}, "struct stat and file_id_t are not layout-compatible"); -#ifndef WIN32 +#ifndef _WIN32 // Get the device ID and inode, to detect common files if(fstat(fd, &statbuf)) { @@ -226,7 +226,7 @@ void hipblaslt_internal_ostream::worker::send(std::string str) // The future indicating when the operation has completed auto future = promise.get_future(); -#ifdef WIN32 +#ifdef _WIN32 // Passing an empty string will make the worker thread exit. // The below flag will be used to handle worker thread exit condition for Windows bool empty_string = str.empty(); @@ -247,7 +247,7 @@ void hipblaslt_internal_ostream::worker::send(std::string str) } // Wait for the task to be completed, to ensure flushed IO -#ifdef WIN32 +#ifdef _WIN32 if(empty_string) // Occassionaly this thread is not getting the promise set by the 'worker' thread during exit condition. // Added a timed wait to exit after one second, if we do not get the promise from worker thread. @@ -313,7 +313,7 @@ void hipblaslt_internal_ostream::worker::thread_function() hipblaslt_internal_ostream::worker::worker(int fd) { // The worker duplicates the file descriptor (RAII) -#ifdef WIN32 +#ifdef _WIN32 fd = _dup(fd); #else fd = fcntl(fd, F_DUPFD_CLOEXEC, 0); diff --git a/library/src/include/hipblaslt_ostream.hpp b/library/src/include/hipblaslt_ostream.hpp index 1b50f01114..6857fd2e52 100644 --- a/library/src/include/hipblaslt_ostream.hpp +++ b/library/src/include/hipblaslt_ostream.hpp @@ -47,7 +47,7 @@ #include #include #include -#ifdef WIN32 +#ifdef _WIN32 #include #include #include diff --git a/tensilelite/Tensile/Source/client/include/DataInitialization.hpp b/tensilelite/Tensile/Source/client/include/DataInitialization.hpp index 2305441974..c30edfb3a1 100644 --- a/tensilelite/Tensile/Source/client/include/DataInitialization.hpp +++ b/tensilelite/Tensile/Source/client/include/DataInitialization.hpp @@ -2601,13 +2601,25 @@ namespace TensileLite template <> inline Float8 DataInitialization::getValue() { +#if _WIN32 + //msvc's STL implementation follows [rand.req.genl](1.5), so Float8 as template arg + //is not allowed + return Float8(rocm_random_narrow_range{}()); +#else return rocm_random_narrow_range{}(); +#endif } template <> inline BFloat8 DataInitialization::getValue() { +#if _WIN32 + //msvc's STL implementation follows [rand.req.genl](1.5), so BFloat8 as template arg + //is not allowed + return BFloat8(rocm_random_narrow_range{}()); +#else return rocm_random_narrow_range{}(); +#endif } template <> diff --git a/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp b/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp index cdf305510e..36235e4346 100644 --- a/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp +++ b/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp @@ -32,8 +32,9 @@ #include #include //Replace std::regex, as it crashes when matching long lines(GCC Bug #86164). -#ifdef WIN32 +#ifdef _WIN32 #include "shlwapi.h" +#pragma comment(lib, "shlwapi.lib") #else #include #endif @@ -73,7 +74,7 @@ namespace TensileLite for(auto condition : ctx->preloaded) { std::string pattern = RegexPattern(condition); -#ifdef WIN32 +#ifdef _WIN32 if(PathMatchSpecA(lib.filePrefix.c_str(), pattern.c_str())) #else if(fnmatch(pattern.c_str(), lib.filePrefix.c_str(), 0) == 0) diff --git a/tensilelite/Tensile/Source/lib/source/ContractionSolution.cpp b/tensilelite/Tensile/Source/lib/source/ContractionSolution.cpp index adee51e8d3..05da6ddaf5 100644 --- a/tensilelite/Tensile/Source/lib/source/ContractionSolution.cpp +++ b/tensilelite/Tensile/Source/lib/source/ContractionSolution.cpp @@ -2285,7 +2285,9 @@ namespace TensileLite gsuTemp++; name += "_PostGSU" - + std::to_string(std::min((unsigned long)gsuTemp, sizeMapping.globalSplitUPGR)); + + std::to_string( + std::min(static_cast(gsuTemp), + sizeMapping.globalSplitUPGR)); name += "_VW" + std::to_string(vw); diff --git a/tensilelite/Tensile/Source/lib/source/MLPNet.cpp b/tensilelite/Tensile/Source/lib/source/MLPNet.cpp index 131e2905fb..30fb461cb5 100644 --- a/tensilelite/Tensile/Source/lib/source/MLPNet.cpp +++ b/tensilelite/Tensile/Source/lib/source/MLPNet.cpp @@ -26,6 +26,7 @@ #include #include +#include #include #include #include diff --git a/tensilelite/Tensile/Source/lib/source/hip/HipSolutionAdapter.cpp b/tensilelite/Tensile/Source/lib/source/hip/HipSolutionAdapter.cpp index 17e501b921..235bcaee1d 100644 --- a/tensilelite/Tensile/Source/lib/source/hip/HipSolutionAdapter.cpp +++ b/tensilelite/Tensile/Source/lib/source/hip/HipSolutionAdapter.cpp @@ -35,7 +35,7 @@ #include //@TODO add alternative for windows -#ifndef WIN32 +#ifndef _WIN32 #include #endif #include diff --git a/tensilelite/rocisa/rocisa/src/helper.cpp b/tensilelite/rocisa/rocisa/src/helper.cpp index c73de0f29d..797cb68561 100644 --- a/tensilelite/rocisa/rocisa/src/helper.cpp +++ b/tensilelite/rocisa/rocisa/src/helper.cpp @@ -27,7 +27,7 @@ // Windows. The alternatives are switched at a whole-file level. Please do // not use inline/fine-grained ifdefs. -#ifdef WIN32 +#ifdef _WIN32 #include // windows.h must be loaded before other windows headers.