Add additional src files

src/targets/gpu/compile_hip_code_object.cpp

@@ -161,7 +161,7 @@ operation compile_hip_code_object(const std::string& content, hip_compile_option
     assert(not options.inputs.empty());
     assert(options.inputs.size() == options.virtual_inputs.size() or
            options.virtual_inputs.empty());
-    std::vector<src_file> srcs;
+    std::vector<src_file> srcs = options.additional_src_files;
     std::transform(migraphx_kernels().begin(),
                    migraphx_kernels().end(),
                    std::back_inserter(srcs),

src/targets/gpu/include/migraphx/gpu/compile_hip_code_object.hpp

@@ -29,6 +29,9 @@
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
+
+struct src_file;
+
 namespace gpu {
 
 struct context;
@@ -42,6 +45,7 @@ struct hip_compile_options
     std::string kernel_name           = "kernel";
     std::string params                = "";
     std::vector<shape> virtual_inputs = {};
+    std::vector<src_file> additional_src_files = {};
 
     /**
      * @brief Set the launch parameters but allow v to override the values

src/targets/gpu/jit/ck_gemm.cpp

@@ -80,15 +80,6 @@ __global__ void ${kernel}(${params})
 
 static bool transposed_matrix(const shape& s) { return s.strides().back() != 1; }
 
-template <class F, class Action>
-auto action_decorate(F f, Action action)
-{
-    return [=](auto&&... xs) {
-        action();
-        f(std::forward<decltype(xs)>(xs)...);
-    };
-}
-
 using tuning_entry = std::pair<std::vector<shape>, size_t>;
 static std::vector<tuning_entry> read_tuning(const std::string& s)
 {
@@ -300,7 +291,9 @@ struct ck_gemm_compiler : compiler<ck_gemm_compiler>
         const auto block_size       = solution.block_size;
 
         hip_compile_options options;
-        options.embedded_headers = ck_headers;
+        std::transform(ck_headers.begin(), ck_headers.end(), std::back_inserter(options.additional_src_files), [&](auto&& p) {
+            return src_file{fs::path{p.first}, p.second};
+        });
         auto grid_size = can_fold_batch ? blocks_per_batch : batch_count * blocks_per_batch;
         options.set_launch_params(v, grid_size * block_size, block_size);
         options.inputs         = inputs;