Add hip compilation (#664)

* Add compiler flags

* Add missing include

* Add filesystem header

* Formatting

* Add tmp_dir to run

* Formatting

* Kernel compilation and launching

* Formatting

* Seperate pack_args

* Formatting

* Add alignment tests

* Formatting

* Add compile test

* Formatting

* Complete compile test

* Formatting

* Use is_regular_file free function

* Fix is_regular_file call

* Fix tidy issues

* Fix tidy

* Fix tidy issue

* Print size in read_buffer to debug issue on jenkins

* Add hip flags before src file

* Fix reading output files

* Fix unsued variable warning

* Formatting

* Formatting

* Disable tidy check
Co-authored-by: Shucai Xiao <shucai.xiao@amd.com>
Co-authored-by: mvermeulen <5479696+mvermeulen@users.noreply.github.com>

src/tmp_dir.cpp

+#include <migraphx/tmp_dir.hpp>
+#include <migraphx/env.hpp>
+#include <migraphx/errors.hpp>
+#include <algorithm>
+#include <random>
+#include <thread>
+#include <sstream>
+#include <iostream>
+#include <string>
+#include <sys/types.h>
+#include <unistd.h>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DEBUG_SAVE_TEMP_DIR)
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_TRACE_CMD_EXECUTE)
+
+std::string random_string(std::string::size_type length)
+{
+    static const std::string& chars = "0123456789"
+                                      "abcdefghijklmnopqrstuvwxyz"
+                                      "ABCDEFGHIJKLMNOPQRSTUVWXYZ";
+
+    std::mt19937 rg{std::random_device{}()};
+    std::uniform_int_distribution<std::string::size_type> pick(0, chars.length() - 1);
+
+    std::string str(length, 0);
+    std::generate(str.begin(), str.end(), [&] { return chars[pick(rg)]; });
+
+    return str;
+}
+
+std::string unique_string(const std::string& prefix)
+{
+    auto pid = getpid();
+    auto tid = std::this_thread::get_id();
+    std::stringstream ss;
+    ss << prefix << "-" << pid << "-" << tid << "-" << random_string(64);
+    return ss.str();
+}
+
+tmp_dir::tmp_dir() : path(fs::temp_directory_path() / unique_string("migraphx"))
+{
+    fs::create_directories(this->path);
+}
+
+void system_cmd(const std::string& cmd)
+{
+// We shouldn't call system commands
+#ifdef MIGRAPHX_USE_CLANG_TIDY
+    (void)cmd;
+#else
+    if(std::system(cmd.c_str()) != 0)
+        MIGRAPHX_THROW("Can't execute " + cmd);
+#endif
+}
+
+void tmp_dir::execute(const std::string& exe, const std::string& args) const
+{
+    std::string cd  = "cd " + this->path.string() + "; ";
+    std::string cmd = cd + exe + " " + args; // + " > /dev/null";
+    if(enabled(MIGRAPHX_TRACE_CMD_EXECUTE{}))
+        std::cout << cmd << std::endl;
+    system_cmd(cmd);
+}
+
+tmp_dir::~tmp_dir()
+{
+    if(!enabled(MIGRAPHX_DEBUG_SAVE_TEMP_DIR{}))
+    {
+        fs::remove_all(this->path);
+    }
+}
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

test/gpu/jit.cpp

+#include <test.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/gpu/kernel.hpp>
+#include <migraphx/gpu/hip.hpp>
+#include <migraphx/gpu/compile_hip.hpp>
+
+// NOLINTNEXTLINE
+const std::string write_2s = R"migraphx(
+#include <hip/hip_runtime.h>
+
+extern "C" {
+__global__ void write(int* data) 
+{
+    int num = threadIdx.x + blockDim.x * blockIdx.x;
+    data[num] = 2;
+}
+    
+}
+
+int main() {}
+
+)migraphx";
+
+migraphx::gpu::src_file make_src_file(const std::string& name, const std::string& content)
+{
+    return {name, std::make_pair(content.data(), content.data() + content.size())};
+}
+
+std::string get_device_name()
+{
+    hipDeviceProp_t props{};
+    int device;
+    EXPECT(hipGetDevice(&device) == hipSuccess);
+    EXPECT(hipGetDeviceProperties(&props, device) == hipSuccess);
+    return "gfx" + std::to_string(props.gcnArch);
+}
+
+TEST_CASE(simple_compile_hip)
+{
+    auto binaries = migraphx::gpu::compile_hip_src(
+        {make_src_file("main.cpp", write_2s)}, "", get_device_name());
+    EXPECT(binaries.size() == 1);
+
+    migraphx::argument input{{migraphx::shape::int32_type, {5}}};
+    auto ginput = migraphx::gpu::to_gpu(input);
+    migraphx::gpu::kernel k{binaries.front(), "write"};
+    k.launch(nullptr, input.get_shape().elements(), 1024)(ginput.cast<int>());
+    auto output = migraphx::gpu::from_gpu(ginput);
+
+    EXPECT(output != input);
+    auto data = output.get<int>();
+    EXPECT(migraphx::all_of(data, [](auto x) { return x == 2; }));
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/gpu/pack_args.cpp

+#include <test.hpp>
+#include <migraphx/gpu/pack_args.hpp>
+
+template <class T>
+std::size_t packed_sizes()
+{
+    return sizeof(T);
+}
+
+template <class T, class U, class... Ts>
+std::size_t packed_sizes()
+{
+    return sizeof(T) + packed_sizes<U, Ts...>();
+}
+
+template <class... Ts>
+std::size_t sizes()
+{
+    return migraphx::gpu::pack_args({Ts{}...}).size();
+}
+
+template <class... Ts>
+std::size_t padding()
+{
+    EXPECT(sizes<Ts...>() >= packed_sizes<Ts...>());
+    return sizes<Ts...>() - packed_sizes<Ts...>();
+}
+
+struct float_struct
+{
+    float x, y;
+};
+
+TEST_CASE(alignment_padding)
+{
+    EXPECT(padding<short, short>() == 0);
+    EXPECT(padding<float, float_struct>() == 0);
+    EXPECT(padding<short, float_struct>() == 2);
+    EXPECT(padding<short, int>() == 2);
+    EXPECT(padding<char, short, int, char>() == 1);
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }