diff --git a/runtime/api/api.cpp b/runtime/api/api.cpp
index e43138134a..9d2400594a 100644
--- a/runtime/api/api.cpp
+++ b/runtime/api/api.cpp
@@ -2585,8 +2585,9 @@ cl_int CL_API_CALL clEnqueueNDRangeKernel(cl_command_queue commandQueue,
 
     auto pKernel = castToObjectOrAbort<Kernel>(kernel);
     TakeOwnershipWrapper<Kernel> kernelOwnership(*pKernel, gtpinIsGTPinInitialized());
-
-    gtpinNotifyKernelSubmit(kernel, pCommandQueue);
+    if (gtpinIsGTPinInitialized()) {
+        gtpinNotifyKernelSubmit(kernel, pCommandQueue);
+    }
 
     retVal = pCommandQueue->enqueueKernel(
         kernel,
diff --git a/runtime/command_queue/enqueue_common.h b/runtime/command_queue/enqueue_common.h
index 713d025212..8a176f1dda 100644
--- a/runtime/command_queue/enqueue_common.h
+++ b/runtime/command_queue/enqueue_common.h
@@ -553,7 +553,9 @@ CompletionStamp CommandQueueHw<GfxFamily>::enqueueNonBlocked(
 
     DEBUG_BREAK_IF(taskLevel >= Event::eventNotReady);
 
-    gtpinNotifyPreFlushTask(this);
+    if (gtpinIsGTPinInitialized()) {
+        gtpinNotifyPreFlushTask(this);
+    }
 
     CompletionStamp completionStamp = commandStreamReceiver.flushTask(
         commandStream,
diff --git a/runtime/command_stream/command_stream_receiver.cpp b/runtime/command_stream/command_stream_receiver.cpp
index eb5390d2b2..93f0a222b6 100644
--- a/runtime/command_stream/command_stream_receiver.cpp
+++ b/runtime/command_stream/command_stream_receiver.cpp
@@ -193,7 +193,9 @@ bool CommandStreamReceiver::waitForCompletionWithTimeout(bool enableTimeout, int
         }
     }
     if (*getTagAddress() >= taskCountToWait) {
-        gtpinNotifyTaskCompletion(taskCountToWait);
+        if (gtpinIsGTPinInitialized()) {
+            gtpinNotifyTaskCompletion(taskCountToWait);
+        }
         return true;
     }
     return false;
diff --git a/runtime/command_stream/command_stream_receiver_hw.inl b/runtime/command_stream/command_stream_receiver_hw.inl
index 5a99bc335a..f99a912f65 100644
--- a/runtime/command_stream/command_stream_receiver_hw.inl
+++ b/runtime/command_stream/command_stream_receiver_hw.inl
@@ -398,7 +398,9 @@ CompletionStamp CommandStreamReceiverHw<GfxFamily>::flushTask(
 
     this->taskLevel += levelClosed ? 1 : 0;
 
-    gtpinNotifyFlushTask(completionStamp.taskCount);
+    if (gtpinIsGTPinInitialized()) {
+        gtpinNotifyFlushTask(completionStamp.taskCount);
+    }
 
     return completionStamp;
 }
diff --git a/runtime/gtpin/gtpin_callback_stubs.cpp b/runtime/gtpin/gtpin_callback_stubs.cpp
index 6aab8f1f28..2ede6b62a2 100644
--- a/runtime/gtpin/gtpin_callback_stubs.cpp
+++ b/runtime/gtpin/gtpin_callback_stubs.cpp
@@ -24,6 +24,7 @@
 #include <cstdint>
 
 namespace OCLRT {
+bool isGTPinInitialized = false;
 
 void gtpinNotifyContextCreate(cl_context context) {
 }
@@ -54,8 +55,4 @@ void gtpinNotifyUpdateResidencyList(void *pKernel, void *pResidencyVector) {
 
 void gtpinNotifyPlatformShutdown() {
 }
-
-bool gtpinIsGTPinInitialized() {
-    return false;
-}
 }
diff --git a/runtime/gtpin/gtpin_callbacks.cpp b/runtime/gtpin/gtpin_callbacks.cpp
index f5d80c64db..6acb727e9f 100644
--- a/runtime/gtpin/gtpin_callbacks.cpp
+++ b/runtime/gtpin/gtpin_callbacks.cpp
@@ -30,6 +30,7 @@
 #include "runtime/device/device_info.h"
 #include "runtime/gtpin/gtpin_defs.h"
 #include "runtime/gtpin/gtpin_hw_helper.h"
+#include "runtime/gtpin/gtpin_notify.h"
 #include "runtime/kernel/kernel.h"
 #include "runtime/mem_obj/buffer.h"
 #include "runtime/memory_manager/surface.h"
@@ -42,7 +43,6 @@ using namespace gtpin;
 
 namespace OCLRT {
 
-extern bool isGTPinInitialized;
 extern gtpin::ocl::gtpin_events_t GTPinCallbacks;
 
 igc_init_t *pIgcInfo = nullptr;
@@ -242,8 +242,4 @@ void gtpinNotifyPlatformShutdown() {
         kernelExecQueue.clear();
     }
 }
-
-bool gtpinIsGTPinInitialized() {
-    return isGTPinInitialized;
-}
 }
diff --git a/runtime/gtpin/gtpin_notify.h b/runtime/gtpin/gtpin_notify.h
index 42d6e158ea..8d874f76cb 100644
--- a/runtime/gtpin/gtpin_notify.h
+++ b/runtime/gtpin/gtpin_notify.h
@@ -23,6 +23,8 @@
 #pragma once
 
 namespace OCLRT {
+extern bool isGTPinInitialized;
+
 void gtpinNotifyContextCreate(cl_context context);
 void gtpinNotifyContextDestroy(cl_context context);
 void gtpinNotifyKernelCreate(cl_kernel kernel);
@@ -33,5 +35,5 @@ void gtpinNotifyTaskCompletion(uint32_t completedTaskCount);
 void gtpinNotifyMakeResident(void *pKernel, void *pCommandStreamReceiver);
 void gtpinNotifyUpdateResidencyList(void *pKernel, void *pResidencyVector);
 void gtpinNotifyPlatformShutdown();
-bool gtpinIsGTPinInitialized();
+inline bool gtpinIsGTPinInitialized() { return isGTPinInitialized; }
 }
diff --git a/unit_tests/gtpin/gtpin_tests.cpp b/unit_tests/gtpin/gtpin_tests.cpp
index e986ca5408..2ef561c90b 100644
--- a/unit_tests/gtpin/gtpin_tests.cpp
+++ b/unit_tests/gtpin/gtpin_tests.cpp
@@ -49,7 +49,6 @@ using namespace OCLRT;
 using namespace gtpin;
 
 namespace OCLRT {
-extern bool isGTPinInitialized;
 extern std::deque<gtpinkexec_t> kernelExecQueue;
 }
 
@@ -2073,4 +2072,17 @@ TEST_F(GTPinTests, checkWhetherGTPinHwHelperGetterWorksWell) {
     EXPECT_NE(nullptr, pGTPinHelper);
 }
 
+TEST(GTPinOfflineTests, givenGtPinInDisabledStateWhenCallbacksFromEnqueuePathAreCalledThenNothingHappens) {
+    ASSERT_FALSE(gtpinIsGTPinInitialized());
+    auto dummyKernel = reinterpret_cast<cl_kernel>(0x1000);
+    auto dummyQueue = reinterpret_cast<void *>(0x1000);
+    uint32_t dummyCompletedTask = 0u;
+
+    //now call gtpin function with dummy data, this must not crash
+    gtpinNotifyKernelSubmit(dummyKernel, dummyQueue);
+    gtpinNotifyPreFlushTask(dummyQueue);
+    gtpinNotifyTaskCompletion(dummyCompletedTask);
+    gtpinNotifyFlushTask(dummyCompletedTask);
+    EXPECT_FALSE(gtpinIsGTPinInitialized());
+}
 } // namespace ULT