43   ostringstream code;  in CeedQFunctionBuildKernel_Cuda_ref()  local
46   code << "// QFunction source\n";  in CeedQFunctionBuildKernel_Cuda_ref()
47   code << "#include <ceed/jit-source/cuda/cuda-ref-qfunction.h>\n\n";  in CeedQFunctionBuildKernel_Cuda_ref()
54     code << "// User QFunction source\n";  in CeedQFunctionBuildKernel_Cuda_ref()
55     code << "#include \"" << source_path << "\"\n\n";  in CeedQFunctionBuildKernel_Cuda_ref()
57 …code << "extern \"C\" __global__ void " << kernel_name << "(void *ctx, CeedInt Q, Fields_Cuda fiel…  in CeedQFunctionBuildKernel_Cuda_ref()
60   code << "  // Input fields\n";  in CeedQFunctionBuildKernel_Cuda_ref()
63     code << "  const CeedInt size_input_" << i << " = " << size << ";\n";  in CeedQFunctionBuildKernel_Cuda_ref()
64     code << "  CeedScalar input_" << i << "[size_input_" << i << "];\n";  in CeedQFunctionBuildKernel_Cuda_ref()
66   code << "  const CeedScalar *inputs[" << CeedIntMax(num_input_fields, 1) << "];\n";  in CeedQFunctionBuildKernel_Cuda_ref()
68     code << "  inputs[" << i << "] = input_" << i << ";\n";  in CeedQFunctionBuildKernel_Cuda_ref()
70   code << "\n";  in CeedQFunctionBuildKernel_Cuda_ref()
73   code << "  // Output fields\n";  in CeedQFunctionBuildKernel_Cuda_ref()
76     code << "  const CeedInt size_output_" << i << " = " << size << ";\n";  in CeedQFunctionBuildKernel_Cuda_ref()
77     code << "  CeedScalar output_" << i << "[size_output_" << i << "];\n";  in CeedQFunctionBuildKernel_Cuda_ref()
79   code << "  CeedScalar *outputs[" << CeedIntMax(num_output_fields, 1) << "];\n";  in CeedQFunctionBuildKernel_Cuda_ref()
81     code << "  outputs[" << i << "] = output_" << i << ";\n";  in CeedQFunctionBuildKernel_Cuda_ref()
83   code << "\n";  in CeedQFunctionBuildKernel_Cuda_ref()
86   code << "  // Loop over quadrature points\n";  in CeedQFunctionBuildKernel_Cuda_ref()
87 …code << "  for (CeedInt q = blockIdx.x * blockDim.x + threadIdx.x; q < Q; q += blockDim.x * gridDi…  in CeedQFunctionBuildKernel_Cuda_ref()
90   code << "    // -- Load inputs\n";  in CeedQFunctionBuildKernel_Cuda_ref()
92 …code << "    readQuads<size_input_" << i << ">(q, Q, fields.inputs[" << i << "], input_" << i << "…  in CeedQFunctionBuildKernel_Cuda_ref()
94   code << "\n";  in CeedQFunctionBuildKernel_Cuda_ref()
97   code << "    // -- Call QFunction\n";  in CeedQFunctionBuildKernel_Cuda_ref()
98   code << "    " << qfunction_name << "(ctx, 1, inputs, outputs);\n\n";  in CeedQFunctionBuildKernel_Cuda_ref()
101   code << "    // -- Write outputs\n";  in CeedQFunctionBuildKernel_Cuda_ref()
103 …code << "    writeQuads<size_output_" << i << ">(q, Q, output_" << i << ", fields.outputs[" << i <…  in CeedQFunctionBuildKernel_Cuda_ref()
105   code << "  }\n";  in CeedQFunctionBuildKernel_Cuda_ref()
106   code << "}\n";  in CeedQFunctionBuildKernel_Cuda_ref()
109   CeedCallBackend(CeedCompile_Cuda(ceed, code.str().c_str(), &data->module, 0));  in CeedQFunctionBuildKernel_Cuda_ref()