mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-06-30 11:47:48 +00:00
2e3183af4f
* updating codegen build for MIOpen access: adding .cmake for codegen component * updating CMake * adding in header guards for some headers due to issues with hiprtc compilation in MIOpen * some more header guards * putting env file in header guard * cleaning up some includes * updated types file for hiprtc purposes * fixed types file: bit-wise/memcpy issue * updating multiple utility files to deal with standard header inclusion for hiprtc * added some more header guards in the utility files, replacing some standard header functionality * added some more header guards * fixing some conflicts in utility files, another round of header guards * fixing errors in data type file * resolved conflict errors in a few utility files * added header guards/replicated functionality in device files * resolved issues with standard headers in device files: device_base and device_grouped_conv_fwd_multiple_abd * resolved issues with standard headers in device files: device_base.hpp, device_grouped_conv_fwd_multiple_abd.hpp, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp * added header guards for gridwise gemm files: gridwise_gemm_multiple_abd_xdl_cshuffle.hpp and gridwise_gemm_multiple_d_xdl_cshuffle.hpp * fixed issue with numerics header, removed from transform_conv_fwd_to_gemm and added to device_column_to_image_impl, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3, device_image_to_column_impl * replaced standard header usage and added header guards in block to ctile map and gridwise_gemm_pipeline_selector * resolved errors in device_gemm_xdl_splitk_c_shuffle files in regards to replacement of standard headers in previous commit * added replicated functionality for standard header methods in utility files * replaced standard header functionality in threadwise tensor slice transfer files and added header guards in element_wise_operation.hpp * temp fix for namespace error in MIOpen * remove standard header usage in codegen device op * removed standard header usage in elementwise files, resolved namespace errors * formatting fix * changed codegen argument to ON for testing * temporarily removing codegen compiler flag for testing purposes * added codegen flag again, set default to ON * set codegen flag default back to OFF * replaced enable_if_t standard header usage in data_type.hpp * added some debug prints to pinpoint issues in MIOpen * added print outs to debug in MIOpen * removed debug print outs from device op * resolved stdexcept include error * formatting fix * adding includes to new fp8 file to resolve ck::enable_if_t errors * made changes to amd_wave_read_first_lane * updated functionality in type utility file * fixed end of file issue * resovled errors in type utility file, added functionality to array utility file * fixed standard header usage replication in data_type file, resolves error with failing examples on navi3x * formatting fix * replaced standard header usage in amd_ck_fp8 file * added include to random_gen file * removed and replicated standard header usage from data_type and type_convert files for fp8 changes * replicated standard unsigned integer types in random_gen * resolved comments from review: put calls to reinterpret_cast for size_t in header guards * updated/added copyright headers * removed duplicate header * fixed typo in header guard * updated copyright headers --------- Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
107 lines
3.2 KiB
C++
107 lines
3.2 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#include <rtc/hip.hpp>
|
|
#include <rtc/compile_kernel.hpp>
|
|
#include <rtc/tmp_dir.hpp>
|
|
#include <stdexcept>
|
|
#include <iostream>
|
|
#include <fstream>
|
|
#include <cassert>
|
|
|
|
namespace rtc {
|
|
|
|
template <class T>
|
|
T generic_read_file(const std::string& filename, size_t offset = 0, size_t nbytes = 0)
|
|
{
|
|
std::ifstream is(filename, std::ios::binary | std::ios::ate);
|
|
if(nbytes == 0)
|
|
{
|
|
// if there is a non-zero offset and nbytes is not set,
|
|
// calculate size of remaining bytes to read
|
|
nbytes = is.tellg();
|
|
if(offset > nbytes)
|
|
throw std::runtime_error("offset is larger than file size");
|
|
nbytes -= offset;
|
|
}
|
|
if(nbytes < 1)
|
|
throw std::runtime_error("Invalid size for: " + filename);
|
|
is.seekg(offset, std::ios::beg);
|
|
|
|
T buffer(nbytes, 0);
|
|
if(not is.read(&buffer[0], nbytes))
|
|
throw std::runtime_error("Error reading file: " + filename);
|
|
return buffer;
|
|
}
|
|
|
|
std::vector<char> read_buffer(const std::string& filename, size_t offset = 0, size_t nbytes = 0)
|
|
{
|
|
return generic_read_file<std::vector<char>>(filename, offset, nbytes);
|
|
}
|
|
|
|
std::string read_string(const std::string& filename)
|
|
{
|
|
return generic_read_file<std::string>(filename);
|
|
}
|
|
|
|
void write_buffer(const std::string& filename, const char* buffer, std::size_t size)
|
|
{
|
|
std::ofstream os(filename);
|
|
os.write(buffer, size);
|
|
}
|
|
void write_buffer(const std::string& filename, const std::vector<char>& buffer)
|
|
{
|
|
write_buffer(filename, buffer.data(), buffer.size());
|
|
}
|
|
void write_string(const std::string& filename, const std::string_view& buffer)
|
|
{
|
|
write_buffer(filename, buffer.data(), buffer.size());
|
|
}
|
|
|
|
std::string compiler() { return "/opt/rocm/llvm/bin/clang++ -x hip --cuda-device-only"; }
|
|
// TODO: undo after extracting the codeobj
|
|
// std::string compiler() { return "/opt/rocm/llvm/bin/clang++ -x hip"; }
|
|
|
|
kernel compile_kernel(const std::vector<src_file>& srcs, compile_options options)
|
|
{
|
|
assert(not srcs.empty());
|
|
tmp_dir td{"compile"};
|
|
options.flags += " -I. -O3";
|
|
options.flags += " -std=c++17";
|
|
options.flags += " --offload-arch=" + get_device_name();
|
|
std::string out;
|
|
|
|
for(const auto& src : srcs)
|
|
{
|
|
fs::path full_path = td.path / src.path;
|
|
fs::path parent_path = full_path.parent_path();
|
|
fs::create_directories(parent_path);
|
|
write_string(full_path.string(), src.content);
|
|
if(src.path.extension().string() == ".cpp")
|
|
{
|
|
options.flags += " -c " + src.path.filename().string();
|
|
if(out.empty())
|
|
out = src.path.stem().string() + ".o";
|
|
}
|
|
}
|
|
|
|
options.flags += " -o " + out;
|
|
td.execute(compiler() + options.flags);
|
|
|
|
auto out_path = td.path / out;
|
|
if(not fs::exists(out_path))
|
|
throw std::runtime_error("Output file missing: " + out);
|
|
|
|
auto obj = read_buffer(out_path.string());
|
|
|
|
std::ofstream ofh("obj.o", std::ios::binary);
|
|
for(auto i : obj)
|
|
ofh << i;
|
|
ofh.close();
|
|
// int s = std::system(("/usr/bin/cp " + out_path.string() + " codeobj.bin").c_str());
|
|
// assert(s == 0);
|
|
return kernel{obj.data(), options.kernel_name};
|
|
}
|
|
|
|
} // namespace rtc
|