Make queue wait() explicit, rather than every kernel launch.

highperformancecoder · Dec 11, 2024 · 9840057 · 9840057
1 parent 7167d78
commit 9840057
Show file tree

Hide file tree

Showing 3 changed files with 16 additions and 7 deletions.
diff --git a/include/ecolab.h b/include/ecolab.h
@@ -246,7 +246,6 @@ namespace ecolab
           }
         });
       });
-      syclQ().wait();
 #else
       auto sz=this->size();
 #ifdef _OPENMP
@@ -257,6 +256,13 @@ namespace ecolab
         cell.m_idx=idx;
         f(cell);
       }
+#endif
+    }
+
+    /// synchronise GPU threads
+    void syncThreads() {
+#ifdef SYCL_LANGUAGE_VERSION
+      syclQ().wait();
 #endif
     }
   };

diff --git a/models/ecolab_model.cc b/models/ecolab_model.cc
@@ -205,7 +205,7 @@ void SpatialModel::mutate()
   DeviceType<array<unsigned,graphcode::Allocator<unsigned>>> cell_ids;
 #ifdef SYCL_LANGUAGE_VERSION
   cell_ids->allocator(graphcode::Allocator<unsigned>(syclQ(),sycl::usm::alloc::shared));
-  syclQ().wait();
+  syncThreads();
 #endif
 
   // TODO - this is a kind of scan - can it be done on device?
@@ -618,4 +618,5 @@ void SpatialModel::makeConsistent()
       }
   });
   ModelData::makeConsistent(nsp);
+  syncThreads();
 }
diff --git a/models/spatial_ecolab.py b/models/spatial_ecolab.py
@@ -26,8 +26,8 @@ def randomList(num, min, max):
 
 ecolab.species(range(nsp))
 
-numX=32
-numY=32
+numX=64
+numY=64
 ecolab.setGrid(numX,numY)
 ecolab.partitionObjects()
 
@@ -54,15 +54,17 @@ def randomList(num, min, max):
 
 def stepImpl():
     ecolab.generate(100)
-    ecolab.mutate()
-#    ecolab.migrate()
-#    ecolab.condense()
+    #    ecolab.mutate()
+    #    ecolab.migrate()
+    #    ecolab.condense()
+    ecolab.syncThreads()
     ecolab.gather()
 
 from timeit import timeit
 print(timeit('stepImpl()', globals=globals(), number=10))
 
 def step():
+    stepImpl()
     if myid()==0:
         nsp=len(ecolab.species)
         statusBar.configure(text=f't={ecolab.tstep()} nsp:{nsp}')