Support bindless sampler kernel argument and SYCL bindless sampler

* For bindless sampler passed as kernel argument, e.g. set by OpenCL
  setKernelArg or L0 zeKernelSetArgumentValue API, bindless sampler
  offset is obtained from kernel argument directly.
* For SYCL bindless sampled image, it is created using a single
  zeImageCreate API call from image and sampler descriptors. NEO sets up
  both bindless image state and bindless sampler state of an image in a
  global bindless surface heap (UseExternalAllocatorForSshAndDsh=1).
  Sampler state is placed after image state and its implicits args
  state. The offset between sampler state and image state is 128 bytes.
* In both of above cases, bindless sampler offset is OR'ed with 1, which
  sets `Sampler State Base Address Select` to use Bindless Sampler State
  Base Address.

Author: wenju-he

IGC/Compiler/Optimizer/OCLBIUtils.cpp

@@ -233,7 +233,7 @@ Argument* CImagesBI::CImagesUtils::findImageFromBufferPtr(const MetaDataUtils& M
     return nullptr;
 }
 
-static bool isBindlessImageOrSamplerLoad(Value *v)
+static bool isBindlessImageLoad(Value *v)
 {
     auto *load = dyn_cast<LoadInst>(v);
     if (!load)
@@ -254,7 +254,7 @@ ConstantInt* CImagesBI::CImagesUtils::getImageIndex(
 {
     ConstantInt* imageIndex = nullptr;
 
-    imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageOrSamplerLoad);
+    imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageLoad);
     IGC_ASSERT(imageParam);
     IGC_ASSERT(isa<Argument>(imageParam) || isa<LoadInst>(imageParam));
     int i = (*pParamMap)[imageParam].index;
@@ -264,7 +264,7 @@ ConstantInt* CImagesBI::CImagesUtils::getImageIndex(
 
 BufferType CImagesBI::CImagesUtils::getImageType(ParamMap* pParamMap, CallInst* pCallInst, unsigned int paramIndex)
 {
-    Value *imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageOrSamplerLoad);
+    Value *imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageLoad);
     IGC_ASSERT(imageParam);
     IGC_ASSERT(isa<Argument>(imageParam) || isa<LoadInst>(imageParam));
     return isa<LoadInst>(imageParam) ? BufferType::BINDLESS : (*pParamMap)[imageParam].type;
@@ -409,7 +409,21 @@ class COCL_sample : public CImagesBI
     Value* getSamplerValue(void)
     {
         ConstantInt* samplerIndex = nullptr;
-        Value* samplerParam = ValueTracker::track(m_pCallInst, 1, m_pMdUtils, m_modMD, isBindlessImageOrSamplerLoad);
+        auto isBindlessSampler = [](Value *v)
+        {
+            // Bindless sampler is computed in ResolveSampledImageBuiltins pass.
+            if (auto *I = dyn_cast<BinaryOperator>(v))
+            {
+                if (I->getOpcode() != BinaryOperator::Or || !I->getType()->isIntegerTy(64))
+                    return false;
+                if (auto *C = dyn_cast<ConstantInt>(I->getOperand(1)))
+                {
+                    return C->isOne();
+                }
+            }
+            return false;
+        };
+        Value* samplerParam = ValueTracker::track(m_pCallInst, 1, m_pMdUtils, m_modMD, isBindlessSampler);
         if (!samplerParam) {
             emitError("There are instructions that use a sampler, but no sampler found in the kernel!", m_pCallInst);
             return nullptr;

IGC/Compiler/Optimizer/OpenCLPasses/ImageFuncs/ResolveSampledImageBuiltins.cpp

@@ -8,6 +8,7 @@ SPDX-License-Identifier: MIT
 
 #include "Compiler/Optimizer/OpenCLPasses/ImageFuncs/ResolveSampledImageBuiltins.hpp"
 #include "Compiler/IGCPassSupport.h"
+#include "common/MDFrameWork.h"
 #include "common/LLVMWarningsPush.hpp"
 #include <llvm/IR/Function.h>
 #include <llvm/IR/Instructions.h>
@@ -24,6 +25,7 @@ using namespace IGC;
 #define PASS_CFG_ONLY false
 #define PASS_ANALYSIS false
 IGC_INITIALIZE_PASS_BEGIN(ResolveSampledImageBuiltins, PASS_FLAG, PASS_DESCRIPTION, PASS_CFG_ONLY, PASS_ANALYSIS)
+IGC_INITIALIZE_PASS_DEPENDENCY(MetaDataUtilsWrapper)
 IGC_INITIALIZE_PASS_END(ResolveSampledImageBuiltins, PASS_FLAG, PASS_DESCRIPTION, PASS_CFG_ONLY, PASS_ANALYSIS)
 
 char ResolveSampledImageBuiltins::ID = 0;
@@ -38,6 +40,7 @@ ResolveSampledImageBuiltins::ResolveSampledImageBuiltins() : ModulePass(ID)
 
 bool ResolveSampledImageBuiltins::runOnModule(Module& M) {
     m_changed = false;
+    modMD = getAnalysis<MetaDataUtilsWrapper>().getModuleMetaData();
     visit(M);
 
     for (auto builtin : m_builtinsToRemove)
@@ -133,6 +136,7 @@ Value* ResolveSampledImageBuiltins::lowerGetSampler(CallInst& CI)
     IGC_ASSERT(callReturningOpaque);
 
     m_builtinsToRemove.insert(callReturningOpaque);
+    auto *Int64Ty = Type::getInt64Ty(CI.getContext());
 
     Value* samplerArg = callReturningOpaque->getArgOperand(1);
     if (CallInst* samplerInitializer = dyn_cast<CallInst>(samplerArg))
@@ -141,18 +145,45 @@ Value* ResolveSampledImageBuiltins::lowerGetSampler(CallInst& CI)
         return ZExtInst::Create(
             Instruction::ZExt,
             samplerInitializer->getArgOperand(0),
-            Type::getInt64Ty(CI.getContext()),
+            Int64Ty,
             "",
             &CI);
     }
+    else if (modMD->UseBindlessImage && !isa<Argument>(samplerArg))
+    {
+        Value* image = callReturningOpaque->getArgOperand(0);
+        IGC_ASSERT(image->getType()->isPointerTy());
+        Value *imageOffset = PtrToIntInst::Create(
+            Instruction::PtrToInt,
+            image,
+            Int64Ty,
+            "",
+            &CI);
+        // When sampled image is created in a single API call, e.g. SYCL bindless image,
+        // bindless surface state heap layout is
+        // | image state | image implicit args state | sampler state | redescribed image state | ...
+        // Sampler state offset is addition of image state offset, size of
+        // image state and size of image implicit args state.
+        // Both size of image state and image implicit args state are 64 bytes.
+        constexpr uint64_t surfaceStateSize = 64;
+        auto *stateSizeValue = ConstantInt::get(Int64Ty, surfaceStateSize * 2);
+        auto *samplerOffset = BinaryOperator::CreateAdd(imageOffset, stateSizeValue, "sampler_offset", &CI);
+        // Set bit-field 0 to 1 to select Bindless Sampler State Base Address.
+        return BinaryOperator::CreateOr(samplerOffset, ConstantInt::get(Int64Ty, 1), "", &CI);
+    }
     else
     {
         IGC_ASSERT(samplerArg->getType()->isPointerTy());
-        return PtrToIntInst::Create(
+        Value *samplerOffset = PtrToIntInst::Create(
             Instruction::PtrToInt,
             samplerArg,
-            Type::getInt64Ty(CI.getContext()),
+            Int64Ty,
             "",
             &CI);
+        if (modMD->UseBindlessImage) {
+            // Set bit-field 0 to 1 to select Bindless Sampler State Base Address.
+            samplerOffset = BinaryOperator::CreateOr(samplerOffset, ConstantInt::get(Int64Ty, 1), "", &CI);
+        }
+        return samplerOffset;
     }
 }

IGC/Compiler/Optimizer/OpenCLPasses/ImageFuncs/ResolveSampledImageBuiltins.hpp

@@ -8,6 +8,7 @@ SPDX-License-Identifier: MIT
 
 #pragma once
 
+#include "Compiler/MetaDataUtilsWrapper.h"
 #include "common/LLVMWarningsPush.hpp"
 #include <llvm/Pass.h>
 #include <llvm/IR/InstVisitor.h>
@@ -34,6 +35,11 @@ namespace IGC
             return "ResolveSampledImageBuiltins";
         }
 
+        virtual void getAnalysisUsage(llvm::AnalysisUsage &AU) const override
+        {
+            AU.addRequired<MetaDataUtilsWrapper>();
+        }
+
         virtual bool runOnModule(llvm::Module& M) override;
         void visitCallInst(llvm::CallInst& CI);
 
@@ -45,6 +51,7 @@ namespace IGC
         llvm::Value* lowerGetSampler(llvm::CallInst& CI);
 
         bool m_changed = false;
+        ModuleMetaData* modMD = nullptr;
         std::unordered_set<llvm::CallInst*> m_builtinsToRemove;
     };
 

IGC/Compiler/tests/OCLBIConverter/sampled_image_handle_2d.ll

@@ -10,46 +10,40 @@
 
 ; RUN: igc_opt %s -S -o - -igc-conv-ocl-to-common | FileCheck %s
 
-%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle" = type { %"struct.sycl::_V1::ext::oneapi::experimental::combined_sampled_image_handle" }
-%"struct.sycl::_V1::ext::oneapi::experimental::combined_sampled_image_handle" = type { i64, i64 }
+%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle" = type { i64 }
 %spirv.Image._void_1_0_0_0_0_0_0 = type opaque
-%spirv.Sampler = type opaque
 
 define spir_kernel void @_ZTS14image_addition(%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle1, %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle2) {
 entry:
   %0 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle1 to i64*
   %__SYCLKernel.1.copyload = load i64, i64* %0, align 8
-  %1 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle1 to i8*
-  %__SYCLKernel.imgHandle1.sroa_idx = getelementptr inbounds i8, i8* %1, i64 8
-  %2 = bitcast i8* %__SYCLKernel.imgHandle1.sroa_idx to i64*
-  %__SYCLKernel.2.copyload = load i64, i64* %2, align 8
-  %3 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle2 to i64*
-  %__SYCLKernel.3.copyload = load i64, i64* %3, align 8
-  %4 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle2 to i8*
-  %__SYCLKernel.imgHandle2.sroa_idx = getelementptr inbounds i8, i8* %4, i64 8
-  %5 = bitcast i8* %__SYCLKernel.imgHandle2.sroa_idx to i64*
-  %__SYCLKernel.4.copyload = load i64, i64* %5, align 8
+  %1 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle2 to i64*
+  %__SYCLKernel.3.copyload = load i64, i64* %1, align 8
   %astype = inttoptr i64 %__SYCLKernel.1.copyload to %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*
-  %astype2 = inttoptr i64 %__SYCLKernel.2.copyload to %spirv.Sampler addrspace(2)*
-  %6 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype to i64
-  %7 = ptrtoint %spirv.Sampler addrspace(2)* %astype2 to i64
-  %conv = trunc i64 %6 to i32
-  %conv2 = trunc i64 %7 to i32
+  %2 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype to i64
+  %3 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype to i64
+  %conv = trunc i64 %2 to i32
+  %sampler_offset = add i64 %3, 128
+  %4 = or i64 %sampler_offset, 1
+  %conv2 = trunc i64 %4 to i32
 
+; CHECK: [[OR:%.*]] = or i64 %sampler_offset, 1
 ; CHECK: %bindless_img = inttoptr i64 %__SYCLKernel.1.copyload to float addrspace(393216)*
-; CHECK-NEXT: %bindless_sampler = inttoptr i64 %__SYCLKernel.2.copyload to float addrspace(655360)*
+; CHECK-NEXT: %bindless_sampler = inttoptr i64 [[OR]] to float addrspace(655360)*
 ; CHECK-NEXT: %call1 = call <4 x float> @llvm.genx.GenISA.sampleLptr.v4f32.f32.p196608f32.p393216f32.p655360f32(float 0.000000e+00, float %CoordX, float %CoordY, float 0.000000e+00, float 0.000000e+00, float addrspace(196608)* undef, float addrspace(393216)* %bindless_img, float addrspace(655360)* %bindless_sampler, i32 0, i32 0, i32 0)
 
   %call = call spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32 %conv, i32 %conv2, <2 x float> zeroinitializer, float 0.000000e+00)
   %astype3 = inttoptr i64 %__SYCLKernel.3.copyload to %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*
-  %astype4 = inttoptr i64 %__SYCLKernel.4.copyload to %spirv.Sampler addrspace(2)*
-  %8 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype3 to i64
-  %9 = ptrtoint %spirv.Sampler addrspace(2)* %astype4 to i64
-  %conv3 = trunc i64 %8 to i32
-  %conv4 = trunc i64 %9 to i32
+  %5 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype3 to i64
+  %6 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype3 to i64
+  %conv3 = trunc i64 %5 to i32
+  %sampler_offset23 = add i64 %6, 128
+  %7 = or i64 %sampler_offset23, 1
+  %conv4 = trunc i64 %7 to i32
 
+; CHECK: [[OR2:%.*]] = or i64 %sampler_offset{{.*}}, 1
 ; CHECK: [[IMG2:%bindless_img[0-9]+]] = inttoptr i64 %__SYCLKernel.3.copyload to float addrspace(393216)*
-; CHECK-NEXT: [[SAMPLER2:%bindless_sampler[0-9]+]] = inttoptr i64 %__SYCLKernel.4.copyload to float addrspace(655360)*
+; CHECK-NEXT: [[SAMPLER2:%bindless_sampler[0-9]+]] = inttoptr i64 [[OR2]] to float addrspace(655360)*
 ; CHECK-NEXT: %call26 = call <4 x float> @llvm.genx.GenISA.sampleLptr.v4f32.f32.p196608f32.p393216f32.p655360f32(float 0.000000e+00, float %CoordX2, float %CoordY3, float 0.000000e+00, float 0.000000e+00, float addrspace(196608)* undef, float addrspace(393216)* [[IMG2]], float addrspace(655360)* [[SAMPLER2]], i32 0, i32 0, i32 0)
 
   %call2 = call spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32 %conv3, i32 %conv4, <2 x float> zeroinitializer, float 0.000000e+00)

IGC/Compiler/tests/ResolveSampledImageBuiltins/bindless-sampler-argument.ll

@@ -0,0 +1,46 @@
+;=========================== begin_copyright_notice ============================
+;
+; Copyright (C) 2024 Intel Corporation
+;
+; SPDX-License-Identifier: MIT
+;
+;============================ end_copyright_notice =============================
+
+; RUN: igc_opt -igc-image-sampler-resolution %s -S -o - | FileCheck %s
+
+; Check offset of bindless sampler as kernel argument is computed as: ptrtoint(%sampler) | 1
+
+%spirv.Image._void_1_0_0_0_0_0_0 = type opaque
+%spirv.Sampler = type opaque
+%spirv.SampledImage._void_1_0_0_0_0_0_0 = type opaque
+
+define spir_kernel void @image_read_sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %img, %spirv.Sampler addrspace(2)* %sampler) {
+entry:
+; CHECK: [[INT:%.*]] = ptrtoint %spirv.Sampler addrspace(2)* %sampler to i64
+; CHECK-NEXT: [[OR:%.*]] = or i64 [[INT]], 1
+; CHECK-NEXT: [[SAMPLER:%.*]] = trunc i64 [[OR]] to i32
+; CHECK: call spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32 {{.*}}, i32 [[SAMPLER]], <2 x float>
+
+  %TempSampledImage = call spir_func %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS133__spirv_Image__void_1_0_0_0_0_0_0PU3AS215__spirv_Sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %img, %spirv.Sampler addrspace(2)* %sampler)
+  %0 = bitcast %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* %TempSampledImage to i8 addrspace(1)*
+  %call1 = call spir_func i64 @__builtin_IB_get_sampler(i8 addrspace(1)* %0)
+  %conv2 = trunc i64 %call1 to i32
+  %call2 = call spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32 0, i32 %conv2, <2 x float> zeroinitializer, float 0.000000e+00)
+  ret void
+}
+
+declare spir_func %spirv.SampledImage._void_1_0_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS133__spirv_Image__void_1_0_0_0_0_0_0PU3AS215__spirv_Sampler(%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)
+
+declare spir_func i64 @__builtin_IB_get_sampler(i8 addrspace(1)*)
+
+declare spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32, i32, <2 x float>, float)
+
+!IGCMetadata = !{!0}
+!igc.functions = !{}
+
+!0 = !{!"ModuleMD", !1, !5}
+!1 = !{!"FuncMD", !2, !3}
+!2 = !{!"FuncMDMap[0]", void (%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @image_read_sampler}
+!3 = !{!"FuncMDValue[0]", !4}
+!4 = !{!"localOffsets"}
+!5 = !{!"UseBindlessImage", i1 true}

IGC/Compiler/tests/ResolveSampledImageBuiltins/sycl-bindless-sampler.ll

@@ -0,0 +1,60 @@
+;=========================== begin_copyright_notice ============================
+;
+; Copyright (C) 2024 Intel Corporation
+;
+; SPDX-License-Identifier: MIT
+;
+;============================ end_copyright_notice =============================
+
+; RUN: igc_opt -igc-image-sampler-resolution %s -S -o - | FileCheck %s
+
+; Check SYCL bindless sampler offset is computed from bindless image offset: (BindlessImageOffset + 128) | 1
+
+%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle" = type { i64 }
+%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
+%"class.sycl::_V1::detail::array" = type { [1 x i64] }
+%spirv.Image._void_0_0_0_0_0_0_0 = type opaque
+%spirv.SampledImage._void_0_0_0_0_0_0_0 = type opaque
+%spirv.Sampler = type opaque
+
+define spir_kernel void @_ZTS14image_addition(%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %_arg_imgHandle, float addrspace(1)* %_arg_outAcc, %"class.sycl::_V1::id"* %_arg_outAcc3) {
+entry:
+; CHECK: [[BC:%.*]] = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %_arg_imgHandle to i64*
+; CHECK: [[LOAD:%.*]] = load i64, i64* [[BC]], align 8
+; CHECK: [[PTR:%.*]] = inttoptr i64 [[LOAD]] to %spirv.Image._void_0_0_0_0_0_0_0 addrspace(1)*
+; CHECK: [[INT1:%.*]] = ptrtoint %spirv.Image._void_0_0_0_0_0_0_0 addrspace(1)* [[PTR]] to i64
+; CHECK: [[INT2:%.*]] = ptrtoint %spirv.Image._void_0_0_0_0_0_0_0 addrspace(1)* [[PTR]] to i64
+; CHECK-NEXT: [[ADD:%.*]] = add i64 [[INT2]], 128
+; CHECK-NEXT: [[OR:%.*]] = or i64 [[ADD]], 1
+; CHECK-NEXT: [[IMG:%.*]] = trunc i64 [[INT1]] to i32
+; CHECK-NEXT: [[SAMPLER:%.*]] = trunc i64 [[OR]] to i32
+; CHECK: call spir_func <4 x float> @__builtin_IB_OCL_1d_sample_l(i32 [[IMG]], i32 [[SAMPLER]], float
+
+  %0 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %_arg_imgHandle to i64*
+  %__SYCLKernel.sroa.0.0.copyload = load i64, i64* %0, align 8
+  %astype.i = inttoptr i64 %__SYCLKernel.sroa.0.0.copyload to %spirv.Image._void_0_0_0_0_0_0_0 addrspace(1)*
+  %call3.i.i = call spir_func %spirv.SampledImage._void_0_0_0_0_0_0_0 addrspace(1)* undef(%spirv.Image._void_0_0_0_0_0_0_0 addrspace(1)* %astype.i, %spirv.Sampler addrspace(2)* null)
+  %1 = bitcast %spirv.SampledImage._void_0_0_0_0_0_0_0 addrspace(1)* %call3.i.i to i8 addrspace(1)*
+  %call.i.i.i5 = call spir_func i64 @__builtin_IB_get_image(i8 addrspace(1)* %1)
+  %call1.i.i.i = call spir_func i64 @__builtin_IB_get_sampler(i8 addrspace(1)* %1)
+  %conv.i.i.i = trunc i64 %call.i.i.i5 to i32
+  %conv2.i.i.i = trunc i64 %call1.i.i.i to i32
+  %call7.i.i.i = call spir_func <4 x float> @__builtin_IB_OCL_1d_sample_l(i32 %conv.i.i.i, i32 %conv2.i.i.i, float 0.000000e+00, float 0.000000e+00)
+  ret void
+}
+
+declare spir_func i64 @__builtin_IB_get_image(i8 addrspace(1)*)
+
+declare spir_func i64 @__builtin_IB_get_sampler(i8 addrspace(1)*)
+
+declare spir_func <4 x float> @__builtin_IB_OCL_1d_sample_l(i32, i32, float, float)
+
+!IGCMetadata = !{!0}
+!igc.functions = !{}
+
+!0 = !{!"ModuleMD", !1, !5}
+!1 = !{!"FuncMD", !2, !3}
+!2 = !{!"FuncMDMap[0]", void (%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"*, float addrspace(1)*, %"class.sycl::_V1::id"*)* @_ZTS14image_addition}
+!3 = !{!"FuncMDValue[0]", !4}
+!4 = !{!"localOffsets"}
+!5 = !{!"UseBindlessImage", i1 true}