Enable has_printf_calls zeinfo field

This commit enables previously added has_printf_calls field and test for it.

Author: KanclerzPiotr

IGC/AdaptorOCL/OCL/sp/zebin_builder.cpp

@@ -471,8 +471,8 @@ void ZEBinaryBuilder::addKernelExecEnv(const SOpenCLKernelInfo &annotations, zeI
       CPlatform(mPlatform).hasScratchSurface() && IGC_IS_FLAG_ENABLED(SeparateSpillPvtScratchSpace);
   env.has_no_stateless_write = (annotations.m_executionEnvironment.StatelessWritesCount == 0);
   env.has_stack_calls = annotations.m_executionEnvironment.HasStackCalls;
-  // env.has_printf_calls = annotations.m_executionEnvironment.HasPrintfCalls;
-  // env.has_indirect_calls = annotations.m_executionEnvironment.HasIndirectCalls;
+  env.has_printf_calls = annotations.m_executionEnvironment.HasPrintfCalls;
+  env.has_indirect_calls = annotations.m_executionEnvironment.HasIndirectCalls;
   env.require_disable_eufusion = annotations.m_executionEnvironment.RequireDisableEUFusion;
   env.indirect_stateless_count = annotations.m_executionEnvironment.IndirectStatelessCount;
   env.inline_data_payload_size = annotations.m_threadPayload.PassInlineDataSize;
@@ -516,8 +516,8 @@ void ZEBinaryBuilder::addFunctionExecEnv(const SOpenCLKernelInfo &annotations,
   env.simd_size = annotations.m_executionEnvironment.CompiledSIMDSize;
   env.barrier_count = zeFuncAttr.f_BarrierCount;
   env.has_rtcalls = zeFuncAttr.f_hasRTCalls;
-  // env.has_printf_calls = zeFuncAttr.f_hasPrintfCalls;
-  // env.has_indirect_calls = zeFuncAttr.f_hasIndirectCalls;
+  env.has_printf_calls = zeFuncAttr.f_hasPrintfCalls;
+  env.has_indirect_calls = zeFuncAttr.f_hasIndirectCalls;
 }
 
 void ZEBinaryBuilder::addLocalIds(uint32_t simdSize, uint32_t grfSize, bool has_local_id_x, bool has_local_id_y,

IGC/ocloc_tests/ZEBinary/has_printf.cl

@@ -0,0 +1,59 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2025 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// The test checks if zeinfo "has_printf_calls" and "has_indirect_calls" and it fields are set.
+
+// UNSUPPORTED: sys32
+// REQUIRES: pvc-supported, regkeys
+
+// RUN: ocloc compile -file %s -options "-igc_opts 'DumpZEInfoToConsole=1'" \
+// RUN:     -device pvc | FileCheck %s
+
+// CHECK:       kernels:
+// CHECK-NEXT:  - name:            test
+// CHECK:       execution_env:
+// CHECK:       has_indirect_calls: true
+
+// CHECK:       functions:
+// CHECK:       - name:            even
+// CHECK:       execution_env
+// CHECK:       has_printf_calls: true
+// CHECK:       - name:            odd
+// CHECK:       execution_env
+// CHECK:       has_printf_calls: true
+
+char *__builtin_IB_get_function_pointer(__constant char *function_name);
+void __builtin_IB_call_function_pointer(char *function_pointer,
+                                        char *argument_structure);
+
+void even(char *argument_structure) {
+  int *pio = (int *)argument_structure;
+  printf("%d", *pio * 2);
+}
+
+void odd(char *argument_structure) {
+  int *pio = (int *)argument_structure;
+  printf("%d", *pio * 3);
+}
+
+#define EVEN_FUNC_NAME "even"
+#define ODD_FUNC_NAME "odd"
+
+__kernel void test() {
+  int gid = (int)get_global_id(0);
+
+  char* fp = 0;
+  if(gid % 2)
+    fp = __builtin_IB_get_function_pointer(ODD_FUNC_NAME);
+  else
+    fp = __builtin_IB_get_function_pointer(EVEN_FUNC_NAME);
+
+  int x = 5;
+
+  __builtin_IB_call_function_pointer(fp, (char *)&x);
+}