Merge pull request #1957 from CHIP-SPV/sycl-fixes

sycl: fix several correctness bugs in sycl-ref backend

Author: jeremylt

backends/sycl-ref/ceed-sycl-ref-basis.sycl.cpp

@@ -97,6 +97,10 @@ static int CeedBasisApplyInterp_Sycl(sycl::queue &sycl_queue, const SyclModule_t
         const CeedScalar *cur_u = u + elem * u_stride + comp * u_comp_stride;
         CeedScalar       *cur_v = v + elem * v_stride + comp * v_comp_stride;
 
+        // Prevent race: idle work items (i >= writeLen) must not overwrite
+        // s_buffer_1 while active work items still read it from the previous comp.
+        sycl::group_barrier(work_group);
+
         for (CeedInt k = i; k < u_size; k += group_size) {
           s_buffer_1[k] = cur_u[k];
         }
@@ -105,9 +109,10 @@ static int CeedBasisApplyInterp_Sycl(sycl::queue &sycl_queue, const SyclModule_t
         CeedInt post = 1;
 
         for (CeedInt d = 0; d < dim; d++) {
-          // Use older version of sycl workgroup barrier for performance reasons
-          // Can be updated in future to align with SYCL2020 spec if performance bottleneck is removed
-          // sycl::group_barrier(work_group);
+          // Full work-group barrier with local-only fence: s_buffer_1/2 are SLM
+          // (local_accessor), so local_space is sufficient and avoids the cost of
+          // a global memory fence. Do not replace with a sub-group barrier —
+          // work_group_size (= Q) can exceed the hardware sub-group size.
           work_item.barrier(sycl::access::fence_space::local_space);
 
           pre /= P;
@@ -206,9 +211,10 @@ static int CeedBasisApplyGrad_Sycl(sycl::queue &sycl_queue, const SyclModule_t &
           CeedScalar       *cur_v = v + elem * v_stride + dim_1 * v_dim_stride + comp * v_comp_stride;
 
           for (CeedInt dim_2 = 0; dim_2 < dim; dim_2++) {
-            // Use older version of sycl workgroup barrier for performance reasons
-            // Can be updated in future to align with SYCL2020 spec if performance bottleneck is removed
-            // sycl::group_barrier(work_group);
+            // Full work-group barrier with local-only fence: s_buffer_1/2 are SLM
+            // (local_accessor), so local_space is sufficient and avoids the cost of
+            // a global memory fence. Do not replace with a sub-group barrier —
+            // work_group_size (= Q) can exceed the hardware sub-group size.
             work_item.barrier(sycl::access::fence_space::local_space);
 
             pre /= P;

backends/sycl-ref/ceed-sycl-vector.sycl.cpp

@@ -690,7 +690,6 @@ static int CeedVectorDestroy_Sycl(const CeedVector vec) {
   CeedCallBackend(CeedVectorGetData(vec, &impl));
   CeedCallBackend(CeedGetData(ceed, &data));
 
-  // Wait for all work to finish before freeing memory
   CeedCallSycl(ceed, data->sycl_queue.wait_and_throw());
   CeedCallSycl(ceed, sycl::free(impl->d_array_owned, data->sycl_context));
   CeedCallSycl(ceed, sycl::free(impl->reduction_norm, data->sycl_context));

backends/sycl/ceed-sycl-common.sycl.cpp

@@ -82,6 +82,9 @@ int CeedDestroy_Sycl(Ceed ceed) {
   Ceed_Sycl *data;
 
   CeedCallBackend(CeedGetData(ceed, &data));
+  // CeedCalloc allocates without calling constructors; explicitly run destructors
+  // before freeing so the sycl::queue destructor waits for pending GPU work.
+  data->~Ceed_Sycl();
   CeedCallBackend(CeedFree(&data));
   return CEED_ERROR_SUCCESS;
 }

backends/sycl/ceed-sycl-compile.sycl.cpp

@@ -60,8 +60,26 @@ static int CeedJitAddDefinitions_Sycl(Ceed ceed, const std::string &kernel_sourc
 //------------------------------------------------------------------------------
 // TODO: Add architecture flags, optimization flags
 //------------------------------------------------------------------------------
-static inline int CeedJitGetFlags_Sycl(std::vector<std::string> &flags) {
+static inline int CeedJitGetFlags_Sycl(Ceed ceed, std::vector<std::string> &flags) {
   flags = {std::string("-cl-std=CL3.0"), std::string("-Dint32_t=int"), std::string("-DCEED_RUNNING_JIT_PASS=1")};
+  // Add JIT source root directories as -I include paths
+  {
+    const char **jit_source_dirs;
+    CeedInt      num_jit_source_dirs;
+
+    CeedCallBackend(CeedGetJitSourceRoots(ceed, &num_jit_source_dirs, &jit_source_dirs));
+    for (CeedInt i = 0; i < num_jit_source_dirs; i++) flags.push_back(std::string("-I") + jit_source_dirs[i]);
+    CeedCallBackend(CeedRestoreJitSourceRoots(ceed, &jit_source_dirs));
+  }
+  // Add user JIT defines as -D flags
+  {
+    const char **jit_defines;
+    CeedInt      num_jit_defines;
+
+    CeedCallBackend(CeedGetJitDefines(ceed, &num_jit_defines, &jit_defines));
+    for (CeedInt i = 0; i < num_jit_defines; i++) flags.push_back(std::string("-D") + jit_defines[i]);
+    CeedCallBackend(CeedRestoreJitDefines(ceed, &jit_defines));
+  }
   return CEED_ERROR_SUCCESS;
 }
 
@@ -129,7 +147,7 @@ int CeedBuildModule_Sycl(Ceed ceed, const std::string &kernel_source, SyclModule
 
   CeedCallBackend(CeedGetData(ceed, &data));
   CeedCallBackend(CeedJitAddDefinitions_Sycl(ceed, kernel_source, jit_source, constants));
-  CeedCallBackend(CeedJitGetFlags_Sycl(flags));
+  CeedCallBackend(CeedJitGetFlags_Sycl(ceed, flags));
   CeedCallBackend(CeedJitCompileSource_Sycl(ceed, data->sycl_device, jit_source, il_binary, flags));
   CeedCallBackend(CeedLoadModule_Sycl(ceed, data->sycl_context, data->sycl_device, il_binary, sycl_module));
   return CEED_ERROR_SUCCESS;

examples/ceed/ex1-volume-f-c.h

@@ -9,7 +9,7 @@
 
 /// libCEED Q-function for building quadrature data for a mass operator
 CEED_QFUNCTION(build_mass)(void *ctx, const CeedInt Q, const CeedScalar *const *in, CeedScalar *const *out) {
-  long long int *build_data = (long long int *)ctx;
+  const long *build_data = (const long *)ctx;
 
   // in[0] is Jacobians with shape [dim, dim, Q]
   // in[1] is quadrature weights with shape [1, Q]

interface/ceed.c

@@ -1679,6 +1679,8 @@ int CeedDestroy(Ceed *ceed) {
     CeedCall(CeedFree(&(*ceed)->obj_delegates));
   }
 
+  CeedCall(CeedWorkVectorsDestroy(*ceed));
+
   if ((*ceed)->Destroy) CeedCall((*ceed)->Destroy(*ceed));
 
   for (CeedInt i = 0; i < (*ceed)->num_jit_source_roots; i++) {
@@ -1699,7 +1701,6 @@ int CeedDestroy(Ceed *ceed) {
   CeedCall(CeedFree(&(*ceed)->f_offsets));
   CeedCall(CeedFree(&(*ceed)->resource));
   CeedCall(CeedDestroy(&(*ceed)->op_fallback_ceed));
-  CeedCall(CeedWorkVectorsDestroy(*ceed));
   CeedCall(CeedObjectDestroy_Private(&(*ceed)->obj));
   CeedCall(CeedFree(ceed));
   return CEED_ERROR_SUCCESS;