Pass gtpinInfo to Igc during link

Change-Id: Ie05ea252d59372dcfb766851efa1311642f29c74
Signed-off-by: Maciej Plewka <maciej.plewka@intel.com>

opencl/source/gtpin/gtpin_callbacks.cpp

@@ -231,4 +231,8 @@ void gtpinNotifyPlatformShutdown() {
 void *gtpinGetIgcInit() {
     return pIgcInit;
 }
+void gtpinSetIgcInit(void *pIgcInitPtr) {
+    pIgcInit = static_cast<igc_init_t *>(pIgcInitPtr);
+}
+
 } // namespace NEO

opencl/source/gtpin/gtpin_notify.h

@@ -26,4 +26,5 @@ void gtpinNotifyUpdateResidencyList(void *pKernel, void *pResidencyVector);
 void gtpinNotifyPlatformShutdown();
 inline bool gtpinIsGTPinInitialized() { return isGTPinInitialized; }
 void *gtpinGetIgcInit();
+void gtpinSetIgcInit(void *pIgcInitPtr);
 } // namespace NEO

opencl/source/program/link.cpp

@@ -15,6 +15,7 @@
 #include "shared/source/utilities/stackvec.h"
 
 #include "opencl/source/device/cl_device.h"
+#include "opencl/source/gtpin/gtpin_notify.h"
 #include "opencl/source/helpers/validators.h"
 #include "opencl/source/platform/platform.h"
 #include "opencl/source/program/kernel_info.h"
@@ -117,6 +118,7 @@ cl_int Program::link(
         inputArgs.src = ArrayRef<const char>(reinterpret_cast<const char *>(clLinkInput.data()), clLinkInput.size());
         inputArgs.apiOptions = ArrayRef<const char>(options.c_str(), options.length());
         inputArgs.internalOptions = ArrayRef<const char>(internalOptions.c_str(), internalOptions.length());
+        inputArgs.GTPinInput = gtpinGetIgcInit();
 
         if (!isCreateLibrary) {
             inputArgs.outType = IGC::CodeType::oclGenBin;

opencl/test/unit_test/program/program_tests.cpp

@@ -26,6 +26,7 @@
 #include "shared/test/unit_test/helpers/debug_manager_state_restore.h"
 #include "shared/test/unit_test/utilities/base_object_utils.h"
 
+#include "opencl/source/gtpin/gtpin_notify.h"
 #include "opencl/source/helpers/hardware_commands_helper.h"
 #include "opencl/source/kernel/kernel.h"
 #include "opencl/source/program/create.inl"
@@ -2921,3 +2922,39 @@ TEST_F(ProgramMultiRootDeviceTests, privateSurfaceHasCorrectRootDeviceIndex) {
     EXPECT_NE(nullptr, privateSurface);
     EXPECT_EQ(expectedRootDeviceIndex, privateSurface->getRootDeviceIndex());
 }
+
+class MockCompilerInterfaceWithGtpinParam : public CompilerInterface {
+  public:
+    TranslationOutput::ErrorCode link(
+        const NEO::Device &device,
+        const TranslationInput &input,
+        TranslationOutput &output) override {
+        gtpinInfoPassed = input.GTPinInput;
+        return CompilerInterface::link(device, input, output);
+    }
+    void *gtpinInfoPassed;
+};
+
+TEST_F(ProgramBinTest, GivenSourceKernelWhenLinkingProgramThengtpinInitInfoIsPassed) {
+    cl_device_id device = pClDevice;
+    void *pIgcInitPtr = reinterpret_cast<void *>(0x1234);
+    gtpinSetIgcInit(pIgcInitPtr);
+    const char *sourceCode = "__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n";
+    pProgram = Program::create<MockProgram>(
+        pContext,
+        1,
+        &sourceCode,
+        &knownSourceSize,
+        retVal);
+    std::unique_ptr<MockCompilerInterfaceWithGtpinParam> mockCompilerInterface(new MockCompilerInterfaceWithGtpinParam);
+
+    retVal = pProgram->compile(1, &device, nullptr, 0, nullptr, nullptr, nullptr, nullptr);
+    EXPECT_EQ(CL_SUCCESS, retVal);
+    pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(mockCompilerInterface.get());
+
+    cl_program programToLink = pProgram;
+    retVal = pProgram->link(1, &device, nullptr, 1, &programToLink, nullptr, nullptr);
+
+    EXPECT_EQ(pIgcInitPtr, mockCompilerInterface->gtpinInfoPassed);
+    mockCompilerInterface.release();
+}

shared/source/compiler_interface/compiler_interface.cpp

@@ -225,7 +225,7 @@ TranslationOutput::ErrorCode CompilerInterface::link(
 
         auto igcTranslationCtx = createIgcTranslationCtx(device, inType, outType);
         currOut = translate(igcTranslationCtx.get(), currSrc.get(),
-                            igcOptions.get(), igcInternalOptions.get());
+                            igcOptions.get(), igcInternalOptions.get(), input.GTPinInput);
 
         if (currOut == nullptr) {
             return TranslationOutput::ErrorCode::UnknownError;