staging

src/opt/memory_coloring_impl.cpp

@@ -136,7 +136,9 @@ void memory_coloring_impl::build()
                     {
                         // input memory is used as this instruction's output.
                         // def is considered as use. Coalesce the live intervals.
-                        DEBUG(assert(def_interval != nullptr));
+#ifndef NDEBUG
+                        assert(def_interval != nullptr);
+#endif
                         def_interval->add_use(cur_points);
                         instr2_live[p_arg] = def_interval;
                     }
@@ -158,7 +160,9 @@ void memory_coloring_impl::build()
                     {
                         interval_ptr interval = instr2_live[p_arg];
                         interval->add_use(cur_points);
-                        DEBUG(assert(live_set.find(interval->id) != live_set.end()));
+#ifndef NDEBUG
+                        assert(live_set.find(interval->id) != live_set.end());
+#endif
                     }
                 }
             }
@@ -186,8 +190,9 @@ void memory_coloring_impl::rewrite()
             interval_ptr interval = instr2_live[p_iter];
             if(interval->get_offset() == InvalidOffset)
             {
-                DEBUG(assert((interval->get_begin() == InvalidOffset) ||
-                             interval->result.bytes() == 0));
+#ifndef NDEBUG
+                assert((interval->get_begin() == InvalidOffset) || interval->result.bytes() == 0);
+#endif
                 continue;
             }
             std::size_t offset = interval->get_offset();

src/opt/memory_coloring_impl.hpp

@@ -94,14 +94,13 @@ struct memory_coloring_impl
         }
         return last_allocate;
     }
+#ifdef DEBUG_OPT    
     bool is_disjoin(live_range& range1, live_range& range2)
     {
         long long end1 = range1.offset + range1.size - 1;
         long long end2 = range2.offset + range2.size - 1;
         return ((end1 < range2.offset) || (end2 < range1.offset));
     }
-
-#ifdef DEBUG_OPT
     void dump(const std::string);
     void dump_program();
     void dump_intervals();

src/targets/gpu/lowering.cpp

@@ -107,7 +107,9 @@ struct miopen_convolution
         workspace_shape = shape{shape::int8_type, {workspace_size}};
 
         auto x         = to_gpu(generate_argument(inputs[0]->get_shape()));
+        gpu_sync();
         auto w         = to_gpu(generate_argument(inputs[1]->get_shape()));
+        gpu_sync();
         auto y         = to_gpu(generate_argument(output_shape));
         auto workspace = allocate_gpu(workspace_shape);