Unused bindless image args treated as bindful 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

@@ -1151,7 +1151,9 @@ namespace IGC
             // the image arg is either bindless or stateful. check from "kernelArg->needsAllocation()"
             // For stateful image argument, the arg has 0 offset and 0 size
             zebin::PreDefinedAttrGetter::ArgAddrMode arg_addrmode =
-                zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
+                m_ModuleMetadata->UseBindlessImage ?
+                    zebin::PreDefinedAttrGetter::ArgAddrMode::bindless :
+                    zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
             uint arg_off = 0;
             uint arg_size = 0;
 
@@ -1302,7 +1304,9 @@ namespace IGC
             // 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;
+                m_ModuleMetadata->UseBindlessImage ?
+                    zebin::PreDefinedAttrGetter::ArgAddrMode::bindless :
+                    zebin::PreDefinedAttrGetter::ArgAddrMode::stateful;
             uint arg_off = 0;
             uint arg_size = 0;
             if (kernelArg->needsAllocation()) {

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

@@ -0,0 +1,30 @@
+/*========================== 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
+
+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;
+}