Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

rename Intrinsic::getDeclaration to Intrinsic::getOrInsertDeclaration #1406

Merged
merged 1 commit into from
Oct 17, 2024
Merged

rename Intrinsic::getDeclaration to Intrinsic::getOrInsertDeclaration #1406

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion deps.json
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@
"subrepo" : "llvm/llvm-project",
"branch" : "main",
"subdir" : "third_party/llvm",
"commit" : "125aa10b3d645bd26523a1bc321bb2e6b1cf04e1"
"commit" : "30deb76d46053c243561c6fa072c5a30407241cb"
},
{
"name" : "SPIRV-Headers",
Expand Down
54 changes: 27 additions & 27 deletions lib/LongVectorLoweringPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,9 +31,9 @@
#include "Constants.h"
#include "clspv/Passes.h"

#include "BitcastUtils.h"
#include "Builtins.h"
#include "LongVectorLoweringPass.h"
#include "BitcastUtils.h"

#include <array>
#include <functional>
Expand Down Expand Up @@ -120,7 +120,8 @@ Function *getIntrinsicScalarVersion(Function &Intrinsic) {
Param = Param->getScalarType();
}

return Intrinsic::getDeclaration(Intrinsic.getParent(), id, ParamTys);
return Intrinsic::getOrInsertDeclaration(Intrinsic.getParent(), id,
ParamTys);
break;
}
}
Expand Down Expand Up @@ -154,30 +155,29 @@ std::string getSpirvCompliantName(const clspv::Builtins::FunctionInfo &IInfo) {
Type *getScalarPointerType(Function &Builtin) {
const auto &Info = clspv::Builtins::Lookup(&Builtin);
switch (Info.getType()) {
case clspv::Builtins::kSincos:
case clspv::Builtins::kModf:
case clspv::Builtins::kFract:
return Builtin.getReturnType()->getScalarType();
case clspv::Builtins::kFrexp:
case clspv::Builtins::kRemquo:
case clspv::Builtins::kLgammaR:
return Type::getInt32Ty(Builtin.getParent()->getContext());
case clspv::Builtins::kVloadHalf:
case clspv::Builtins::kVloadaHalf:
case clspv::Builtins::kVstoreHalf:
case clspv::Builtins::kVstoreaHalf:
return Type::getHalfTy(Builtin.getParent()->getContext());
case clspv::Builtins::kVstore:
return Builtin.getArg(0)->getType()->getScalarType();
case clspv::Builtins::kVload:
return Builtin.getReturnType()->getScalarType();
default:
// What about llvm intrinsics (e.g. memcpy) or other OpenCL builtins?
return nullptr;
case clspv::Builtins::kSincos:
case clspv::Builtins::kModf:
case clspv::Builtins::kFract:
return Builtin.getReturnType()->getScalarType();
case clspv::Builtins::kFrexp:
case clspv::Builtins::kRemquo:
case clspv::Builtins::kLgammaR:
return Type::getInt32Ty(Builtin.getParent()->getContext());
case clspv::Builtins::kVloadHalf:
case clspv::Builtins::kVloadaHalf:
case clspv::Builtins::kVstoreHalf:
case clspv::Builtins::kVstoreaHalf:
return Type::getHalfTy(Builtin.getParent()->getContext());
case clspv::Builtins::kVstore:
return Builtin.getArg(0)->getType()->getScalarType();
case clspv::Builtins::kVload:
return Builtin.getReturnType()->getScalarType();
default:
// What about llvm intrinsics (e.g. memcpy) or other OpenCL builtins?
return nullptr;
}
}


/// Get the scalar overload for the given OpenCL builtin function @p Builtin.
Function *getBIFScalarVersion(Function &Builtin) {
assert(!Builtin.isIntrinsic());
Expand Down Expand Up @@ -520,7 +520,8 @@ Function *createFunctionWithMappedTypes(Function &F,
B.CreateRetVoid();
} else {
auto *EquivalentReturnTy = EquivalentFunctionTy->getReturnType();
Value *ReturnValue = convertEquivalentValue(B, Call, EquivalentReturnTy, DL);
Value *ReturnValue =
convertEquivalentValue(B, Call, EquivalentReturnTy, DL);
B.CreateRet(ReturnValue);
}

Expand Down Expand Up @@ -916,8 +917,7 @@ Value *clspv::LongVectorLoweringPass::visitCastInst(CastInst &I) {
V = B.CreatePtrToInt(EquivalentValue, EquivalentDestTy, I.getName());
break;
}
case Instruction::IntToPtr:
{
case Instruction::IntToPtr: {
IRBuilder<> B(&I);
V = B.CreateIntToPtr(EquivalentValue, EquivalentDestTy, I.getName());
break;
Expand Down Expand Up @@ -1049,7 +1049,7 @@ void clspv::LongVectorLoweringPass::reworkIndices(
for (unsigned i = 1; i < Idxs.size(); i++) {
Indices.push_back(Idxs[i]);
// Get original indices up to ith element for below:
auto CumulativeOldIdxs = ArrayRef<Value*>(Idxs.begin(), Idxs.begin() + i);
auto CumulativeOldIdxs = ArrayRef<Value *>(Idxs.begin(), Idxs.begin() + i);
auto IndexedTy = GetElementPtrInst::getIndexedType(Ty, CumulativeOldIdxs);
if (getEquivalentType(IndexedTy)) {
auto Idx = Indices.pop_back_val();
Expand Down
7 changes: 4 additions & 3 deletions lib/ReplaceLLVMIntrinsicsPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,8 +95,9 @@ bool clspv::ReplaceLLVMIntrinsicsPass::runOnFunction(Function &F) {
return replaceAddSubSat(F, true, true);
case Intrinsic::is_fpclass:
return replaceIsFpClass(F);
// SPIR-V OpAssumeTrueKHR requires ExpectAssumeKHR capability in SPV_KHR_expect_assume extension.
// Vulkan doesn't support that, so remove assume declaration.
// SPIR-V OpAssumeTrueKHR requires ExpectAssumeKHR capability in
// SPV_KHR_expect_assume extension. Vulkan doesn't support that, so remove
// assume declaration.
case Intrinsic::assume:
// SPIR-V OpLifetimeStart and OpLifetimeEnd require Kernel capability.
// Vulkan doesn't support that, so remove all lifteime bounds declarations.
Expand Down Expand Up @@ -465,7 +466,7 @@ bool clspv::ReplaceLLVMIntrinsicsPass::replaceCountZeroes(Function &F,
IRBuilder<> builder(Call);
auto ty = Call->getType()->getWithNewBitWidth(32);
auto c32 = ConstantInt::get(ty, 32);
auto func_32bit = Intrinsic::getDeclaration(
auto func_32bit = Intrinsic::getOrInsertDeclaration(
F.getParent(), leading ? Intrinsic::ctlz : Intrinsic::cttz, ty);
if (bitwidth < 32) {
// Extend the input to 32-bits and perform a clz/ctz.
Expand Down
68 changes: 35 additions & 33 deletions lib/ReplaceOpenCLBuiltinPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -878,7 +878,8 @@ Value *ReplaceOpenCLBuiltinPass::InsertOpMulExtended(Instruction *InsertPoint,
res_neg = Builder.CreateICmpSLT(a_b_xor, ConstantInt::get(Ty, 0, true));

auto F = InsertPoint->getFunction();
auto abs = Intrinsic::getDeclaration(F->getParent(), Intrinsic::abs, Ty);
auto abs =
Intrinsic::getOrInsertDeclaration(F->getParent(), Intrinsic::abs, Ty);
a = Builder.CreateCall(abs, {a, Builder.getInt1(false)});
b = Builder.CreateCall(abs, {b, Builder.getInt1(false)});
}
Expand Down Expand Up @@ -1190,7 +1191,7 @@ bool ReplaceOpenCLBuiltinPass::replaceCopysign(Function &F) {
return replaceCallsWithValue(F, [&F](CallInst *Call) {
const auto x = Call->getArgOperand(0);
const auto y = Call->getArgOperand(1);
auto intrinsic = Intrinsic::getDeclaration(
auto intrinsic = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::copysign, Call->getType());
return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {x, y}, "",
Call);
Expand Down Expand Up @@ -1308,8 +1309,8 @@ bool ReplaceOpenCLBuiltinPass::replaceLog1p(Function &F) {
auto ArgP1 = BinaryOperator::Create(
Instruction::FAdd, ConstantFP::get(Arg->getType(), 1.0), Arg, "", CI);

auto log =
Intrinsic::getDeclaration(F.getParent(), Intrinsic::log, CI->getType());
auto log = Intrinsic::getOrInsertDeclaration(F.getParent(), Intrinsic::log,
CI->getType());
return CallInst::Create(log, ArgP1, "", CI);
});
}
Expand Down Expand Up @@ -1728,8 +1729,8 @@ bool ReplaceOpenCLBuiltinPass::replaceRotate(Function &F) {

// Replace with LLVM's funnel shift left intrinsic because it is more
// generic than rotate.
Function *intrinsic =
Intrinsic::getDeclaration(F.getParent(), Intrinsic::fshl, SrcType);
Function *intrinsic = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::fshl, SrcType);
return CallInst::Create(intrinsic->getFunctionType(), intrinsic,
{SrcValue, SrcValue, RotAmount}, "", CI);
});
Expand Down Expand Up @@ -2883,7 +2884,8 @@ llvm::Value *ReplaceOpenCLBuiltinPass::createVstoreHalf(llvm::Module &M,
CallInst::Create(NewF, Params, "store_halfword_xor_trick", CI);

// Return a Nop so the old Call is removed
Function *donothing = Intrinsic::getDeclaration(&M, Intrinsic::donothing);
Function *donothing =
Intrinsic::getOrInsertDeclaration(&M, Intrinsic::donothing);
V = CallInst::Create(donothing, {}, "", CI);
}

Expand Down Expand Up @@ -3633,8 +3635,8 @@ bool ReplaceOpenCLBuiltinPass::replaceAddSubSat(Function &F, bool is_signed,
: (is_add ? Intrinsic::uadd_sat : Intrinsic::usub_sat);
auto a = Call->getArgOperand(0);
auto b = Call->getArgOperand(1);
auto intrinsic = Intrinsic::getDeclaration(F.getParent(), intrinsic_type,
Call->getType());
auto intrinsic = Intrinsic::getOrInsertDeclaration(
F.getParent(), intrinsic_type, Call->getType());
return CallInst::Create(intrinsic->getFunctionType(), intrinsic, {a, b}, "",
Call);
});
Expand Down Expand Up @@ -3821,7 +3823,7 @@ bool ReplaceOpenCLBuiltinPass::replaceCountZeroes(Function &F, bool leading) {
return false;

return replaceCallsWithValue(F, [&F, leading](CallInst *Call) {
Function *intrinsic = Intrinsic::getDeclaration(
Function *intrinsic = Intrinsic::getOrInsertDeclaration(
F.getParent(), leading ? Intrinsic::ctlz : Intrinsic::cttz,
Call->getType());
const auto c_false = ConstantInt::getFalse(Call->getContext());
Expand Down Expand Up @@ -4104,13 +4106,13 @@ bool ReplaceOpenCLBuiltinPass::replaceRound(Function &F) {
clspv_fract_fn->setCallingConv(CallingConv::SPIR_FUNC);
}

auto ceil = Intrinsic::getDeclaration(F.getParent(), Intrinsic::ceil,
Call->getType());
auto floor = Intrinsic::getDeclaration(F.getParent(), Intrinsic::floor,
Call->getType());
auto fabs = Intrinsic::getDeclaration(F.getParent(), Intrinsic::fabs,
Call->getType());
auto copysign = Intrinsic::getDeclaration(
auto ceil = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::ceil, Call->getType());
auto floor = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::floor, Call->getType());
auto fabs = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::fabs, Call->getType());
auto copysign = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::copysign, {Call->getType(), Call->getType()});

IRBuilder<> builder(Call);
Expand Down Expand Up @@ -4138,21 +4140,21 @@ bool ReplaceOpenCLBuiltinPass::replaceTrigPi(Function &F,
auto mul = builder.CreateFMul(x, pi);
switch (type) {
case Builtins::kSinpi: {
auto func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin,
x->getType());
auto func = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::sin, x->getType());
return builder.CreateCall(func->getFunctionType(), func, {mul});
}
case Builtins::kCospi: {
auto func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos,
x->getType());
auto func = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::cos, x->getType());
return builder.CreateCall(func->getFunctionType(), func, {mul});
}
case Builtins::kTanpi: {
auto sin = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin,
x->getType());
auto sin = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::sin, x->getType());
auto sin_call = builder.CreateCall(sin->getFunctionType(), sin, {mul});
auto cos = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos,
x->getType());
auto cos = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::cos, x->getType());
auto cos_call = builder.CreateCall(cos->getFunctionType(), cos, {mul});
return builder.CreateFDiv(sin_call, cos_call);
}
Expand All @@ -4166,10 +4168,10 @@ bool ReplaceOpenCLBuiltinPass::replaceTrigPi(Function &F,

bool ReplaceOpenCLBuiltinPass::replaceSincos(Function &F) {
return replaceCallsWithValue(F, [&F](CallInst *Call) {
auto sin_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::sin,
Call->getType());
auto cos_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::cos,
Call->getType());
auto sin_func = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::sin, Call->getType());
auto cos_func = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::cos, Call->getType());

IRBuilder<> builder(Call);
auto sin = builder.CreateCall(sin_func->getFunctionType(), sin_func,
Expand All @@ -4183,8 +4185,8 @@ bool ReplaceOpenCLBuiltinPass::replaceSincos(Function &F) {

bool ReplaceOpenCLBuiltinPass::replaceExpm1(Function &F) {
return replaceCallsWithValue(F, [&F](CallInst *Call) {
auto exp_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::exp,
Call->getType());
auto exp_func = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::exp, Call->getType());

IRBuilder<> builder(Call);
auto exp = builder.CreateCall(exp_func->getFunctionType(), exp_func,
Expand All @@ -4195,8 +4197,8 @@ bool ReplaceOpenCLBuiltinPass::replaceExpm1(Function &F) {

bool ReplaceOpenCLBuiltinPass::replacePown(Function &F) {
return replaceCallsWithValue(F, [&F](CallInst *Call) {
auto pow_func = Intrinsic::getDeclaration(F.getParent(), Intrinsic::pow,
Call->getType());
auto pow_func = Intrinsic::getOrInsertDeclaration(
F.getParent(), Intrinsic::pow, Call->getType());

IRBuilder<> builder(Call);
auto conv = builder.CreateSIToFP(Call->getArgOperand(1), Call->getType());
Expand Down
2 changes: 1 addition & 1 deletion test/ProgramScopeConstants/in_storage_buffer_descriptor_map.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: clspv %target %s -o %t.spv -module-constants-in-storage-buffer -int8=0
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@rjodinchr observed that Clang is inserting explicit padding members.

In this case we now have an i8 first member, a [3 x i8] inserted second member. Everything breaks when artificially trying to remap i8 to u32. Use the default, which is actually support i8.

// RUN: clspv %target %s -o %t.spv -module-constants-in-storage-buffer
// RUN: clspv-reflection %t.spv -o %t.map
// RUN: FileCheck -check-prefix=MAP %s < %t.map
// RUN: spirv-val --target-env vulkan1.0 %t.spv
Expand Down
7 changes: 3 additions & 4 deletions test/UBO/transform_global.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,11 +54,10 @@ __kernel void foo(__global data_type *data, __constant data_type *c_arg,
// CHECK-DAG: [[c_var_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Private [[array]]
// CHECK-DAG: [[c_var_ele_ptr:%[0-9a-zA-Z_]+]] = OpTypePointer Private [[int]]
// CHECK-DAG: [[zero:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 0
// CHECK-DAG: [[undef:%[0-9a-zA-Z_]+]] = OpUndef [[int]]
// CHECK-DAG: [[zero_undef:%[0-9a-zA-Z_]+]] = OpConstantComposite [[data_type]] [[zero]] [[undef]]
// CHECK-DAG: [[zero_zero:%[0-9a-zA-Z_]+]] = OpConstantNull [[data_type]]
// CHECK-DAG: [[one:%[0-9a-zA-Z_]+]] = OpConstant [[int]] 1
// CHECK-DAG: [[one_undef:%[0-9a-zA-Z_]+]] = OpConstantComposite [[data_type]] [[one]] [[undef]]
// CHECK-DAG: [[array_const:%[0-9a-zA-Z_]+]] = OpConstantComposite [[array]] [[zero_undef]] [[one_undef]]
// CHECK-DAG: [[one_zero:%[0-9a-zA-Z_]+]] = OpConstantComposite [[data_type]] [[one]] [[zero]]
// CHECK-DAG: [[array_const:%[0-9a-zA-Z_]+]] = OpConstantComposite [[array]] [[zero_zero]] [[one_zero]]
// CHECK-DAG: [[c_var:%[0-9a-zA-Z_]+]] = OpVariable [[c_var_ptr]] Private [[array_const]]
// CHECK-DAG: [[data]] = OpVariable [[data_ptr]] StorageBuffer
// CHECK: [[c_arg]] = OpVariable [[c_arg_ptr]] Uniform
Expand Down