Sort based on weights

src/schedule.cpp

@@ -218,9 +218,13 @@ void schedule::apply(program& p) const
 
     // Topo sort
     fix([&](auto self, auto ins) {
-        for(auto i : ins->inputs())
+        auto args = ins->inputs();
+        std::sort(args.begin(), args.end(), [&](auto x, auto y) {
+            return si.weights[x] < si.weights[y];
+        });
+        for(auto i : args)
             p.move_instruction(i, p.begin());
-        for(auto i : ins->inputs())
+        for(auto i : args)
             self(i);
     })(last);
 

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

@@ -36,7 +36,7 @@ struct hip_device
         static hip_stream_ptr create_stream()
         {
             hipStream_t result = nullptr;
-            auto status        = hipStreamCreate(&result);
+            auto status        = hipStreamCreateWithFlags(&result, hipStreamNonBlocking);
             if(status != hipSuccess)
                 MIGRAPHX_THROW("Failed to allocate stream");
             return hip_stream_ptr{result};

src/targets/gpu/schedule_model.cpp

@@ -13,7 +13,8 @@ using hip_event_ptr = MIGRAPHX_MANAGE_PTR(hipEvent_t, hipEventDestroy);
 hip_event_ptr create_event()
 {
     hipEvent_t event;
-    auto status = hipEventCreateWithFlags(&event, hipEventDisableTiming);
+    // Default is hipEventReleaseToDevice
+    auto status = hipEventCreateWithFlags(&event, hipEventDisableTiming | hipEventReleaseToSystem | hipEventBlockingSync);
     if(status != hipSuccess)
         MIGRAPHX_THROW("Failed to create event");
     return hip_event_ptr{event};