Skip to content

Commit 918d599

Browse files
abhinavgabaabhinavgaba
committed
Merge from 'master' to 'sycl-web' (#2)
 CONFLICT (content): Merge conflict in clang/lib/CodeGen/CGCall.cpp
2 parents e17e513 + 9142c0b commit 918d599

File tree

3 files changed

+45
-30
lines changed

3 files changed

+45
-30
lines changed

‎clang/lib/CodeGen/CGCall.cpp‎

Lines changed: 37 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -1022,8 +1022,8 @@ static void forConstantArrayExpansion(CodeGenFunction &CGF,
10221022
}
10231023
}
10241024

1025-
void CodeGenFunction::ExpandTypeFromArgs(
1026-
QualType Ty, LValue LV, SmallVectorImpl<llvm::Value *>::iterator &AI) {
1025+
void CodeGenFunction::ExpandTypeFromArgs(QualType Ty, LValue LV,
1026+
llvm::Function::arg_iterator &AI) {
10271027
assert(LV.isSimple() &&
10281028
"Unexpected non-simple lvalue during struct expansion.");
10291029

@@ -1052,17 +1052,17 @@ void CodeGenFunction::ExpandTypeFromArgs(
10521052
ExpandTypeFromArgs(FD->getType(), SubLV, AI);
10531053
}
10541054
} else if (isa<ComplexExpansion>(Exp.get())) {
1055-
auto realValue = *AI++;
1056-
auto imagValue = *AI++;
1055+
auto realValue = &*AI++;
1056+
auto imagValue = &*AI++;
10571057
EmitStoreOfComplex(ComplexPairTy(realValue, imagValue), LV, /*init*/ true);
10581058
} else {
10591059
// Call EmitStoreOfScalar except when the lvalue is a bitfield to emit a
10601060
// primitive store.
10611061
assert(isa<NoExpansion>(Exp.get()));
10621062
if (LV.isBitField())
1063-
EmitStoreThroughLValue(RValue::get(*AI++), LV);
1063+
EmitStoreThroughLValue(RValue::get(&*AI++), LV);
10641064
else
1065-
EmitStoreOfScalar(*AI++, LV);
1065+
EmitStoreOfScalar(&*AI++, LV);
10661066
}
10671067
}
10681068

@@ -2329,27 +2329,21 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
23292329
// simplify.
23302330

23312331
ClangToLLVMArgMapping IRFunctionArgs(CGM.getContext(), FI);
2332-
// Flattened function arguments.
2333-
SmallVector<llvm::Value *, 16> FnArgs;
2334-
FnArgs.reserve(IRFunctionArgs.totalIRArgs());
2335-
for (auto &Arg : Fn->args()) {
2336-
FnArgs.push_back(&Arg);
2337-
}
2338-
assert(FnArgs.size() == IRFunctionArgs.totalIRArgs());
2332+
assert(Fn->arg_size() == IRFunctionArgs.totalIRArgs());
23392333

23402334
// If we're using inalloca, all the memory arguments are GEPs off of the last
23412335
// parameter, which is a pointer to the complete memory area.
23422336
Address ArgStruct = Address::invalid();
23432337
if (IRFunctionArgs.hasInallocaArg()) {
2344-
ArgStruct = Address(FnArgs[IRFunctionArgs.getInallocaArgNo()],
2338+
ArgStruct = Address(Fn->getArg(IRFunctionArgs.getInallocaArgNo()),
23452339
FI.getArgStructAlignment());
23462340

23472341
assert(ArgStruct.getType() == FI.getArgStruct()->getPointerTo());
23482342
}
23492343

23502344
// Name the struct return parameter.
23512345
if (IRFunctionArgs.hasSRetArg()) {
2352-
auto AI = cast<llvm::Argument>(FnArgs[IRFunctionArgs.getSRetArgNo()]);
2346+
auto AI = Fn->getArg(IRFunctionArgs.getSRetArgNo());
23532347
AI->setName("agg.result");
23542348
AI->addAttr(llvm::Attribute::NoAlias);
23552349
}
@@ -2400,7 +2394,8 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
24002394

24012395
case ABIArgInfo::Indirect: {
24022396
assert(NumIRArgs == 1);
2403-
Address ParamAddr = Address(FnArgs[FirstIRArg], ArgI.getIndirectAlign());
2397+
Address ParamAddr =
2398+
Address(Fn->getArg(FirstIRArg), ArgI.getIndirectAlign());
24042399

24052400
if (!hasScalarEvaluationKind(Ty)) {
24062401
// Aggregates and complex variables are accessed by reference. All we
@@ -2436,16 +2431,18 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
24362431

24372432
case ABIArgInfo::Extend:
24382433
case ABIArgInfo::Direct: {
2439-
2440-
// If we have the trivial case, handle it with no muss and fuss.
2441-
if (!isa<llvm::StructType>(ArgI.getCoerceToType()) &&
2442-
ArgI.getCoerceToType() == ConvertType(Ty) &&
2443-
ArgI.getDirectOffset() == 0) {
2434+
auto AI = Fn->getArg(FirstIRArg);
2435+
llvm::Type *LTy = ConvertType(Arg->getType());
2436+
2437+
// Prepare parameter attributes. So far, only attributes for pointer
2438+
// parameters are prepared. See
2439+
// http://llvm.org/docs/LangRef.html#paramattrs.
2440+
if (ArgI.getDirectOffset() == 0 && LTy->isPointerTy() &&
2441+
ArgI.getCoerceToType()->isPointerTy()) {
24442442
assert(NumIRArgs == 1);
2445-
llvm::Value *V = FnArgs[FirstIRArg];
2446-
auto AI = cast<llvm::Argument>(V);
24472443

24482444
if (const ParmVarDecl *PVD = dyn_cast<ParmVarDecl>(Arg)) {
2445+
// Set `nonnull` attribute if any.
24492446
if (getNonNullAttr(CurCodeDecl, PVD, PVD->getType(),
24502447
PVD->getFunctionScopeIndex()) &&
24512448
!CGM.getCodeGenOpts().NullPointerIsValid)
@@ -2483,6 +2480,7 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
24832480
AI->addAttr(llvm::Attribute::NonNull);
24842481
}
24852482

2483+
// Set `align` attribute if any.
24862484
const auto *AVAttr = PVD->getAttr<AlignValueAttr>();
24872485
if (!AVAttr)
24882486
if (const auto *TOTy = dyn_cast<TypedefType>(OTy))
@@ -2500,14 +2498,24 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
25002498
}
25012499
}
25022500

2501+
// Set 'noalias' if an argument type has the `restrict` qualifier.
25032502
if (Arg->getType().isRestrictQualified() ||
25042503
(CurCodeDecl &&
25052504
CurCodeDecl->hasAttr<SYCLIntelKernelArgsRestrictAttr>() &&
25062505
Arg->getType()->isPointerType()))
25072506
AI->addAttr(llvm::Attribute::NoAlias);
2507+
}
2508+
2509+
// Prepare the argument value. If we have the trivial case, handle it
2510+
// with no muss and fuss.
2511+
if (!isa<llvm::StructType>(ArgI.getCoerceToType()) &&
2512+
ArgI.getCoerceToType() == ConvertType(Ty) &&
2513+
ArgI.getDirectOffset() == 0) {
2514+
assert(NumIRArgs == 1);
25082515

25092516
// LLVM expects swifterror parameters to be used in very restricted
25102517
// ways. Copy the value into a less-restricted temporary.
2518+
llvm::Value *V = AI;
25112519
if (FI.getExtParameterInfo(ArgNo).getABI()
25122520
== ParameterABI::SwiftErrorResult) {
25132521
QualType pointeeTy = Ty->getPointeeType();
@@ -2569,7 +2577,7 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
25692577

25702578
assert(STy->getNumElements() == NumIRArgs);
25712579
for (unsigned i = 0, e = STy->getNumElements(); i != e; ++i) {
2572-
auto AI = FnArgs[FirstIRArg + i];
2580+
auto AI = Fn->getArg(FirstIRArg + i);
25732581
AI->setName(Arg->getName() + ".coerce" + Twine(i));
25742582
Address EltPtr = Builder.CreateStructGEP(AddrToStoreInto, i);
25752583
Builder.CreateStore(AI, EltPtr);
@@ -2582,7 +2590,7 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
25822590
} else {
25832591
// Simple case, just do a coerced store of the argument into the alloca.
25842592
assert(NumIRArgs == 1);
2585-
auto AI = FnArgs[FirstIRArg];
2593+
auto AI = Fn->getArg(FirstIRArg);
25862594
AI->setName(Arg->getName() + ".coerce");
25872595
CreateCoercedStore(AI, Ptr, /*DstIsVolatile=*/false, *this);
25882596
}
@@ -2615,7 +2623,7 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
26152623
continue;
26162624

26172625
auto eltAddr = Builder.CreateStructGEP(alloca, i);
2618-
auto elt = FnArgs[argIndex++];
2626+
auto elt = Fn->getArg(argIndex++);
26192627
Builder.CreateStore(elt, eltAddr);
26202628
}
26212629
assert(argIndex == FirstIRArg + NumIRArgs);
@@ -2630,11 +2638,11 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI,
26302638
LValue LV = MakeAddrLValue(Alloca, Ty);
26312639
ArgVals.push_back(ParamValue::forIndirect(Alloca));
26322640

2633-
auto FnArgIter = FnArgs.begin() + FirstIRArg;
2641+
auto FnArgIter = Fn->arg_begin() + FirstIRArg;
26342642
ExpandTypeFromArgs(Ty, LV, FnArgIter);
2635-
assert(FnArgIter == FnArgs.begin() + FirstIRArg + NumIRArgs);
2643+
assert(FnArgIter == Fn->arg_begin() + FirstIRArg + NumIRArgs);
26362644
for (unsigned i = 0, e = NumIRArgs; i != e; ++i) {
2637-
auto AI = FnArgs[FirstIRArg + i];
2645+
auto AI = Fn->getArg(FirstIRArg + i);
26382646
AI->setName(Arg->getName() + "." + Twine(i));
26392647
}
26402648
break;

‎clang/lib/CodeGen/CodeGenFunction.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4367,7 +4367,7 @@ class CodeGenFunction : public CodeGenTypeCache {
43674367
///
43684368
/// \param AI - The first function argument of the expansion.
43694369
void ExpandTypeFromArgs(QualType Ty, LValue Dst,
4370-
SmallVectorImpl<llvm::Value *>::iterator &AI);
4370+
llvm::Function::arg_iterator &AI);
43714371

43724372
/// ExpandTypeToArgs - Expand an CallArg \arg Arg, with the LLVM type for \arg
43734373
/// Ty, into individual arguments on the provided vector \arg IRCallArgs,

‎clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -67,3 +67,10 @@ __global__ void kernel6(struct T t) {
6767
t.x[0][0] += 1.f;
6868
t.x[1][0] += 2.f;
6969
}
70+
71+
// Check that coerced pointers retain the noalias attribute when qualified with __restrict.
72+
// CHECK: define amdgpu_kernel void @_Z7kernel7Pi(i32 addrspace(1)* noalias %x.coerce)
73+
// HOST: define void @_Z22__device_stub__kernel7Pi(i32* noalias %x)
74+
__global__ void kernel7(int *__restrict x) {
75+
x[0]++;
76+
}

0 commit comments

Comments
 (0)