Skip to content

Commit

Permalink
Apply comments
Browse files Browse the repository at this point in the history
  • Loading branch information
AlexeySachkov committed Nov 24, 2020
1 parent 09fb54a commit 4c9ce32
Show file tree
Hide file tree
Showing 4 changed files with 52 additions and 45 deletions.
Original file line number Diff line number Diff line change
@@ -1,6 +1,10 @@
; RUN: sycl-post-link -spec-const=default --ir-output-only %s -S -o - \
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
;
; This test checks that composite specialization constants can be correctly
; initialized by sycl-post-link tool for AOT use-case (default initialization
; should be used according to the type of constant)
;
; TODO: consider adding a test case with vector type: the pass itself already
; supports this, but at the moment, sycl::vec type is not a POD type, which
; means we can't have it within a spec constant, i.e. we can't generate LLVM IR
Expand Down
2 changes: 1 addition & 1 deletion llvm/test/tools/sycl-post-link/composite-spec-constant.ll
Original file line number Diff line number Diff line change
Expand Up @@ -83,7 +83,7 @@ attributes #4 = { convergent }
!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 1, i32 2}
!2 = !{i32 4, i32 100000}
!3 = !{!"clang version 12.0.0 (/data/github.com/intel/llvm/clang 56ee5b054b5a1f2f703722fc414fcb05af18b40a)"}
!3 = !{!"clang version 12.0.0 "}
!4 = !{i32 -1, i32 -1, i32 -1, i32 -1}
!5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6}
!6 = !{!7, !7, i64 0}
Expand Down
77 changes: 35 additions & 42 deletions llvm/tools/sycl-post-link/SpecConstants.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -109,8 +109,6 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo,
return Res;
}

// TODO support spec constant types other than integer or
// floating-point.
Value *getDefaultCPPValue(Type *T) {
if (T->isIntegerTy())
return Constant::getIntegerValue(T, APInt(T->getScalarSizeInBits(), 0));
Expand Down Expand Up @@ -167,16 +165,13 @@ std::string manglePrimitiveType(const Type *T) {
// llvm-spirv doesn't care about the mangling and the only intent here is to
// make sure that we won't encounter redefinition error when we proceed two
// spec constants with different types.
if (T->isStructTy()) {
if (T->isStructTy())
return T->getStructName().str();
}
if (T->isArrayTy()) {
if (T->isArrayTy())
return "A" + manglePrimitiveType(T->getArrayElementType());
}
if (auto *VecTy = dyn_cast<FixedVectorType>(T)) {
if (auto *VecTy = dyn_cast<FixedVectorType>(T))
return "Dv" + std::to_string(VecTy->getNumElements()) + "_" +
manglePrimitiveType(VecTy->getElementType());
}
llvm_unreachable("unsupported spec const type");
return "";
}
Expand All @@ -196,10 +191,9 @@ void setSpecConstSymIDMetadata(Instruction *I, StringRef SymID,
LLVMContext &Ctx = I->getContext();
SmallVector<Metadata *, 4> MDOperands;
MDOperands.push_back(MDString::get(Ctx, SymID));
for (unsigned ID : IntIDs) {
for (unsigned ID : IntIDs)
MDOperands.push_back(
ConstantAsMetadata::get(ConstantInt::get(Ctx, APInt(32, ID))));
}
MDNode *Entry = MDNode::get(Ctx, MDOperands);
I->setMetadata(SPEC_CONST_SYM_ID_MD_STRING, Entry);
}
Expand All @@ -218,7 +212,7 @@ getScalarSpecConstMetadata(const Instruction *I) {

void collectCompositeElementsInfoRecursive(
const Type *Ty, unsigned &Index, unsigned &Offset,
std::vector<CompositeSpecConstDescriptor> &Result) {
std::vector<CompositeSpecConstElementDescriptor> &Result) {
if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
for (size_t I = 0; I < ArrTy->getNumElements(); ++I) {
// TODO: this is a spot for potential optimization: for arrays we could
Expand All @@ -233,14 +227,14 @@ void collectCompositeElementsInfoRecursive(
}
} else if (auto *VecTy = dyn_cast<FixedVectorType>(Ty)) {
for (size_t I = 0; I < VecTy->getNumElements(); ++I) {
// TODO: this is a spot for potential optimization: for arrays we could
// TODO: this is a spot for potential optimization: for vectors we could
// just make a single recursive call here and use it to populate Result
// in a loop.
collectCompositeElementsInfoRecursive(VecTy->getElementType(), Index,
Offset, Result);
}
} else { // Assume that we encountered some scalar element
CompositeSpecConstDescriptor Desc;
CompositeSpecConstElementDescriptor Desc;
Desc.ID = 0; // To be filled later
Desc.Offset = Offset;
Desc.Size = Ty->getPrimitiveSizeInBits() / 8;
Expand All @@ -249,14 +243,16 @@ void collectCompositeElementsInfoRecursive(
}
}

std::pair<StringRef, std::vector<CompositeSpecConstDescriptor>>
std::pair<StringRef, std::vector<CompositeSpecConstElementDescriptor>>
getCompositeSpecConstMetadata(const Instruction *I) {
const MDNode *N = I->getMetadata(SPEC_CONST_SYM_ID_MD_STRING);
if (!N)
return std::make_pair("", std::vector<CompositeSpecConstDescriptor>{});
return std::make_pair("",
std::vector<CompositeSpecConstElementDescriptor>{});
const auto *MDSym = cast<MDString>(N->getOperand(0));

std::vector<CompositeSpecConstDescriptor> Result(N->getNumOperands() - 1);
std::vector<CompositeSpecConstElementDescriptor> Result(N->getNumOperands() -
1);
unsigned Index = 0, Offset = 0;
collectCompositeElementsInfoRecursive(I->getType(), Index, Offset, Result);

Expand All @@ -269,6 +265,21 @@ getCompositeSpecConstMetadata(const Instruction *I) {
return std::make_pair(MDSym->getString(), Result);
}

Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
ArrayRef<Value *> Args, Instruction *InsertBefore) {
SmallVector<Type *, 8> ArgTys(Args.size());
for (unsigned I = 0; I < Args.size(); ++I) {
ArgTys[I] = Args[I]->getType();
}
auto *FT = FunctionType::get(RetTy, ArgTys, false /*isVarArg*/);
std::string FunctionName = mangleFuncItanium(BaseFunctionName, FT);
Module *M = InsertBefore->getFunction()->getParent();
FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT);
assert(FC.getCallee() && "SPIRV intrinsic creation failed");
auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
return Call;
}

Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,
Instruction *InsertBefore) {
Function *F = InsertBefore->getFunction();
Expand All @@ -279,37 +290,17 @@ Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,
Value *Def = getDefaultCPPValue(Ty);
// ... Now replace the call with SPIRV intrinsic version.
Value *Args[] = {ID, Def};
constexpr size_t NArgs = sizeof(Args) / sizeof(Args[0]);
Type *ArgTys[NArgs] = {nullptr};
for (unsigned int I = 0; I < NArgs; ++I)
ArgTys[I] = Args[I]->getType();
FunctionType *FT = FunctionType::get(Ty, ArgTys, false /*isVarArg*/);
Module *M = F->getParent();
std::string SPIRVName = mangleFuncItanium(SPIRV_GET_SPEC_CONST_VAL, FT);
FunctionCallee FC = M->getOrInsertFunction(SPIRVName, FT);
assert(FC.getCallee() && "SPIRV intrinsic creation failed");
CallInst *SpecConstant =
CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
return SpecConstant;
return emitCall(Ty, SPIRV_GET_SPEC_CONST_VAL, Args, InsertBefore);
}

Instruction *emitSpecConstantComposite(Type *Ty,
ArrayRef<Instruction *> Elements,
Instruction *InsertBefore) {
SmallVector<Type *, 8> ArgTys(Elements.size());
SmallVector<Value *, 8> Args(Elements.size());
for (unsigned I = 0; I < Elements.size(); ++I) {
ArgTys[I] = Elements[I]->getType();
Args[I] = cast<Value>(Elements[I]);
}
auto *FT = FunctionType::get(Ty, ArgTys, false /*isVarArg*/);
Module *M = InsertBefore->getFunction()->getParent();
std::string SPIRVName = mangleFuncItanium(SPIRV_GET_SPEC_CONST_COMPOSITE, FT);
FunctionCallee FC = M->getOrInsertFunction(SPIRVName, FT);
assert(FC.getCallee() && "SPIRV intrinsic creation failed");
CallInst *SpecConstant =
CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
return SpecConstant;
return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Args, InsertBefore);
}

Instruction *
Expand Down Expand Up @@ -416,13 +407,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
NextID += GeneratedIDs.size();
}

if (IsComposite)
if (IsComposite) {
// __sycl_getCompositeSpecConstant returns through argument, so, the
// only thing we need to do here is to store into a memory pointed by
// that argument
new StoreInst(SPIRVCall, CI->getArgOperand(0), CI);
else
} else {
CI->replaceAllUsesWith(SPIRVCall);
}

// Mark the instruction with <symbolic_id, int_ids...> list for later
// recollection by collectSpecConstantMetadata method.
Expand All @@ -442,13 +434,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
// 2a. Spec constant must be resolved at compile time - just replace
// the intrinsic with default C++ value for the spec constant type.
Value *Default = getDefaultCPPValue(SCTy);
if (IsComposite)
if (IsComposite) {
// __sycl_getCompositeSpecConstant returns through argument, so, the
// only thing we need to do here is to store into a memory pointed by
// that argument
new StoreInst(Default, CI->getArgOperand(0), CI);
else
} else {
CI->replaceAllUsesWith(Default);
}
}

for (auto *I : DelInsts) {
Expand Down
14 changes: 12 additions & 2 deletions llvm/tools/sycl-post-link/SpecConstants.h
Original file line number Diff line number Diff line change
Expand Up @@ -25,13 +25,23 @@
using namespace llvm;

using ScalarSpecIDMapTy = std::map<StringRef, unsigned>;
struct CompositeSpecConstDescriptor {
// Represents an element of a composite speciailization constant - at SYCL RT
// level composite specialization constants are being represented as a single
// byte-array, while at SPIR-V level they are represented by a number of scalar
// specialization constants.
struct CompositeSpecConstElementDescriptor {
// Encodes ID of a scalar specialization constants which is a leaf of some
// composite specialization constant.
unsigned ID;
// Encodes offset from the beginning of composite, where scalar resides, i.e.
// location of the scalar value within a byte-array containing the whole
// composite specialization constant.
unsigned Offset;
// Encodes size of scalar specialization constant.
unsigned Size;
};
using CompositeSpecIDMapTy =
std::map<StringRef, std::vector<CompositeSpecConstDescriptor>>;
std::map<StringRef, std::vector<CompositeSpecConstElementDescriptor>>;

class SpecConstantsPass : public PassInfoMixin<SpecConstantsPass> {
public:
Expand Down

0 comments on commit 4c9ce32

Please sign in to comment.