46   ostringstream code;  in CeedQFunctionBuildKernel_Hip_ref()  local
49   code << "// QFunction source\n";  in CeedQFunctionBuildKernel_Hip_ref()
50   code << "#include <ceed/jit-source/hip/hip-ref-qfunction.h>\n\n";  in CeedQFunctionBuildKernel_Hip_ref()
57     code << "// User QFunction source\n";  in CeedQFunctionBuildKernel_Hip_ref()
58     code << "#include \"" << source_path << "\"\n\n";  in CeedQFunctionBuildKernel_Hip_ref()
60   code << "extern \"C\" __launch_bounds__(BLOCK_SIZE)\n";  in CeedQFunctionBuildKernel_Hip_ref()
61   code << "__global__ void " << kernel_name << "(void *ctx, CeedInt Q, Fields_Hip fields) {\n";  in CeedQFunctionBuildKernel_Hip_ref()
64   code << "  // Input fields\n";  in CeedQFunctionBuildKernel_Hip_ref()
67     code << "  const CeedInt size_input_" << i << " = " << size << ";\n";  in CeedQFunctionBuildKernel_Hip_ref()
68     code << "  CeedScalar input_" << i << "[size_input_" << i << "];\n";  in CeedQFunctionBuildKernel_Hip_ref()
70   code << "  const CeedScalar *inputs[" << CeedIntMax(num_input_fields, 1) << "];\n";  in CeedQFunctionBuildKernel_Hip_ref()
72     code << "  inputs[" << i << "] = input_" << i << ";\n";  in CeedQFunctionBuildKernel_Hip_ref()
74   code << "\n";  in CeedQFunctionBuildKernel_Hip_ref()
77   code << "  // Output fields\n";  in CeedQFunctionBuildKernel_Hip_ref()
80     code << "  const CeedInt size_output_" << i << " = " << size << ";\n";  in CeedQFunctionBuildKernel_Hip_ref()
81     code << "  CeedScalar output_" << i << "[size_output_" << i << "];\n";  in CeedQFunctionBuildKernel_Hip_ref()
83   code << "  CeedScalar *outputs[" << CeedIntMax(num_output_fields, 1) << "];\n";  in CeedQFunctionBuildKernel_Hip_ref()
85     code << "  outputs[" << i << "] = output_" << i << ";\n";  in CeedQFunctionBuildKernel_Hip_ref()
87   code << "\n";  in CeedQFunctionBuildKernel_Hip_ref()
90   code << "  // Loop over quadrature points\n";  in CeedQFunctionBuildKernel_Hip_ref()
91 …code << "  for (CeedInt q = blockIdx.x * blockDim.x + threadIdx.x; q < Q; q += blockDim.x * gridDi…  in CeedQFunctionBuildKernel_Hip_ref()
94   code << "    // -- Load inputs\n";  in CeedQFunctionBuildKernel_Hip_ref()
96 …code << "    readQuads<size_input_" << i << ">(q, Q, fields.inputs[" << i << "], input_" << i << "…  in CeedQFunctionBuildKernel_Hip_ref()
98   code << "\n";  in CeedQFunctionBuildKernel_Hip_ref()
101   code << "    // -- Call QFunction\n";  in CeedQFunctionBuildKernel_Hip_ref()
102   code << "    " << qfunction_name << "(ctx, 1, inputs, outputs);\n\n";  in CeedQFunctionBuildKernel_Hip_ref()
105   code << "    // -- Write outputs\n";  in CeedQFunctionBuildKernel_Hip_ref()
107 …code << "    writeQuads<size_output_" << i << ">(q, Q, output_" << i << ", fields.outputs[" << i <…  in CeedQFunctionBuildKernel_Hip_ref()
109   code << "  }\n";  in CeedQFunctionBuildKernel_Hip_ref()
110   code << "}\n";  in CeedQFunctionBuildKernel_Hip_ref()
113 …CeedCallBackend(CeedCompile_Hip(ceed, code.str().c_str(), &data->module, 1, "BLOCK_SIZE", ceed_Hip…  in CeedQFunctionBuildKernel_Hip_ref()