Skip to content

Commit f2034a4

Browse files
wenju-heigcbot
wenju-he
authored and
igcbot
committed
[Bindless Image] Pass bindless sampler directly to sampleLptr
Bindless sampler handle should be directly passed to sampleLptr intrinsic, same as bindless image. This PR handles following cases: * sampler is passed as kernel argument in OpenCL * sampler is passed either in sampled image handle or through memory in SYCL bindless image. OpenCL Inline sampler isn't handled yet.
1 parent 7daf84c commit f2034a4

File tree

4 files changed

+168
-16
lines changed

4 files changed

+168
-16
lines changed

‎IGC/Compiler/Optimizer/OCLBIUtils.cpp‎

Lines changed: 34 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -233,7 +233,7 @@ Argument* CImagesBI::CImagesUtils::findImageFromBufferPtr(const MetaDataUtils& M
233233
return nullptr;
234234
}
235235

236-
static bool isBindlessImageLoad(Value *v)
236+
static bool isBindlessImageOrSamplerLoad(Value *v)
237237
{
238238
auto *load = dyn_cast<LoadInst>(v);
239239
if (!load)
@@ -254,7 +254,7 @@ ConstantInt* CImagesBI::CImagesUtils::getImageIndex(
254254
{
255255
ConstantInt* imageIndex = nullptr;
256256

257-
imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageLoad);
257+
imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageOrSamplerLoad);
258258
IGC_ASSERT(imageParam);
259259
IGC_ASSERT(isa<Argument>(imageParam) || isa<LoadInst>(imageParam));
260260
int i = (*pParamMap)[imageParam].index;
@@ -264,7 +264,7 @@ ConstantInt* CImagesBI::CImagesUtils::getImageIndex(
264264

265265
BufferType CImagesBI::CImagesUtils::getImageType(ParamMap* pParamMap, CallInst* pCallInst, unsigned int paramIndex)
266266
{
267-
Value *imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageLoad);
267+
Value *imageParam = ValueTracker::track(pCallInst, paramIndex, nullptr, nullptr, isBindlessImageOrSamplerLoad);
268268
IGC_ASSERT(imageParam);
269269
IGC_ASSERT(isa<Argument>(imageParam) || isa<LoadInst>(imageParam));
270270
return isa<LoadInst>(imageParam) ? BufferType::BINDLESS : (*pParamMap)[imageParam].type;
@@ -406,15 +406,31 @@ class COCL_sample : public CImagesBI
406406
public:
407407
COCL_sample(ParamMap* paramMap, InlineMap* inlineMap, int* nextSampler, Dimension Dim, MetaDataUtils* pMdUtils, ModuleMetaData* modMD) : CImagesBI(paramMap, inlineMap, nextSampler, Dim), m_pMdUtils(pMdUtils), m_modMD(modMD) {}
408408

409-
ConstantInt* getSamplerIndex(void)
409+
Value* getSamplerValue(void)
410410
{
411411
ConstantInt* samplerIndex = nullptr;
412-
Value* samplerParam = ValueTracker::track(m_pCallInst, 1, m_pMdUtils, m_modMD);
412+
Value* samplerParam = ValueTracker::track(m_pCallInst, 1, m_pMdUtils, m_modMD, isBindlessImageOrSamplerLoad);
413413
if (!samplerParam) {
414414
emitError("There are instructions that use a sampler, but no sampler found in the kernel!", m_pCallInst);
415415
return nullptr;
416416
}
417417

418+
auto modMD = m_pCodeGenContext->getModuleMetaData();
419+
420+
// If bindless image is preferred, map the bindless pointer
421+
if (modMD->UseBindlessImage)
422+
{
423+
// If sampler is argument, look up index in the parameter map.
424+
int i = isa<Argument>(samplerParam) ? (*m_pParamMap)[samplerParam].index : 0;
425+
samplerIndex = ConstantInt::get(m_pIntType, i);
426+
unsigned int addressSpace = IGC::EncodeAS4GFXResource(*samplerIndex, BufferType::BINDLESS_SAMPLER);
427+
Type* ptrTy = llvm::PointerType::get(m_pFloatType, addressSpace);
428+
Value* bindlessSampler = isa<IntegerType>(samplerParam->getType()) ?
429+
BitCastInst::CreateBitOrPointerCast(samplerParam, ptrTy, "bindless_sampler", m_pCallInst) :
430+
BitCastInst::CreatePointerCast(samplerParam, ptrTy, "bindless_sampler", m_pCallInst);
431+
return bindlessSampler;
432+
}
433+
418434
// Argument samplers are looked up in the parameter map
419435
if (isa<Argument>(samplerParam))
420436
{
@@ -609,13 +625,17 @@ class COCL_sample : public CImagesBI
609625
}
610626
}
611627

612-
bool prepareSamplerIndex()
628+
bool prepareSamplerValue()
613629
{
614-
ConstantInt* samplerIndex = getSamplerIndex();
615-
if (!samplerIndex) return false;
616-
unsigned int addrSpace = EncodeAS4GFXResource(*samplerIndex, SAMPLER);
617-
Value* sampler = ConstantPointerNull::get(PointerType::get(samplerIndex->getType(), addrSpace));
618-
m_args.push_back(sampler);
630+
Value* samplerValue = getSamplerValue();
631+
if (!samplerValue) return false;
632+
if (isa<ConstantInt>(samplerValue))
633+
{
634+
unsigned int addrSpace = EncodeAS4GFXResource(*samplerValue, SAMPLER);
635+
samplerValue = ConstantPointerNull::get(PointerType::get(samplerValue->getType(), addrSpace));
636+
}
637+
638+
m_args.push_back(samplerValue);
619639
return true;
620640
}
621641

@@ -711,8 +731,8 @@ class COCL_sample_l : public COCL_sample
711731
m_args.push_back(m_pFloatZero); // ai (?)
712732
preparePairedResource();
713733
createGetBufferPtr();
714-
bool samplerIndexFound = prepareSamplerIndex();
715-
if (!samplerIndexFound) return;
734+
bool samplerValueFound = prepareSamplerValue();
735+
if (!samplerValueFound) return;
716736

717737
prepareZeroOffsets();
718738
Type* types[] = {
@@ -749,7 +769,7 @@ class COCL_sample_d : public COCL_sample
749769
m_args.push_back(m_pFloatZero); // minLOD (?)
750770
preparePairedResource();
751771
prepareImageBTI();
752-
prepareSamplerIndex();
772+
prepareSamplerValue();
753773
prepareZeroOffsets();
754774
Type* types[] = {
755775
m_pCallInst->getType(),

‎IGC/Compiler/Optimizer/OCLBIUtils.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -228,8 +228,8 @@ namespace IGC
228228
void prepareLOD(CoordType Coord);
229229

230230

231-
/// @brief push the sampler index into the function argument list
232-
void prepareSamplerIndex(void);
231+
/// @brief push the sampler value into the function argument list
232+
void prepareSamplerValue(void);
233233

234234
/// @brief create a call to the GetBufferPtr intrinsic pseudo-instruction
235235
/// @brief push the image index into the function argument list

‎IGC/Compiler/tests/OCLBIConverter/ocl_bindless_sampler_2d.ll‎

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2024 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; Check that bindless image and sampler are passed via kernel arguments.
10+
11+
; RUN: igc_opt %s -S -o - -igc-conv-ocl-to-common | FileCheck %s
12+
13+
%spirv.Image._void_1_0_0_0_0_0_0 = type opaque
14+
%spirv.Sampler = type opaque
15+
16+
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) {
17+
entry:
18+
%0 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %img to i64
19+
%1 = ptrtoint %spirv.Sampler addrspace(2)* %sampler to i64
20+
%conv = trunc i64 %0 to i32
21+
%conv2 = trunc i64 %1 to i32
22+
23+
; CHECK: %bindless_img = addrspacecast %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %img to float addrspace(393218)*
24+
; CHECK-NEXT: %bindless_sampler = addrspacecast %spirv.Sampler addrspace(2)* %sampler to float addrspace(655360)*
25+
; CHECK-NEXT: call <4 x float> @llvm.genx.GenISA.sampleLptr.v4f32.f32.p196610f32.p393218f32.p655360f32(float 0.000000e+00, float %CoordX, float %CoordY, float 0.000000e+00, float 0.000000e+00, float addrspace(196610)* undef, float addrspace(393218)* %bindless_img, float addrspace(655360)* %bindless_sampler, i32 0, i32 0, i32 0)
26+
27+
%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)
28+
ret void
29+
}
30+
31+
declare spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32, i32, <2 x float>, float)
32+
33+
!igc.functions = !{!0}
34+
!IGCMetadata = !{!2}
35+
36+
!0 = !{void (%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @image_read_sampler, !1}
37+
!1 = !{}
38+
!2 = !{!"ModuleMD", !3, !19}
39+
!3 = !{!"FuncMD", !4, !5}
40+
!4 = !{!"FuncMDMap[0]", void (%spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(2)*)* @image_read_sampler}
41+
!5 = !{!"FuncMDValue[0]", !6}
42+
!6 = !{!"resAllocMD", !7, !8, !9, !10}
43+
!7 = !{!"uavsNumType", i32 3}
44+
!8 = !{!"srvsNumType", i32 0}
45+
!9 = !{!"samplersNumType", i32 1}
46+
!10 = !{!"argAllocMDList", !11, !15}
47+
!11 = !{!"argAllocMDListVec[0]", !12, !13, !14}
48+
!12 = !{!"type", i32 4}
49+
!13 = !{!"extensionType", i32 0}
50+
!14 = !{!"indexType", i32 2}
51+
!15 = !{!"argAllocMDListVec[1]", !16, !17, !18}
52+
!16 = !{!"type", i32 5}
53+
!17 = !{!"extensionType", i32 -1}
54+
!18 = !{!"indexType", i32 0}
55+
!19 = !{!"UseBindlessImage", i1 true}

‎IGC/Compiler/tests/OCLBIConverter/sampled_image_handle_2d.ll‎

Lines changed: 77 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,77 @@
1+
;=========================== begin_copyright_notice ============================
2+
;
3+
; Copyright (C) 2024 Intel Corporation
4+
;
5+
; SPDX-License-Identifier: MIT
6+
;
7+
;============================ end_copyright_notice =============================
8+
9+
; Check that SYCL bindless image and sampler are passed via kernel arguments.
10+
11+
; RUN: igc_opt %s -S -o - -igc-conv-ocl-to-common | FileCheck %s
12+
13+
%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle" = type { %"struct.sycl::_V1::ext::oneapi::experimental::combined_sampled_image_handle" }
14+
%"struct.sycl::_V1::ext::oneapi::experimental::combined_sampled_image_handle" = type { i64, i64 }
15+
%spirv.Image._void_1_0_0_0_0_0_0 = type opaque
16+
%spirv.Sampler = type opaque
17+
18+
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) {
19+
entry:
20+
%0 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle1 to i64*
21+
%__SYCLKernel.1.copyload = load i64, i64* %0, align 8
22+
%1 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle1 to i8*
23+
%__SYCLKernel.imgHandle1.sroa_idx = getelementptr inbounds i8, i8* %1, i64 8
24+
%2 = bitcast i8* %__SYCLKernel.imgHandle1.sroa_idx to i64*
25+
%__SYCLKernel.2.copyload = load i64, i64* %2, align 8
26+
%3 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle2 to i64*
27+
%__SYCLKernel.3.copyload = load i64, i64* %3, align 8
28+
%4 = bitcast %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"* %imgHandle2 to i8*
29+
%__SYCLKernel.imgHandle2.sroa_idx = getelementptr inbounds i8, i8* %4, i64 8
30+
%5 = bitcast i8* %__SYCLKernel.imgHandle2.sroa_idx to i64*
31+
%__SYCLKernel.4.copyload = load i64, i64* %5, align 8
32+
%astype = inttoptr i64 %__SYCLKernel.1.copyload to %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*
33+
%astype2 = inttoptr i64 %__SYCLKernel.2.copyload to %spirv.Sampler addrspace(2)*
34+
%6 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype to i64
35+
%7 = ptrtoint %spirv.Sampler addrspace(2)* %astype2 to i64
36+
%conv = trunc i64 %6 to i32
37+
%conv2 = trunc i64 %7 to i32
38+
39+
; CHECK: %bindless_img = inttoptr i64 %__SYCLKernel.1.copyload to float addrspace(393216)*
40+
; CHECK-NEXT: %bindless_sampler = inttoptr i64 %__SYCLKernel.2.copyload to float addrspace(655360)*
41+
; 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)
42+
43+
%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)
44+
%astype3 = inttoptr i64 %__SYCLKernel.3.copyload to %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*
45+
%astype4 = inttoptr i64 %__SYCLKernel.4.copyload to %spirv.Sampler addrspace(2)*
46+
%8 = ptrtoint %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)* %astype3 to i64
47+
%9 = ptrtoint %spirv.Sampler addrspace(2)* %astype4 to i64
48+
%conv3 = trunc i64 %8 to i32
49+
%conv4 = trunc i64 %9 to i32
50+
51+
; CHECK: [[IMG2:%bindless_img[0-9]+]] = inttoptr i64 %__SYCLKernel.3.copyload to float addrspace(393216)*
52+
; CHECK-NEXT: [[SAMPLER2:%bindless_sampler[0-9]+]] = inttoptr i64 %__SYCLKernel.4.copyload to float addrspace(655360)*
53+
; 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)
54+
55+
%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)
56+
ret void
57+
}
58+
59+
declare spir_func <4 x float> @__builtin_IB_OCL_2d_sample_l(i32, i32, <2 x float>, float)
60+
61+
!igc.functions = !{!0}
62+
!IGCMetadata = !{!2}
63+
64+
!0 = !{void (%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"*, %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"*)* @_ZTS14image_addition, !1}
65+
!1 = !{}
66+
!2 = !{!"ModuleMD", !3, !13}
67+
!3 = !{!"FuncMD", !4, !5}
68+
!4 = !{!"FuncMDMap[0]", void (%"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"*, %"struct.sycl::_V1::ext::oneapi::experimental::sampled_image_handle"*)* @_ZTS14image_addition}
69+
!5 = !{!"FuncMDValue[0]", !6}
70+
!6 = !{!"resAllocMD", !7}
71+
!7 = !{!"argAllocMDList", !8, !12}
72+
!8 = !{!"argAllocMDListVec[0]", !9, !10, !11}
73+
!9 = !{!"type", i32 0}
74+
!10 = !{!"extensionType", i32 -1}
75+
!11 = !{!"indexType", i32 -1}
76+
!12 = !{!"argAllocMDListVec[1]", !9, !10, !11}
77+
!13 = !{!"UseBindlessImage", i1 true}

0 commit comments

Comments
 (0)