Skip to content

Commit af68a7b

Browse files
wenju-heigcbot
authored andcommitted
Fix assert in GetVariableLocation for bindless image -g build
Fix assert failure of unhandled BindlessUAVResourceType.
1 parent c1c9b84 commit af68a7b

File tree

2 files changed

+25
-0
lines changed

2 files changed

+25
-0
lines changed

IGC/Compiler/DebugInfo/ScalarVISAModule.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -456,6 +456,9 @@ ScalarVisaModule::GetVariableLocation(const llvm::Instruction* pInst) const
456456
index = m_pShader->m_pBtiLayout->GetTextureIndex(index);
457457
IGC_ASSERT_MESSAGE(index < TEXTURE_REGISTER_NUM, "Bad texture index");
458458
return VISAVariableLocation(TEXTURE_REGISTER_BEGIN + index, this);
459+
case BindlessUAVResourceType:
460+
IGC_ASSERT_MESSAGE(index < TEXTURE_REGISTER_NUM, "Bad texture index");
461+
return VISAVariableLocation(TEXTURE_REGISTER_BEGIN + index, this);
459462
default:
460463
IGC_ASSERT_MESSAGE(0, "Unknown texture resource");
461464
return VISAVariableLocation(this);

IGC/ocloc_tests/basic-image.cl

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// UNSUPPORTED: system-windows
2+
3+
// RUN: ocloc compile -file %s -options "-g" -internal_options "-cl-intel-use-bindless-mode" -device dg2 2>&1 | FileCheck %s
4+
5+
// Check that kernel build is successful in "-g" mode.
6+
7+
// CHECK: Build succeeded
8+
9+
const sampler_t sampler =
10+
CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
11+
12+
kernel void test(const global float *a,
13+
global float *c,
14+
read_only image2d_t input,
15+
sampler_t sampler
16+
) {
17+
const int gid = get_global_id(0);
18+
19+
int2 coord = {get_global_id(0), get_global_id(1)};
20+
float4 data = read_imagef(input, coord);
21+
c[gid] = a[gid] + data.x;
22+
}

0 commit comments

Comments
 (0)