Merge

src/include/migraphx/matcher.hpp

@@ -754,10 +754,16 @@ auto skip_broadcasts(Ms... ms)
     return skip(name("broadcast", "multibroadcast", "contiguous"))(ms...);
 }
 
+template <class... Ms>
+auto skip_broadcasts_converts(Ms... ms)
+{
+    return skip(name("broadcast", "multibroadcast", "contiguous", "convert"))(ms...);
+}
+
 template <class T>
 inline auto has_value(T x, float tolerance = 1e-6)
 {
-    return skip_broadcasts(make_basic_pred_matcher([=](instruction_ref ins) {
+    return skip_broadcasts_converts(make_basic_pred_matcher([=](instruction_ref ins) {
         if(ins->name() != "@literal")
             return false;
         auto l = ins->get_literal();

src/targets/gpu/jit/pointwise.cpp

@@ -30,7 +30,7 @@ namespace migraphx {
 ${preamble}
 
 extern "C" {
-__global__ void kernel(${params}) 
+__global__ void ${kernel}(${params}) 
 {
     auto idx = make_index();
     pointwise(idx, ${transformers})(${lambda}, ${args});
@@ -42,6 +42,18 @@ __global__ void kernel(${params})
 
 )__migraphx__";
 
+static std::vector<std::string> get_op_names(const module& m)
+{
+    std::vector<std::string> result;
+    for(auto& ins : m)
+    {
+        if(starts_with(ins.name(), "@"))
+            continue;
+        result.push_back(ins.name());
+    }
+    return result;
+}
+
 struct pointwise_compiler : compiler<pointwise_compiler>
 {
     std::vector<std::string> names() const { return {"pointwise"}; }
@@ -69,7 +81,8 @@ struct pointwise_compiler : compiler<pointwise_compiler>
                                options.output.elements() / vec.size,
                                oversubscribe_if(not preloads.is_preloading())));
         auto src = interpolate_string(pointwise_kernel,
-                                      {{"params", enum_params(inputs.size(), "void * private_p")},
+                                      {{"kernel", options.kernel_name},
+                                       {"params", enum_params(inputs.size(), "void * private_p")},
                                        {"args", enum_params(inputs.size(), "private_p")},
                                        {"lambda", v.at("lambda").to<std::string>()},
                                        {"transformers", make_transformer_args(preloads, vec)},
@@ -98,8 +111,13 @@ struct pointwise_compiler : compiler<pointwise_compiler>
         auto name = g.create_function(
             g.generate_module(*pm).set_attributes({"__device__"}).set_generic_types(*pm));
         std::string lambda = "MIGRAPHX_LIFT(" + name + ")";
+        auto op_names      = get_op_names(*pm);
+        op_names.push_back("kernel");
+        auto op_name_string = join_strings(op_names, "_");
         return replace(
-            compile_op(ctx, to_shapes(ins->inputs()), {{"lambda", lambda}, {"preamble", g.str()}}));
+            compile_op(ctx,
+                       to_shapes(ins->inputs()),
+                       {{"lambda", lambda}, {"preamble", g.str()}, {"kernel", op_name_string}}));
     }
 };
 } // namespace gpu