IntenseWebs/openvino
3
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

W/A for compilation issue (#12346)

Browse Source 
Previously INPUT/OUTPUT_SIZES was declared as an unnamed array macro.
However current ocl compiler deals with the value obtained from such an array non-constant, so that it results in a side effect in compiler optimization.
Here, a workaround for this issue is applied, i.e., declare the INPUT_OUTPUT_SIZES as a __const array instead of jit constants.
This commit is contained in:
Taylor Yeonbok Lee 2022-08-09 14:42:32 +09:00 committed by GitHub
parent f6a7d72a92
commit 751337d428
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
3 changed files with 18 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/plugins/intel_gpu/src/kernel_selector/core/common/jitter.cpp
View File

@ -237,8 +237,8 @@ public:
definitions.push_back({_name + "_SIZE", toCodeString(t.GetDims().size())});
definitions.push_back(
{_name + "_SIZES",
toVectorString(t.GetDims(), "size_t", KERNEL_SELECTOR_TENSOR_DIM_MAX, 1, [](const Tensor::Dim& d) { return d.v; })});
{_name + "_SIZES_DATA",
toVectorString(t.GetDims(), "", KERNEL_SELECTOR_TENSOR_DIM_MAX, 1, [](const Tensor::Dim& d) { return d.v; })});
definitions.push_back(
{_name + "_PITCHES",
toVectorString(t.GetDims(), "size_t", KERNEL_SELECTOR_TENSOR_DIM_MAX, 1, [](const Tensor::Dim& d) { return d.pitch; })});

5
src/plugins/intel_gpu/src/kernel_selector/core/common/jitter.h
View File

@ -105,7 +105,10 @@ inline std::string toVectorString(const VecT& vec,
ValT padFillingVal,
Func fetchFunc) {
std::stringstream ss;
ss << "(" << vectorType << " []){ ";
if (vectorType.length())
ss << "(" << vectorType << " []){ ";
else
ss << "{ ";
for (size_t i = 0; i < vec.size(); i++) ss << toCodeString(fetchFunc(vec[i])) << ",";
for (size_t i = vec.size(); i < maxDim; i++) ss << padFillingVal << ",";
ss << " } ";

13
src/plugins/intel_gpu/src/kernel_selector/core/common/kernel_base_opencl.cpp
View File

@ -92,14 +92,25 @@ std::pair<std::string, std::string> KernelBaseOpenCL::CreateJit(const std::strin
.add_line("// Kernel name: " + kernel_id)
.value_macro("KERNEL(name)", "__kernel void " + kernel_id)
.decoration_macro("FUNC", "", kernel_id)
.decoration_macro("FUNC_CALL", "", kernel_id);
.decoration_macro("FUNC_CALL", "", kernel_id)
.decoration_macro("CONST_ARRAY_DECL", "__constant size_t ", kernel_id + " []")
.decoration_macro("CONST_ARRAY_REF", "", kernel_id);
undefs += "#undef KERNEL\n";
undefs += "#undef FUNC\n";
undefs += "#undef FUNC_CALL\n";
undefs += "#undef CONST_ARRAY_DECL\n";
undefs += "#undef CONST_ARRAY_REF\n";
for (auto& definition : constants.GetDefinitions()) {
code.value_macro(definition.first, definition.second);
if (definition.first.find("SIZES_DATA") != std::string::npos) {
auto size_arr_data = definition.first;
auto size_arr = size_arr_data.erase(size_arr_data.find("_DATA") , 5);
code.add_line("CONST_ARRAY_DECL(" + size_arr + ") = " + definition.first + ";");
code.value_macro(size_arr, "CONST_ARRAY_REF(" + size_arr + ")");
}
undefs += "#ifdef " + definition.first.substr(0, definition.first.find('(')) + "\n";
undefs += "#undef " + definition.first.substr(0, definition.first.find('(')) + "\n";
undefs += "#endif\n";