openvino: [Bug] GPU extension

@wang7393 I’ve tried the following code based on @mbencer 's test onnx_op_extension_mixed_legacy_and_new_api

class CustomRelu : public ov::op::Op {
public:
    OPENVINO_OP("CustomRelu");
    OPENVINO_FRAMEWORK_MAP(onnx)

    CustomRelu() = default;

    CustomRelu(const ov::Output<ov::Node>& input) : Op({input}) {
        constructor_validate_and_infer_types();
    }

    void validate_and_infer_types() override {
        set_output_size(1);
        set_output_type(0, get_input_element_type(0), get_input_partial_shape(0));
    }

    bool visit_attributes(ov::AttributeVisitor& visitor) override {
        return true;
    }

    std::shared_ptr<ov::Node> clone_with_new_inputs(const ov::OutputVector& inputs) const override {
        return std::make_shared<CustomRelu>(inputs[0]);
    }
};

TEST(ONNXOpExtensionViaCommonConstructor, onnx_op_extension_mixed_legacy_and_new_api) {
    const auto input_model_path = CommonTestUtils::getModelFromTestModelZoo(
        ov::util::path_join({TEST_ONNX_MODELS_DIRNAME, "relu_custom_domain.onnx"}));
    ov::Core core;
    const auto new_api_ext = std::make_shared<ov::frontend::onnx::OpExtension<CustomRelu>>("CustomRelu", "my_custom_domain");
    core.add_extension(new_api_ext);
    std::shared_ptr<ov::Model> model = core.read_model(input_model_path);

    {
        auto compiled_model = core.compile_model(model, "GPU",  {{"CONFIG_FILE", "custom_layer_example.xml"}});

        auto request = compiled_model.create_infer_request();
        request.infer();
    }

    {
        core.set_property("GPU", {{"CONFIG_FILE", "custom_layer_example.xml"}});
        auto compiled_model = core.compile_model(model, "GPU");

        auto request = compiled_model.create_infer_request();
        request.infer();
    }

}

custom_layer_example.xml:

<CustomLayer name="CustomRelu" type="SimpleGPU" version="1">
    <Kernel entry="custom_relu">
        <Source filename="custom_relu.cl"/>
    </Kernel>
    <Buffers>
        <Tensor arg-index="0" type="input" port-index="0"/>
        <Tensor arg-index="1" type="output" port-index="0" format="BFYX"/>
    </Buffers>
    <CompilerOptions options="-cl-mad-enable"/>
    <WorkSizes global="2,3,4"/>
</CustomLayer>

custom_relu.cl:

__kernel void custom_relu(__global const float* in, __global float* out) {
    printf("%d %d %d\n", get_global_id(0), get_global_id(1), get_global_id(2));
}

And both versions of setting custom config for GPU (via Core::set_property() and CompiledModel c-tor worked fine and printed global ids.

I’d suggest checking IR after read_model() as I don’t see any issues with GPU extensibility mechanism so far. You can try to save the model via ov::serialize(model, "ir.xml", "ir.bin"); and check if your custom op has been generated correctly by FE extension. Also, if you’re able to build OV from sources, I’d suggest building it with -DENABLE_DEBUG_CAPS=ON cmake option and then running your app with OV_GPU_VERBOSE=2 environment variable. In the verbose logs you’re supposed to see something like this:

GPU_Debug: program.cpp:426:CreateSingleLayerPrimitive: Process op::vopset1::Parameter operation (friendly_name=A)
GPU_Debug: program.cpp:426:CreateSingleLayerPrimitive: Process op::vextension::CustomRelu operation (friendly_name=B) // Creation of custom op
GPU_Debug: program.cpp:426:CreateSingleLayerPrimitive: Process op::vopset1::Result operation (friendly_name=B/sink_port_0)

...
GPU_Debug: program.cpp:119:program: Program config
internal properties:
...
        CONFIG_FILE: custom_layer_example.xml  // path to your config file
...
GPU_Debug: primitive_inst.cpp:681:execute: customrelu:B: execute // Execution of the custom layer