87   std::ostringstream code;  in CeedQFunctionBuildKernel_Sycl()  local
88   code << rw_source_view;  in CeedQFunctionBuildKernel_Sycl()
89   code << qf_source_view;  in CeedQFunctionBuildKernel_Sycl()
90   code << "\n";  in CeedQFunctionBuildKernel_Sycl()
95 …code << "__attribute__((intel_reqd_sub_group_size(" << SUB_GROUP_SIZE_QF << "))) __kernel void " <…  in CeedQFunctionBuildKernel_Sycl()
102     code << "  "  in CeedQFunctionBuildKernel_Sycl()
107   code << "  "  in CeedQFunctionBuildKernel_Sycl()
110     code << "\n,  "  in CeedQFunctionBuildKernel_Sycl()
114   code << ") {\n\n";  in CeedQFunctionBuildKernel_Sycl()
117   code << "  // Input fields\n";  in CeedQFunctionBuildKernel_Sycl()
119     code << "  CeedScalar U_" << i << "[" << input_sizes[i] << "];\n";  in CeedQFunctionBuildKernel_Sycl()
121   code << "  const CeedScalar *inputs[" << CeedIntMax(num_input_fields, 1) << "] = {U_0";  in CeedQFunctionBuildKernel_Sycl()
123     code << ", U_" << i << "\n";  in CeedQFunctionBuildKernel_Sycl()
125   code << "};\n\n";  in CeedQFunctionBuildKernel_Sycl()
128   code << "  // Output fields\n";  in CeedQFunctionBuildKernel_Sycl()
130     code << "  CeedScalar V_" << i << "[" << output_sizes[i] << "];\n";  in CeedQFunctionBuildKernel_Sycl()
132   code << "  CeedScalar *outputs[" << CeedIntMax(num_output_fields, 1) << "] = {V_0";  in CeedQFunctionBuildKernel_Sycl()
134     code << ", V_" << i << "\n";  in CeedQFunctionBuildKernel_Sycl()
136   code << "};\n\n";  in CeedQFunctionBuildKernel_Sycl()
138   code << "  const CeedInt q = get_global_linear_id();\n\n";  in CeedQFunctionBuildKernel_Sycl()
140   code << "if(q < Q){ \n\n";  in CeedQFunctionBuildKernel_Sycl()
143   code << "  // -- Load inputs\n";  in CeedQFunctionBuildKernel_Sycl()
145     code << "  readQuads(" << input_sizes[i] << ", Q, q, "  in CeedQFunctionBuildKernel_Sycl()
148   code << "\n";  in CeedQFunctionBuildKernel_Sycl()
151   code << "  // -- Call QFunction\n";  in CeedQFunctionBuildKernel_Sycl()
152   code << "  " << qf_name_view << "(ctx, 1, inputs, outputs);\n\n";  in CeedQFunctionBuildKernel_Sycl()
155   code << "  // -- Write outputs\n";  in CeedQFunctionBuildKernel_Sycl()
157     code << "  writeQuads(" << output_sizes[i] << ", Q, q, "  in CeedQFunctionBuildKernel_Sycl()
160   code << "\n";  in CeedQFunctionBuildKernel_Sycl()
163   code << "}\n";  in CeedQFunctionBuildKernel_Sycl()
164   code << "}\n";  in CeedQFunctionBuildKernel_Sycl()
168   CeedDebug(ceed, code.str().c_str());  in CeedQFunctionBuildKernel_Sycl()
171   CeedCallBackend(CeedBuildModule_Sycl(ceed, code.str(), &impl->sycl_module));  in CeedQFunctionBuildKernel_Sycl()