Skip to content

Commit

Permalink
The magic numbers from the nextBaseAddress are incorrectly compared with
Browse files Browse the repository at this point in the history
constant values. They are usually simplified as unequal values.
  • Loading branch information
Oleksandr Demydenko committed Oct 3, 2024
1 parent 11ea6b5 commit 17c7305
Show file tree
Hide file tree
Showing 2 changed files with 26 additions and 1 deletion.
10 changes: 9 additions & 1 deletion lib/LogicalPointerToIntPass.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -139,10 +139,18 @@ clspv::LogicalPointerToIntPass::run(Module &M, ModuleAnalysisManager &MAM) {
: IntegerType::get(M.getContext(), 32);
IRBuilder<> B(ICmp);
auto LHS = B.CreatePtrToInt(ICmp->getOperand(0), SizeT);
auto RHS = B.CreatePtrToInt(ICmp->getOperand(1), SizeT);

// The magic numbers from the nextBaseAddress are incorrectly compared with
// constant values. They are usually simplified as unequal values.
if (dyn_cast<ConstantInt>(LHS) != nullptr ||
dyn_cast<ConstantInt>(RHS) != nullptr) {
continue;
}

if (auto *Cast = dyn_cast<CastInst>(LHS)) {
InstrsToProcess.push_back(Cast);
}
auto RHS = B.CreatePtrToInt(ICmp->getOperand(1), SizeT);
if (auto *Cast = dyn_cast<CastInst>(RHS)) {
InstrsToProcess.push_back(Cast);
}
Expand Down
17 changes: 17 additions & 0 deletions test/LogicalPtrToInt/compare_constant.cl
Original file line number Diff line number Diff line change
@@ -0,0 +1,17 @@
// RUN: clspv %target %s -o %t.spv --hack-logical-ptrtoint -cl-kernel-arg-info -spv-version=1.6
// RUN: spirv-dis %t.spv -o %t.spvasm
// RUN: FileCheck %s < %t.spvasm
// RUN: spirv-val --target-env vulkan1.3 %t.spv

// CHECK: [[uint:%[a-zA-Z]+]] = OpTypeInt 32 0
// CHECK: [[uint0:%[a-zA-Z0-9_]+]] = OpConstant [[uint]] 0
// CHECK: [[nullptr:%[0-9]+]] = OpConstantNull %{{[a-zA-Z0-9_]+}}
// CHECK: [[bool:%[a-zA-Z]+]] = OpTypeBool
// CHECK: [[ptr:%[a-zA-Z0-9_]+]] = OpAccessChain %{{[a-zA-Z0-9_]+}} %{{[a-zA-Z0-9_]+}} [[uint0]] [[uint0]]
// CHECK: [[cmp:%[a-zA-Z0-9_]+]] = OpPtrNotEqual [[bool]] [[ptr]] [[nullptr]]

kernel void test_kernel(global float *src, global long *dst)
{
uint tid = get_global_id(0);
dst[tid] = (long)(src != 0);
}

0 comments on commit 17c7305

Please sign in to comment.