[dl-cifar][SYCL] Add CMake option to use in-order queue

Add the queue properties in_order and, if available, discard_events.
The addition is steered by a CMake build option IN_ORDER_QUEUE.
Set the default value to ON for NVIDIA and AMD backends and keep
as OFF for other backends.

Author: rafbiels

dl-cifar/README.md

@@ -96,6 +96,11 @@ CC=clang CXX=clang++ **cmake** -DUSE_AMD_BACKEND=YES -DUSE_AMD_ARCH=gfx90a .. \
 **ONEAPI_DEVICE_SELECTOR=hip:gpu ./dl-cifar_sycl** 
 
 ---------------------------------------------------------------------------------------------------------
+## In-order queue
+The CMake option `-DIN_ORDER_QUEUE` adds the `in_order` property to the SYCL
+queue, as well as `discard_events` if available. The default value of this
+option is `ON` for NVIDIA and AMD backends, and `OFF` otherwise.
+
 ## Workload logging/tracing
 
 **DL-CIFAR provides function tracing:**

dl-cifar/SYCL/CMakeLists.txt

@@ -34,6 +34,12 @@ option(USE_NVIDIA_BACKEND      "Build for NVIDIA backend"     OFF)
 option(USE_AMDHIP_BACKEND      "Build for AMD HIP backend"    OFF)
 option(DEVICE_TIMER            "Build using Device Timer" OFF)
 
+set(IN_ORDER_QUEUE_DEFAULT OFF)
+if (${USE_NVIDIA_BACKEND} OR ${USE_AMDHIP_BACKEND})
+    set(IN_ORDER_QUEUE_DEFAULT ON)
+endif()
+option(IN_ORDER_QUEUE "Use in-order SYCL queue" ${IN_ORDER_QUEUE_DEFAULT})
+
 set(DEF_INTEL_WL_CXX_FLAGS  " -DMKL_ILP64 ")
 set(DEF_NVIDIA_WL_CXX_FLAGS " -DUSE_CUBLAS ")
 set(DEF_AMD_WL_CXX_FLAGS    " -DUSE_ROCBLAS -D__HIP_PLATFORM_AMD__ ")
@@ -42,6 +48,11 @@ set(DEF_INTEL_GENERAL_CXX_FLAGS  " -O3 -fsycl -ffast-math ")
 set(DEF_NVIDIA_GENERAL_CXX_FLAGS " -O3 -fsycl -ffast-math ")
 set(DEF_AMD_GENERAL_CXX_FLAGS    " -O3 -fsycl -ffast-math ")
 
+if (${IN_ORDER_QUEUE})
+    string(APPEND DEF_INTEL_GENERAL_CXX_FLAGS " -DIN_ORDER_QUEUE ")
+    string(APPEND DEF_NVIDIA_GENERAL_CXX_FLAGS " -DIN_ORDER_QUEUE ")
+    string(APPEND DEF_AMD_GENERAL_CXX_FLAGS " -DIN_ORDER_QUEUE ")
+endif()
 
 # -DCMAKE_CXX_FLAGS=" -blah -blah " overrides the default flags (BOTH general and WL specific)
 # -DOVERRIDE_GENERAL_CXX_FLAGS=" -blah -blah " overrides the general flags only (and not the workload specific flags)

dl-cifar/SYCL/basic-dl/lnorm_layer.cpp

@@ -80,7 +80,8 @@ void LNormLayer::doFw() {
                 }
             }
         });
-    }).wait();
+    });
+    langHandle_->getSyclQueue()->wait();
 
 
     Tracer::func_end("LNormLayer::doFw");   
@@ -157,7 +158,8 @@ void LNormLayer::doBw() {
 
             }
         });
-    }).wait();    
+    });
+    langHandle_->getSyclQueue()->wait();
 
     Tracer::func_end("LNormLayer::doBw");   
 }

dl-cifar/SYCL/basic-dl/lnorm_layer.h

@@ -88,7 +88,8 @@ class LNormLayerController {
             float *d_input, *d_d_input;
             d_input   = (float *)sycl::malloc_device(inputSize*sizeof(float),   sycl_queue);
             d_d_input   = (float *)sycl::malloc_device(inputSize*sizeof(float),   sycl_queue);
-            sycl_queue.memcpy(d_input, h_input, sizeof(float) * inputSize).wait();
+            sycl_queue.memcpy(d_input, h_input, sizeof(float) * inputSize);
+            sycl_queue.wait();
 
             int outputSize   = inputSize;
             float *h_d_output   = (float*)calloc(outputSize,   sizeof(float));  
@@ -104,12 +105,14 @@ class LNormLayerController {
             for(int i=0; i<iterCount; i++) {
                 // for some reason the compiler is not liking calls to ImageProcessor::initImage() from here
                 //ImageProcessor::initImage(h_input, inputSize);
-                sycl_queue.memcpy(d_input, h_input, sizeof(float) * inputSize).wait();
+                sycl_queue.memcpy(d_input, h_input, sizeof(float) * inputSize);
+                sycl_queue.wait();
                 lNormLayer->doFw();
 
                 // for some reason the compiler is not liking calls to ImageProcessor::initImage() from here
                 //ImageProcessor::initImage(h_d_output, outputSize);
-                sycl_queue.memcpy(d_d_output, h_d_output, sizeof(float) * outputSize).wait();
+                sycl_queue.memcpy(d_d_output, h_d_output, sizeof(float) * outputSize);
+                sycl_queue.wait();
                 lNormLayer->doBw();
             }
 

dl-cifar/SYCL/handle.h

@@ -104,11 +104,18 @@ class LangHandle {
 #ifdef DEVICE_TIMER  
             timer->recordOpTimeTaken(1000, calculate_op_time_taken(start), "CREATE_SYCL_CONTEXT");
 #endif    
-            //auto propList = sycl::property_list{sycl::property::queue::in_order()};    
+            auto propList = sycl::property_list{
+                #ifdef IN_ORDER_QUEUE
+                sycl::property::queue::in_order{},
+                #ifdef SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS
+                sycl::ext::oneapi::property::queue::discard_events{}
+                #endif
+                #endif
+            };
 #ifdef DEVICE_TIMER  
             start = get_time_now();  
 #endif     
-            sycl_queue_ = new sycl::queue(*context_, *dht_);
+            sycl_queue_ = new sycl::queue(*context_, *dht_, propList);
 #ifdef DEVICE_TIMER  
             timer->recordOpTimeTaken(1000, calculate_op_time_taken(start), "CREATE_SYCL_QUEUE");
 #endif    
@@ -246,11 +253,18 @@ class LangHandle {
 #ifdef DEVICE_TIMER  
             timer->recordOpTimeTaken(1000, calculate_op_time_taken(start), "CREATE_SYCL_CONTEXT");
 #endif    
-            //auto propList = sycl::property_list{sycl::property::queue::in_order()};    
+            auto propList = sycl::property_list{
+                #ifdef IN_ORDER_QUEUE
+                sycl::property::queue::in_order{},
+                #ifdef SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS
+                sycl::ext::oneapi::property::queue::discard_events{}
+                #endif
+                #endif
+            };
 #ifdef DEVICE_TIMER  
             start = get_time_now();  
 #endif     
-            sycl_queue_ = new sycl::queue(*context_, *dht_);
+            sycl_queue_ = new sycl::queue(*context_, *dht_, propList);
 #ifdef DEVICE_TIMER  
             timer->recordOpTimeTaken(1000, calculate_op_time_taken(start), "CREATE_SYCL_QUEUE");
 #endif    
@@ -343,11 +357,18 @@ class LangHandle {
 #ifdef DEVICE_TIMER  
             timer->recordOpTimeTaken(1000, calculate_op_time_taken(start), "CREATE_SYCL_CONTEXT");
 #endif    
-            //auto propList = sycl::property_list{sycl::property::queue::in_order()};    
+            auto propList = sycl::property_list{
+                #ifdef IN_ORDER_QUEUE
+                sycl::property::queue::in_order{},
+                #ifdef SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS
+                sycl::ext::oneapi::property::queue::discard_events{}
+                #endif
+                #endif
+            };
 #ifdef DEVICE_TIMER  
             start = get_time_now();  
 #endif     
-            sycl_queue_ = new sycl::queue(*context_, *dht_);
+            sycl_queue_ = new sycl::queue(*context_, *dht_, propList);
 #ifdef DEVICE_TIMER  
             timer->recordOpTimeTaken(1000, calculate_op_time_taken(start), "CREATE_SYCL_QUEUE");
 #endif    

dl-cifar/SYCL/upsample.cpp

@@ -78,7 +78,8 @@ void Upsampler::upsample(LangHandle *langHandle, float *d_src, float *d_dst, int
                 d_dst[dst_pixCntUntilSegment + relPixIdxInSegment] = d_src[src_imgIdx*src_noOfPixelsPerImg + (pixelIY*srcWidth) + pixelIX];
             }
         });
-    }).wait();
+    });
+    langHandle->getSyclQueue()->wait();
 
     Tracer::func_end("Upsampler::upsample");   
 }