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 { 74 char *tmp; 75 CeedCallBackend(CeedLoadSourceToBuffer(ceed, read_write_kernel_path, &tmp)); 76 read_write_kernel_source = tmp; 77 } 78 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading QFunction Read/Write Kernel Source Complete! -----\n"); 79 80 std::string_view qf_name_view(qfunction_name); 81 std::string_view qf_source_view(qfunction_source); 82 std::string_view rw_source_view(read_write_kernel_source); 83 const std::string kernel_name = "CeedKernelSyclRefQFunction_" + std::string(qf_name_view); 84 85 // Defintions 86 std::ostringstream code; 87 code << rw_source_view; 88 code << qf_source_view; 89 code << "\n"; 90 91 // Kernel function 92 // Here we are fixing a lower sub-group size value to avoid register spills 93 // This needs to be revisited if all qfunctions require this. 94 code << "__attribute__((intel_reqd_sub_group_size(" << SUB_GROUP_SIZE_QF << "))) __kernel void " << kernel_name 95 << "(__global void *ctx, CeedInt Q,\n"; 96 97 // OpenCL doesn't allow for structs with pointers. 98 // We will need to pass all of the arguments individually. 99 // Input parameters 100 for (CeedInt i = 0; i < num_input_fields; ++i) { 101 code << " " 102 << "__global const CeedScalar *in_" << i << ",\n"; 103 } 104 105 // Output parameters 106 code << " " 107 << "__global CeedScalar *out_0"; 108 for (CeedInt i = 1; i < num_output_fields; ++i) { 109 code << "\n, " 110 << "__global CeedScalar *out_" << i; 111 } 112 // Begin kernel function body 113 code << ") {\n\n"; 114 115 // Inputs 116 code << " // Input fields\n"; 117 for (CeedInt i = 0; i < num_input_fields; ++i) { 118 code << " CeedScalar U_" << i << "[" << input_sizes[i] << "];\n"; 119 } 120 code << " const CeedScalar *inputs[" << num_input_fields << "] = {U_0"; 121 for (CeedInt i = 1; i < num_input_fields; i++) { 122 code << ", U_" << i << "\n"; 123 } 124 code << "};\n\n"; 125 126 // Outputs 127 code << " // Output fields\n"; 128 for (CeedInt i = 0; i < num_output_fields; i++) { 129 code << " CeedScalar V_" << i << "[" << output_sizes[i] << "];\n"; 130 } 131 code << " CeedScalar *outputs[" << num_output_fields << "] = {V_0"; 132 for (CeedInt i = 1; i < num_output_fields; i++) { 133 code << ", V_" << i << "\n"; 134 } 135 code << "};\n\n"; 136 137 code << " const CeedInt q = get_global_linear_id();\n\n"; 138 139 code << "if(q < Q){ \n\n"; 140 141 // Load inputs 142 code << " // -- Load inputs\n"; 143 for (CeedInt i = 0; i < num_input_fields; i++) { 144 code << " readQuads(" << input_sizes[i] << ", Q, q, " 145 << "in_" << i << ", U_" << i << ");\n"; 146 } 147 code << "\n"; 148 149 // QFunction 150 code << " // -- Call QFunction\n"; 151 code << " " << qf_name_view << "(ctx, 1, inputs, outputs);\n\n"; 152 153 // Write outputs 154 code << " // -- Write outputs\n"; 155 for (CeedInt i = 0; i < num_output_fields; i++) { 156 code << " writeQuads(" << output_sizes[i] << ", Q, q, " 157 << "V_" << i << ", out_" << i << ");\n"; 158 } 159 code << "\n"; 160 161 // End kernel function body 162 code << "}\n"; 163 code << "}\n"; 164 165 // View kernel for debugging 166 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Generated QFunction Kernels:\n"); 167 CeedDebug(ceed, code.str().c_str()); 168 169 // Compile kernel 170 CeedCallBackend(CeedBuildModule_Sycl(ceed, code.str(), &impl->sycl_module)); 171 CeedCallBackend(CeedGetKernel_Sycl(ceed, impl->sycl_module, kernel_name, &impl->QFunction)); 172 173 // Cleanup 174 CeedCallBackend(CeedFree(&qfunction_source)); 175 CeedCallBackend(CeedFree(&read_write_kernel_path)); 176 CeedCallBackend(CeedFree(&read_write_kernel_source)); 177 return CEED_ERROR_SUCCESS; 178 } 179 180 //------------------------------------------------------------------------------ 181