[NVIDIA] Change IsCudaGraphCompatible() interface to GetCudaGraphComp…

…atibility() using enum
openvinotoolkit · Nov 15, 2023 · 9a8412c · 9a8412c
1 parent e68fc18
commit 9a8412c
Show file tree

Hide file tree

Showing 109 changed files with 167 additions and 143 deletions.
diff --git a/modules/nvidia_plugin/src/cuda_graph_topology_runner.cpp b/modules/nvidia_plugin/src/cuda_graph_topology_runner.cpp
@@ -35,12 +35,12 @@ CudaGraphTopologyRunner::CudaGraphTopologyRunner(const CreationContext& context,
     const auto totalSize = origSequence.size();
     OPENVINO_ASSERT(totalSize != 0, "ExecSequence size is 0");
 
-    bool isLastOpCompatible = origSequence[0]->IsCudaGraphCompatible();
+    CudaGraphCompatibility lastOpCompatibility = origSequence[0]->GetCudaGraphCompatibility();
     currentSequence.push_back(origSequence[0]);
     for (size_t i = 1; i < totalSize; ++i) {
         const auto& op = origSequence[i];
-        if (std::dynamic_pointer_cast<const TensorIteratorOp>(op) || op->IsCudaGraphCompatible() != isLastOpCompatible) {
-            isLastOpCompatible = !isLastOpCompatible;
+        if (auto c = op->GetCudaGraphCompatibility(); c != lastOpCompatibility) {
+            lastOpCompatibility = c;
             sequences.emplace_back(std::move(currentSequence));
             currentSequence.clear();
         }
@@ -51,7 +51,7 @@ CudaGraphTopologyRunner::CudaGraphTopologyRunner(const CreationContext& context,
     const auto& memoryManager = orig_subgraph_.memoryManager();
     for (auto&& sequence : sequences) {
         subgraphs_.emplace_back(context, model, std::move(sequence), memoryManager);
-        if (subgraphs_[subgraphs_.size() - 1].IsCudaGraphCompatible()) {
+        if (subgraphs_.back().GetCudaGraphCompatibility() != CudaGraphCompatibility::NONE) {
             ++cuda_graphs_count_;
         }
     }
@@ -62,7 +62,14 @@ void CudaGraphTopologyRunner::Run(InferenceRequestContext& context, const Device
     auto& graphContext = context.getCudaGraphContext();
     std::size_t graphIndex = 0;
     for (auto& subgraph : subgraphs_) {
-        if (auto ti = getTI(subgraph)) {
+        auto compatibility = subgraph.GetCudaGraphCompatibility();
+        if (compatibility == CudaGraphCompatibility::FULL) {
+            graphContext.select_current_graph(graphIndex);
+            graphContext.get_current_graph_info().launch(stream);
+            graphIndex++;
+        } else if (compatibility == CudaGraphCompatibility::SPECIAL) {
+            // TODO: remove
+            auto ti = getTI(subgraph);
             CUDA::DevicePointer<void*> mutableBuffer{memoryBlock.view().data()};
             const auto& memoryManager = *subgraph.memoryManager();
             const auto& inputTensors = memoryManager.inputTensorPointers(*ti, mutableBuffer);
@@ -71,10 +78,6 @@ void CudaGraphTopologyRunner::Run(InferenceRequestContext& context, const Device
             graphContext.select_current_graph(graphIndex);
             ti->ExecuteGraph(context, inputTensors, outputTensors, workBuffers);
             graphIndex++;
-        } else if (subgraph.IsCudaGraphCompatible()) {
-            graphContext.select_current_graph(graphIndex);
-            graphContext.get_current_graph_info().launch(stream);
-            graphIndex++;
         } else {
             Workbuffers workbuffers{};
             workbuffers.mutable_buffers.emplace_back(memoryBlock.view().data());
@@ -92,10 +95,8 @@ void CudaGraphTopologyRunner::Capture(InferenceRequestContext& context,
     for (const auto& subgraph : subgraphs_) {
         Workbuffers workbuffers{};
         workbuffers.mutable_buffers.emplace_back(memoryBlock.view().data());
-        if (getTI(subgraph)) {
-            graphContext.add_new_graph_info();
-            subgraph.Capture(context, {}, {}, workbuffers);
-        } else if (subgraph.IsCudaGraphCompatible()) {
+        auto compatibility = subgraph.GetCudaGraphCompatibility();
+        if (compatibility == CudaGraphCompatibility::FULL) {
             graphContext.add_new_graph_info();
             CUDA::GraphCapture capture{stream};
             {
@@ -104,6 +105,9 @@ void CudaGraphTopologyRunner::Capture(InferenceRequestContext& context,
             }
             const auto& graph = capture.getGraph();
             graphContext.set_current_graph(graph);
+        } else if (compatibility == CudaGraphCompatibility::SPECIAL) {
+            graphContext.add_new_graph_info();
+            subgraph.Capture(context, {}, {}, workbuffers);
         }
     }
 }

diff --git a/modules/nvidia_plugin/src/cuda_operation_base.hpp b/modules/nvidia_plugin/src/cuda_operation_base.hpp
@@ -30,6 +30,8 @@ namespace nvidia_gpu {
 template <typename T>
 using DevicePointer = CUDA::DevicePointer<T>;
 
+enum class CudaGraphCompatibility { NONE, FULL, SPECIAL };
+
 class IOperationExec {
 public:
     using Inputs = gsl::span<const CUDA::DevicePointer<const void*>>;
@@ -46,7 +48,7 @@ class IOperationExec {
                          Inputs inputTensors,
                          Outputs outputTensors,
                          const Workbuffers& workbuffers) const = 0;
-    virtual bool IsCudaGraphCompatible() const = 0;
+    virtual CudaGraphCompatibility GetCudaGraphCompatibility() const = 0;
     virtual void InitSharedImmutableWorkbuffers(const Buffers&) = 0;
     virtual WorkbufferRequest GetWorkBufferRequest() const = 0;
     virtual const WorkbufferIds& GetWorkbufferIds() const = 0;
@@ -81,7 +83,7 @@ class OperationBase : public IOperationExec, public IOperationMeta, public std::
                   IndexCollection&& inputIds,
                   IndexCollection&& outputIds);
 
-    bool IsCudaGraphCompatible() const override { return false; }
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override { return CudaGraphCompatibility::NONE; }
 
     WorkbufferRequest GetWorkBufferRequest() const override {
         return {};  // Most operators do not need workbuffers

diff --git a/modules/nvidia_plugin/src/ops/activation_forward_cudnn_base.cpp b/modules/nvidia_plugin/src/ops/activation_forward_cudnn_base.cpp
@@ -59,7 +59,9 @@ void ActivationForwardCuDnnOpBase::Execute(const InferenceRequestContext& contex
                                                              outputTensors[0].get());
 }
 
-bool ActivationForwardCuDnnOpBase::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility ActivationForwardCuDnnOpBase::GetCudaGraphCompatibility() const {
+    return CudaGraphCompatibility::FULL;
+}
 
 }  // namespace nvidia_gpu
 }  // namespace ov
diff --git a/modules/nvidia_plugin/src/ops/activation_forward_cudnn_base.hpp b/modules/nvidia_plugin/src/ops/activation_forward_cudnn_base.hpp
@@ -31,7 +31,7 @@ class ActivationForwardCuDnnOpBase : public OperationCuDnn {
                  Outputs outputTensors,
                  const Workbuffers&) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 protected:
     std::unique_ptr<CUDA::DnnActivationDescriptor> op_desc_;

diff --git a/modules/nvidia_plugin/src/ops/avgpool.cpp b/modules/nvidia_plugin/src/ops/avgpool.cpp
@@ -30,7 +30,7 @@ void AvgPoolOp::Execute(const InferenceRequestContext& context,
                   outputs[PoolingImpl::output_index].get());
 }
 
-bool AvgPoolOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility AvgPoolOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 OPERATION_REGISTER(AvgPoolOp, AvgPool);
 

diff --git a/modules/nvidia_plugin/src/ops/avgpool.hpp b/modules/nvidia_plugin/src/ops/avgpool.hpp
@@ -23,7 +23,7 @@ class AvgPoolOp : public OperationCuDnn {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     PoolingImpl impl_;

diff --git a/modules/nvidia_plugin/src/ops/broadcast.cpp b/modules/nvidia_plugin/src/ops/broadcast.cpp
@@ -65,7 +65,7 @@ void BroadcastOp::Execute(const InferenceRequestContext& context,
     (*kernel_)(stream, inputs[0].get(), broadcast_params_->mapper(workbuffers.immutable_buffers), outputs[0].get());
 }
 
-bool BroadcastOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility BroadcastOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 WorkbufferRequest BroadcastOp::GetWorkBufferRequest() const { return {immutable_buffer_sizes_, {}}; }
 

diff --git a/modules/nvidia_plugin/src/ops/broadcast.hpp b/modules/nvidia_plugin/src/ops/broadcast.hpp
@@ -27,7 +27,7 @@ class BroadcastOp : public OperationBase {
     WorkbufferRequest GetWorkBufferRequest() const override;
     void InitSharedImmutableWorkbuffers(const Buffers& buffers) override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     std::vector<WorkbufferRequest::size_in_bytes_t> immutable_buffer_sizes_;

diff --git a/modules/nvidia_plugin/src/ops/clamp_cuda.cpp b/modules/nvidia_plugin/src/ops/clamp_cuda.cpp
@@ -51,7 +51,7 @@ void ClampCudaOp::Execute(const InferenceRequestContext& context,
     (*kernel_)(context.getThreadContext().stream().get(), inputTensors[0].get(), outputTensors[0].get());
 }
 
-bool ClampCudaOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility ClampCudaOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 }  // namespace nvidia_gpu
 }  // namespace ov
diff --git a/modules/nvidia_plugin/src/ops/clamp_cuda.hpp b/modules/nvidia_plugin/src/ops/clamp_cuda.hpp
@@ -26,7 +26,7 @@ class ClampCudaOp : public OperationBase {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     std::optional<kernel::Clamp> kernel_;

diff --git a/modules/nvidia_plugin/src/ops/clamp_cudnn.cpp b/modules/nvidia_plugin/src/ops/clamp_cudnn.cpp
@@ -97,7 +97,7 @@ void ClampCuDnnOp::Execute(const InferenceRequestContext& context,
                                                     outputTensors[0].get());
 }
 
-bool ClampCuDnnOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility ClampCuDnnOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 void ClampCuDnnOp::InitSharedImmutableWorkbuffers(const Buffers& buffers) {
     switch (data_type_) {

diff --git a/modules/nvidia_plugin/src/ops/clamp_cudnn.hpp b/modules/nvidia_plugin/src/ops/clamp_cudnn.hpp
@@ -33,7 +33,7 @@ class ClampCuDnnOp : public OperationCuDnn {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
     void InitSharedImmutableWorkbuffers(const Buffers& buffers) override;
     WorkbufferRequest GetWorkBufferRequest() const override;
 

diff --git a/modules/nvidia_plugin/src/ops/comparison.cpp b/modules/nvidia_plugin/src/ops/comparison.cpp
@@ -84,7 +84,7 @@ Comparison::Comparison(const CreationContext& context,
                                  threads_per_block};
 }
 
-bool Comparison::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility Comparison::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 void Comparison::Execute(const InferenceRequestContext& context,
                          Inputs inputs,

diff --git a/modules/nvidia_plugin/src/ops/comparison.hpp b/modules/nvidia_plugin/src/ops/comparison.hpp
@@ -18,7 +18,7 @@ class Comparison : public OperationBase {
                IndexCollection&& outputIds,
                kernel::Comparison::Op_t operation_type);
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     void calculateOffsets();

diff --git a/modules/nvidia_plugin/src/ops/concat.cpp b/modules/nvidia_plugin/src/ops/concat.cpp
@@ -95,7 +95,7 @@ void ConcatOp::Execute(const InferenceRequestContext& context,
                       outputs[0].get());
 }
 
-bool ConcatOp::IsCudaGraphCompatible() const { return false; }
+CudaGraphCompatibility ConcatOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::NONE; }
 
 OPERATION_REGISTER(ConcatOp, Concat);
 }  // namespace nvidia_gpu

diff --git a/modules/nvidia_plugin/src/ops/concat.hpp b/modules/nvidia_plugin/src/ops/concat.hpp
@@ -28,7 +28,7 @@ class ConcatOp : public OperationBase {
     WorkbufferRequest GetWorkBufferRequest() const override;
     void InitSharedImmutableWorkbuffers(const Buffers&) override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     size_t immutableWbSize() const { return concat_kernel_.value().immutableWbSize(); }

diff --git a/modules/nvidia_plugin/src/ops/convert.cpp b/modules/nvidia_plugin/src/ops/convert.cpp
@@ -55,7 +55,7 @@ void ConvertOp::Execute(const InferenceRequestContext& context,
     (*convert_kernel_)(stream.get(), outputs[0].get(), inputs[0].get());
 }
 
-bool ConvertOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility ConvertOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 OPERATION_REGISTER(ConvertOp, Convert);
 

diff --git a/modules/nvidia_plugin/src/ops/convert.hpp b/modules/nvidia_plugin/src/ops/convert.hpp
@@ -24,7 +24,7 @@ class ConvertOp : public OperationBase {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
     using Type_t = ov::element::Type_t;
     using convert_t = void (*)(

diff --git a/modules/nvidia_plugin/src/ops/convert_color_i420.hpp b/modules/nvidia_plugin/src/ops/convert_color_i420.hpp
@@ -91,7 +91,7 @@ class I420ConvertColorBase : public OperationBase {
         }
     }
 
-    bool IsCudaGraphCompatible() const override { return true; }
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override { return CudaGraphCompatibility::FULL; }
 
 private:
     std::optional<TKernel> kernel_;

diff --git a/modules/nvidia_plugin/src/ops/convert_color_nv12.hpp b/modules/nvidia_plugin/src/ops/convert_color_nv12.hpp
@@ -90,7 +90,7 @@ class NV12ConvertColorBase : public OperationBase {
         }
     }
 
-    bool IsCudaGraphCompatible() const override { return true; }
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override { return CudaGraphCompatibility::FULL; }
 
 private:
     std::optional<TKernel> kernel_;

diff --git a/modules/nvidia_plugin/src/ops/convolution_backprop_data.cpp b/modules/nvidia_plugin/src/ops/convolution_backprop_data.cpp
@@ -43,8 +43,8 @@ void ConvBackpropDataOp<T>::Execute(const InferenceRequestContext& context,
 }
 
 template <typename T>
-bool ConvBackpropDataOp<T>::IsCudaGraphCompatible() const {
-    return true;
+CudaGraphCompatibility ConvBackpropDataOp<T>::GetCudaGraphCompatibility() const {
+    return CudaGraphCompatibility::FULL;
 }
 
 OPERATION_REGISTER(ConvolutionBackpropDataOp, ConvolutionBackpropData);

diff --git a/modules/nvidia_plugin/src/ops/convolution_backprop_data.hpp b/modules/nvidia_plugin/src/ops/convolution_backprop_data.hpp
@@ -32,7 +32,7 @@ class ConvBackpropDataOp : public OperationCuDnn {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
     void InitSharedImmutableWorkbuffers(const IOperationExec::Buffers&) override;
     WorkbufferRequest GetWorkBufferRequest() const override;

diff --git a/modules/nvidia_plugin/src/ops/convolution_cudnn.cpp b/modules/nvidia_plugin/src/ops/convolution_cudnn.cpp
@@ -44,7 +44,7 @@ void ConvolutionCuDnn::Execute(const InferenceRequestContext& context,
     throwIfError(status);
 }
 
-bool ConvolutionCuDnn::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility ConvolutionCuDnn::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 WorkbufferRequest ConvolutionCuDnn::GetWorkBufferRequest() const {
     if (descs_.Algo().memory != 0)

diff --git a/modules/nvidia_plugin/src/ops/convolution_cudnn.hpp b/modules/nvidia_plugin/src/ops/convolution_cudnn.hpp
@@ -29,7 +29,7 @@ class ConvolutionCuDnn : public OperationCuDnn {
 
     WorkbufferRequest GetWorkBufferRequest() const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     Convolution::Details::ConvolutionDescriptorsCuDnn descs_;

diff --git a/modules/nvidia_plugin/src/ops/convolution_cudnn_be.cpp b/modules/nvidia_plugin/src/ops/convolution_cudnn_be.cpp
@@ -148,7 +148,7 @@ void ConvolutionCuDnnBE::Execute(const InferenceRequestContext& context,
     throwIfError(::cudnnBackendExecute(context.getThreadContext().dnnHandle().get(), plan->get(), variantPack->get()));
 }
 
-bool ConvolutionCuDnnBE::IsCudaGraphCompatible() const { return false; }
+CudaGraphCompatibility ConvolutionCuDnnBE::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::NONE; }
 
 std::shared_ptr<CUDA::DnnBETensorDescriptor> ConvolutionCuDnnBE::MakeTensorDescriptor(int64_t id,
                                                                                       cudnnDataType_t element_type,

diff --git a/modules/nvidia_plugin/src/ops/convolution_cudnn_be.hpp b/modules/nvidia_plugin/src/ops/convolution_cudnn_be.hpp
@@ -33,7 +33,7 @@ class ConvolutionCuDnnBE : public OperationCuDnn {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
     WorkbufferRequest GetWorkBufferRequest() const override;
 

diff --git a/modules/nvidia_plugin/src/ops/cudnn_tensor_op_base.cpp b/modules/nvidia_plugin/src/ops/cudnn_tensor_op_base.cpp
@@ -166,7 +166,7 @@ void CuDnnTensorOpBase::Execute(const InferenceRequestContext& context,
                                                     outputTensors[0].get());
 }
 
-bool CuDnnTensorOpBase::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility CuDnnTensorOpBase::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 CuDnnTensorOpBase::IoParams::IoParams(const ov::Node& node, const Type& io_type, int index)
     : type_(convertDataType<cudnnDataType_t>(io_type == Type::INPUT ? node.get_input_element_type(index)

diff --git a/modules/nvidia_plugin/src/ops/cudnn_tensor_op_base.hpp b/modules/nvidia_plugin/src/ops/cudnn_tensor_op_base.hpp
@@ -24,7 +24,7 @@ class CuDnnTensorOpBase : public OperationCuDnn {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     struct IoParams {

diff --git a/modules/nvidia_plugin/src/ops/detection_output.cpp b/modules/nvidia_plugin/src/ops/detection_output.cpp
@@ -107,7 +107,7 @@ void DetectionOutputOp::Execute(const InferenceRequestContext& context,
     }
 }
 
-bool DetectionOutputOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility DetectionOutputOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 void DetectionOutputOp::InitSharedImmutableWorkbuffers(const Buffers& buffers) {
     kernel_.value().initSharedImmutableWorkbuffers(buffers);

diff --git a/modules/nvidia_plugin/src/ops/detection_output.hpp b/modules/nvidia_plugin/src/ops/detection_output.hpp
@@ -25,7 +25,7 @@ class DetectionOutputOp : public OperationBase {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
     void InitSharedImmutableWorkbuffers(const Buffers& buffers) override;
     WorkbufferRequest GetWorkBufferRequest() const override;

diff --git a/modules/nvidia_plugin/src/ops/elementwise_binary.hpp b/modules/nvidia_plugin/src/ops/elementwise_binary.hpp
@@ -59,7 +59,7 @@ class ElementwiseBinaryOp : public OperationBase {
                    static_cast<void*>(outputTensors[0].get()));
     }
 
-    bool IsCudaGraphCompatible() const override { return true; }
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override { return CudaGraphCompatibility::FULL; }
 
     void InitSharedImmutableWorkbuffers(const IOperationExec::Buffers& buffers) override {
         in0_broadcast_params_->initWorkbuffers(buffers);

diff --git a/modules/nvidia_plugin/src/ops/elementwise_unary.hpp b/modules/nvidia_plugin/src/ops/elementwise_unary.hpp
@@ -46,7 +46,7 @@ class ElementwiseUnaryOp : public OperationBase {
         (*kernel_)(stream.get(), inputTensors[0].get(), outputTensors[0].get());
     }
 
-    bool IsCudaGraphCompatible() const override { return true; }
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override { return CudaGraphCompatibility::FULL; }
 
 private:
     std::optional<Kernel> kernel_;

diff --git a/modules/nvidia_plugin/src/ops/elu.cpp b/modules/nvidia_plugin/src/ops/elu.cpp
@@ -45,7 +45,7 @@ void EluOp::Execute(const InferenceRequestContext& context,
     (*kernel_)(stream.get(), inputTensors[0].get(), outputTensors[0].get());
 }
 
-bool EluOp::IsCudaGraphCompatible() const { return true; }
+CudaGraphCompatibility EluOp::GetCudaGraphCompatibility() const { return CudaGraphCompatibility::FULL; }
 
 OPERATION_REGISTER(EluOp, Elu);
 }  // namespace nvidia_gpu

diff --git a/modules/nvidia_plugin/src/ops/elu.hpp b/modules/nvidia_plugin/src/ops/elu.hpp
@@ -23,7 +23,7 @@ class EluOp : public OperationBase {
                  Outputs outputTensors,
                  const Workbuffers& workbuffers) const override;
 
-    bool IsCudaGraphCompatible() const override;
+    CudaGraphCompatibility GetCudaGraphCompatibility() const override;
 
 private:
     std::optional<kernel::Elu> kernel_;