Unused bindless image args treated as bindless fix.

When emitting zeinfo IGC tags addr mode of images with no users as
stateful even if the module is compiled to use bindless images. This
caused NEO to throw an error as it disallows the use of both bindless
and bindful mode in the same module.

This commit sets the default addr mode to bindless for modules that have
UseBindlessImage set to true.

Author: smilczek

IGC/Compiler/CISACodeGen/OpenCLKernelCodeGen.cpp

@@ -771,19 +771,20 @@ bool COpenCLKernel::CreateZEPayloadArguments(IGC::KernelArg *kernelArg, uint pay
   case KernelArg::ArgType::BINDLESS_IMAGE_CUBE_ARRAY:
   case KernelArg::ArgType::IMAGE_CUBE_DEPTH_ARRAY:
   case KernelArg::ArgType::BINDLESS_IMAGE_CUBE_DEPTH_ARRAY: {
-    // the image arg is either bindless or stateful. check from "kernelArg->needsAllocation()"
+    // the image arg is either bindless or stateful.
     // For stateful image argument, the arg has 0 offset and 0 size
-    zebin::PreDefinedAttrGetter::ArgAddrMode arg_addrmode = zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
+    zebin::PreDefinedAttrGetter::ArgAddrMode arg_addrmode = m_ModuleMetadata->UseBindlessImage
+                                                                ? zebin::PreDefinedAttrGetter::ArgAddrMode::bindless
+                                                                : zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
     uint arg_off = 0;
     uint arg_size = 0;
 
     int arg_idx = kernelArg->getAssociatedArgNo();
     if (kernelArg->needsAllocation()) {
-      // set to bindless
-      arg_addrmode = zebin::PreDefinedAttrGetter::ArgAddrMode::bindless;
+      // bindless
       arg_off = payloadPosition;
       arg_size = kernelArg->getSize();
-    } else {
+    } else if (arg_addrmode == zebin::PreDefinedAttrGetter::ArgAddrMode::stateful) {
       // add bti index for this arg if it's stateful
       SOpenCLKernelInfo::SResourceInfo resInfo = getResourceInfo(arg_idx);
       zebin::ZEInfoBuilder::addBindingTableIndex(m_kernelInfo.m_zeBTIArgs, getBTI(resInfo), arg_idx);
@@ -904,15 +905,16 @@ bool COpenCLKernel::CreateZEPayloadArguments(IGC::KernelArg *kernelArg, uint pay
   // sampler
   case KernelArg::ArgType::SAMPLER:
   case KernelArg::ArgType::BINDLESS_SAMPLER: {
-    // the sampler arg is either bindless or stateful. check from "kernelArg->needsAllocation()"
+    // the sampler arg is either bindless or stateful.
     // For stateful image argument, the arg has 0 offset and 0 size
     // NOTE: we only have stateful sampler now
-    zebin::PreDefinedAttrGetter::ArgAddrMode arg_addrmode = zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
+    zebin::PreDefinedAttrGetter::ArgAddrMode arg_addrmode = m_ModuleMetadata->UseBindlessImage
+                                                                ? zebin::PreDefinedAttrGetter::ArgAddrMode::bindless
+                                                                : zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
     uint arg_off = 0;
     uint arg_size = 0;
     if (kernelArg->needsAllocation()) {
-      // set to bindless
-      arg_addrmode = zebin::PreDefinedAttrGetter::ArgAddrMode::bindless;
+      // bindless
       arg_off = payloadPosition;
       arg_size = kernelArg->getSize();
     }

IGC/ocloc_tests/ZEBinary/unused-image-zeinfo.cl

@@ -0,0 +1,31 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2025 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// UNSUPPORTED: system-windows
+// RUN: ocloc compile -file %s -options "-igc_opts 'DumpZEInfoToConsole=1'" -internal_options "-cl-intel-use-bindless-mode" -device mtl | FileCheck %s
+
+// Check that unused image args arent tagged with 'stateful' addrmode in bindless mode.
+
+// CHECK-NOT: addrmode:        stateful
+// CHECK-NOT: binding_table_indices:
+
+const sampler_t sampler =
+    CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
+
+kernel void test(const global float *a,
+                 global float *c,
+                 read_only image2d_t input,
+                 read_only image2d_t unused,
+                 sampler_t sampler
+                 ) {
+    const int gid = get_global_id(0);
+
+    int2 coord = {get_global_id(0), get_global_id(1)};
+    float4 data = read_imagef(input, coord);
+    c[gid] = a[gid] + data.x;
+}