1 // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors. 2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3 // 4 // SPDX-License-Identifier: BSD-2-Clause 5 // 6 // This file is part of CEED: http://github.com/ceed 7 8 #include <ceed/backend.h> 9 #include <ceed/ceed.h> 10 #include <ceed/jit-tools.h> 11 12 #include <iostream> 13 #include <sstream> 14 #include <string> 15 #include <string_view> 16 #include <sycl/sycl.hpp> 17 #include <vector> 18 19 #include "../sycl/ceed-sycl-compile.hpp" 20 #include "ceed-sycl-ref.hpp" 21 22 #define SUB_GROUP_SIZE_QF 16 23 24 //------------------------------------------------------------------------------ 25 // Build QFunction kernel 26 // 27 // TODO: Refactor 28 //------------------------------------------------------------------------------ 29 extern "C" int CeedQFunctionBuildKernel_Sycl(CeedQFunction qf) { 30 Ceed ceed; 31 Ceed_Sycl *data; 32 const char *read_write_kernel_path, *read_write_kernel_source; 33 const char *qfunction_name, *qfunction_source; 34 CeedInt num_input_fields, num_output_fields; 35 CeedQFunctionField *input_fields, *output_fields; 36 CeedQFunction_Sycl *impl; 37 38 CeedCallBackend(CeedQFunctionGetData(qf, (void **)&impl)); 39 // QFunction is built 40 if (impl->QFunction) return CEED_ERROR_SUCCESS; 41 42 CeedQFunctionGetCeed(qf, &ceed); 43 CeedCallBackend(CeedGetData(ceed, &data)); 44 45 // QFunction kernel generation 46 CeedCallBackend(CeedQFunctionGetFields(qf, &num_input_fields, &input_fields, &num_output_fields, &output_fields)); 47 48 std::vector<CeedInt> input_sizes(num_input_fields); 49 CeedQFunctionField *input_i = input_fields; 50 51 for (auto &size_i : input_sizes) { 52 CeedCallBackend(CeedQFunctionFieldGetSize(*input_i, &size_i)); 53 ++input_i; 54 } 55 56 std::vector<CeedInt> output_sizes(num_output_fields); 57 CeedQFunctionField *output_i = output_fields; 58 59 for (auto &size_i : output_sizes) { 60 CeedCallBackend(CeedQFunctionFieldGetSize(*output_i, &size_i)); 61 ++output_i; 62 } 63 64 CeedCallBackend(CeedQFunctionGetKernelName(qf, &qfunction_name)); 65 66 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading QFunction User Source -----\n"); 67 CeedCallBackend(CeedQFunctionLoadSourceToBuffer(qf, &qfunction_source)); 68 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading QFunction User Source Complete! -----\n"); 69 70 CeedCallBackend(CeedGetJitAbsolutePath(ceed, "ceed/jit-source/sycl/sycl-ref-qfunction.h", &read_write_kernel_path)); 71 72 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading QFunction Read/Write Kernel Source -----\n"); 73 CeedCallBackend(CeedLoadSourceToBuffer(ceed, read_write_kernel_path, &read_write_kernel_source)); 74 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading QFunction Read/Write Kernel Source Complete! -----\n"); 75 76 std::string_view qf_name_view(qfunction_name); 77 std::string_view qf_source_view(qfunction_source); 78 std::string_view rw_source_view(read_write_kernel_source); 79 const std::string kernel_name = "CeedKernelSyclRefQFunction_" + std::string(qf_name_view); 80 81 // Defintions 82 std::ostringstream code; 83 code << rw_source_view; 84 code << qf_source_view; 85 code << "\n"; 86 87 // Kernel function 88 // Here we are fixing a lower sub-group size value to avoid register spills 89 // This needs to be revisited if all qfunctions require this. 90 code << "__attribute__((intel_reqd_sub_group_size(" << SUB_GROUP_SIZE_QF << "))) __kernel void " << kernel_name 91 << "(__global void *ctx, CeedInt Q,\n"; 92 93 // OpenCL doesn't allow for structs with pointers. 94 // We will need to pass all of the arguments individually. 95 // Input parameters 96 for (CeedInt i = 0; i < num_input_fields; ++i) { 97 code << " " 98 << "__global const CeedScalar *in_" << i << ",\n"; 99 } 100 101 // Output parameters 102 code << " " 103 << "__global CeedScalar *out_0"; 104 for (CeedInt i = 1; i < num_output_fields; ++i) { 105 code << "\n, " 106 << "__global CeedScalar *out_" << i; 107 } 108 // Begin kernel function body 109 code << ") {\n\n"; 110 111 // Inputs 112 code << " // Input fields\n"; 113 for (CeedInt i = 0; i < num_input_fields; ++i) { 114 code << " CeedScalar U_" << i << "[" << input_sizes[i] << "];\n"; 115 } 116 code << " const CeedScalar *inputs[" << num_input_fields << "] = {U_0"; 117 for (CeedInt i = 1; i < num_input_fields; i++) { 118 code << ", U_" << i << "\n"; 119 } 120 code << "};\n\n"; 121 122 // Outputs 123 code << " // Output fields\n"; 124 for (CeedInt i = 0; i < num_output_fields; i++) { 125 code << " CeedScalar V_" << i << "[" << output_sizes[i] << "];\n"; 126 } 127 code << " CeedScalar *outputs[" << num_output_fields << "] = {V_0"; 128 for (CeedInt i = 1; i < num_output_fields; i++) { 129 code << ", V_" << i << "\n"; 130 } 131 code << "};\n\n"; 132 133 code << " const CeedInt q = get_global_linear_id();\n\n"; 134 135 code << "if(q < Q){ \n\n"; 136 137 // Load inputs 138 code << " // -- Load inputs\n"; 139 for (CeedInt i = 0; i < num_input_fields; i++) { 140 code << " readQuads(" << input_sizes[i] << ", Q, q, " 141 << "in_" << i << ", U_" << i << ");\n"; 142 } 143 code << "\n"; 144 145 // QFunction 146 code << " // -- Call QFunction\n"; 147 code << " " << qf_name_view << "(ctx, 1, inputs, outputs);\n\n"; 148 149 // Write outputs 150 code << " // -- Write outputs\n"; 151 for (CeedInt i = 0; i < num_output_fields; i++) { 152 code << " writeQuads(" << output_sizes[i] << ", Q, q, " 153 << "V_" << i << ", out_" << i << ");\n"; 154 } 155 code << "\n"; 156 157 // End kernel function body 158 code << "}\n"; 159 code << "}\n"; 160 161 // View kernel for debugging 162 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Generated QFunction Kernels:\n"); 163 CeedDebug(ceed, code.str().c_str()); 164 165 // Compile kernel 166 CeedCallBackend(CeedBuildModule_Sycl(ceed, code.str(), &impl->sycl_module)); 167 CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, kernel_name, &impl->QFunction)); 168 169 // Cleanup 170 CeedCallBackend(CeedFree(&qfunction_source)); 171 CeedCallBackend(CeedFree(&read_write_kernel_path)); 172 CeedCallBackend(CeedFree(&read_write_kernel_source)); 173 return CEED_ERROR_SUCCESS; 174 } 175 176 //------------------------------------------------------------------------------ 177