Skip to content

Commit ff69436

Browse files
jdoerfertjdoerfert
committed
Ambiguous call trace support
1 parent 35c135b commit ff69436

File tree

7 files changed

+109
-32
lines changed

7 files changed

+109
-32
lines changed

‎llvm/lib/Passes/PassBuilderPipelines.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2068,11 +2068,11 @@ PassBuilder::buildLTODefaultPipeline(OptimizationLevel Level,
20682068
if (PTO.CallGraphProfile)
20692069
MPM.addPass(CGProfilePass(/*InLTOPostLink=*/true));
20702070

2071-
invokeFullLinkTimeOptimizationLastEPCallbacks(MPM, Level);
2072-
20732071
if (EnableOffloadSanitizer)
20742072
MPM.addPass(GPUSanPass());
20752073

2074+
invokeFullLinkTimeOptimizationLastEPCallbacks(MPM, Level);
2075+
20762076
// Emit annotation remarks.
20772077
addAnnotationRemarksPass(MPM);
20782078

‎llvm/lib/Transforms/Instrumentation/GPUSan.cpp‎

Lines changed: 71 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -778,42 +778,99 @@ bool GPUSanImpl::instrument() {
778778
return false;
779779
}();
780780

781-
for (Function &Fn : M)
781+
SmallVector<Function *> Kernels;
782+
for (Function &Fn : M) {
783+
if (Fn.hasFnAttribute("kernel"))
784+
Kernels.push_back(&Fn);
782785
if (!Fn.getName().contains("ompx") && !Fn.getName().contains("__kmpc") &&
783786
!Fn.getName().starts_with("rpc_"))
784787
if (!Fn.hasFnAttribute(Attribute::DisableSanitizerInstrumentation))
785788
Changed |= instrumentFunction(Fn);
789+
}
786790

787-
SmallVector<std::pair<CallBase *, ConstantInt *>> AmbiguousCallsNumbered;
791+
SmallVector<CallBase *> AmbiguousCallsOrdered;
792+
SmallVector<Constant *> AmbiguousCallsMapping;
793+
if (LocationMap.empty())
794+
AmbiguousCalls.clear();
788795
for (size_t I = 0; I < AmbiguousCalls.size(); ++I) {
789796
CallBase &CB = *AmbiguousCalls[I];
790-
AmbiguousCallsNumbered.push_back({&CB, getSourceIndex(CB)});
797+
AmbiguousCallsOrdered.push_back(&CB);
798+
AmbiguousCallsMapping.push_back(getSourceIndex(CB));
791799
}
792-
IntegerType *ITy = nullptr;
800+
801+
uint64_t AmbiguousCallsBitWidth =
802+
llvm::PowerOf2Ceil(AmbiguousCalls.size() + 1);
803+
804+
new GlobalVariable(M, Int64Ty, /*isConstant=*/true,
805+
GlobalValue::ExternalLinkage,
806+
ConstantInt::get(Int64Ty, AmbiguousCallsBitWidth),
807+
"__san.num_ambiguous_calls", nullptr,
808+
GlobalValue::ThreadLocalMode::NotThreadLocal, 1);
809+
793810
if (size_t NumAmbiguousCalls = AmbiguousCalls.size()) {
794-
ITy = IntegerType::get(Ctx, llvm::PowerOf2Ceil(NumAmbiguousCalls));
795-
auto *ArrayTy = ArrayType::get(ITy, 1024);
811+
{
812+
auto *ArrayTy = ArrayType::get(Int64Ty, NumAmbiguousCalls);
813+
auto *GV = new GlobalVariable(
814+
M, ArrayTy, /*isConstant=*/true, GlobalValue::ExternalLinkage,
815+
ConstantArray::get(ArrayTy, AmbiguousCallsMapping),
816+
"__san.ambiguous_calls_mapping", nullptr,
817+
GlobalValue::ThreadLocalMode::NotThreadLocal, 4);
818+
GV->setVisibility(GlobalValue::ProtectedVisibility);
819+
}
820+
821+
auto *ArrayTy = ArrayType::get(Int64Ty, 1024);
796822
LocationsArray = new GlobalVariable(
797823
M, ArrayTy, /*isConstant=*/false, GlobalValue::PrivateLinkage,
798824
UndefValue::get(ArrayTy), "__san.calls", nullptr,
799825
GlobalValue::ThreadLocalMode::NotThreadLocal, 3);
800826

827+
auto *OldFn = M.getFunction("__san_get_location_value");
828+
if (OldFn)
829+
OldFn->setName("");
801830
Function *LocationGetter = Function::Create(
802-
FunctionType::get(Int64Ty, false), llvm::GlobalValue::ExternalLinkage,
831+
FunctionType::get(Int64Ty, false), GlobalValue::ExternalLinkage,
803832
"__san_get_location_value", M);
833+
if (OldFn) {
834+
OldFn->replaceAllUsesWith(LocationGetter);
835+
OldFn->eraseFromParent();
836+
}
804837
auto *EntryBB = BasicBlock::Create(Ctx, "entry", LocationGetter);
805838
IRBuilder<> IRB(EntryBB);
806839
Value *Idx = IRB.CreateCall(getThreadIdFn(), {}, "san.gtid");
807-
Value *Ptr = IRB.CreateGEP(ITy, LocationsArray, {Idx});
808-
auto *LocationValue = IRB.CreateLoad(ITy, Ptr);
809-
IRB.CreateRet(IRB.CreateZExt(LocationValue, Int64Ty));
840+
Value *Ptr = IRB.CreateGEP(Int64Ty, LocationsArray, {Idx});
841+
auto *LocationValue = IRB.CreateLoad(Int64Ty, Ptr);
842+
IRB.CreateRet(LocationValue);
843+
}
844+
845+
Function *InitSharedFn =
846+
Function::Create(FunctionType::get(VoidTy, false),
847+
GlobalValue::PrivateLinkage, "__san.init_shared", &M);
848+
auto *EntryBB = BasicBlock::Create(Ctx, "entry", InitSharedFn);
849+
IRBuilder<> IRB(EntryBB);
850+
if (!AmbiguousCalls.empty()) {
851+
Value *Idx = IRB.CreateCall(getThreadIdFn(), {}, "san.gtid");
852+
Value *Ptr = IRB.CreateGEP(Int64Ty, LocationsArray, {Idx});
853+
IRB.CreateStore(ConstantInt::get(Int64Ty, 0), Ptr);
854+
}
855+
IRB.CreateRetVoid();
856+
857+
for (auto *KernelFn : Kernels) {
858+
IRBuilder<> IRB(&*KernelFn->getEntryBlock().getFirstNonPHIOrDbgOrAlloca());
859+
IRB.CreateCall(InitSharedFn, {});
810860
}
811861

812-
for (auto &It : AmbiguousCallsNumbered) {
813-
IRBuilder<> IRB(It.first);
862+
for (const auto &It : llvm::enumerate(AmbiguousCallsOrdered)) {
863+
IRBuilder<> IRB(It.value());
814864
Value *Idx = IRB.CreateCall(getThreadIdFn(), {}, "san.gtid");
815-
Value *Ptr = IRB.CreateGEP(ITy, LocationsArray, {Idx});
816-
IRB.CreateStore(It.second, Ptr);
865+
Value *Ptr = IRB.CreateGEP(Int64Ty, LocationsArray, {Idx});
866+
Value *OldVal = IRB.CreateLoad(Int64Ty, Ptr);
867+
Value *OldValShifted = IRB.CreateShl(
868+
OldVal, ConstantInt::get(Int64Ty, AmbiguousCallsBitWidth));
869+
Value *NewVal = IRB.CreateBinOp(Instruction::Or, OldValShifted,
870+
ConstantInt::get(Int64Ty, It.index() + 1));
871+
IRB.CreateStore(NewVal, Ptr);
872+
IRB.SetInsertPoint(It.value()->getNextNode());
873+
IRB.CreateStore(OldVal, Ptr);
817874
}
818875

819876
auto *NamesTy = ArrayType::get(Int8Ty, ConcatenatedString.size() + 1);

‎offload/DeviceRTL/src/Mapping.cpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -364,7 +364,9 @@ _TGT_KERNEL_LANGUAGE(block_id, getBlockIdInKernel)
364364
_TGT_KERNEL_LANGUAGE(block_dim, getNumberOfThreadsInBlock)
365365
_TGT_KERNEL_LANGUAGE(grid_dim, getNumberOfBlocksInKernel)
366366

367-
extern "C" int ompx_global_thread_id() {
367+
extern "C" [[clang::disable_sanitizer_instrumentation, gnu::flatten,
368+
gnu::always_inline, gnu::used, gnu::retain]] int
369+
ompx_global_thread_id() {
368370
return ompx_thread_id(0) + ompx_thread_id(1) * ompx_block_dim(0) +
369371
ompx_thread_id(2) * ompx_block_dim(0) * ompx_block_dim(1);
370372
}

‎offload/DeviceRTL/src/Sanitizer.cpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -416,6 +416,11 @@ ompx_get_allocation_info_global(_AS_PTR(void, AllocationKind::GLOBAL) P) {
416416
ompx_leak_check() {
417417
AllocationTracker<AllocationKind::GLOBAL>::leakCheck();
418418
}
419+
420+
[[gnu::weak, gnu::noinline, gnu::used, gnu::retain]] int64_t
421+
__san_get_location_value() {
422+
return -1;
423+
}
419424
}
420425

421426
#pragma omp end declare target

‎offload/include/Shared/Sanitizer.h‎

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -18,10 +18,7 @@ extern "C" {
1818
int ompx_block_id(int Dim);
1919
int ompx_block_dim(int Dim);
2020
int ompx_thread_id(int Dim);
21-
[[clang::disable_sanitizer_instrumentation, gnu::noinline]] inline int64_t
22-
__san_get_location_value() {
23-
return -1;
24-
}
21+
int64_t __san_get_location_value();
2522
}
2623

2724
enum class AllocationKind { LOCAL, GLOBAL, LAST = GLOBAL };
@@ -182,7 +179,7 @@ struct SanitizerTrapInfoTy {
182179
uint32_t ThreadId[3];
183180
uint64_t PC;
184181
uint64_t LocationId;
185-
uint64_t CallId;
182+
int64_t CallId;
186183
/// }
187184

188185
[[clang::disable_sanitizer_instrumentation]] void

‎offload/plugins-nextgen/common/src/PluginInterface.cpp‎

Lines changed: 22 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -2198,13 +2198,11 @@ void GPUSanTy::checkAndReportError() {
21982198
DeviceImageTy *Image = nullptr;
21992199
for (auto &It : Device.SanitizerTrapInfos) {
22002200
STI = It.second;
2201-
errs() << "STI " << STI << "\n";
22022201
if (!STI || STI->ErrorCode == SanitizerTrapInfoTy::None)
22032202
continue;
22042203
Image = It.first;
22052204
break;
22062205
}
2207-
errs() << "Img " << Image << "\n";
22082206
if (!Image)
22092207
return;
22102208

@@ -2214,19 +2212,26 @@ void GPUSanTy::checkAndReportError() {
22142212
auto Default = []() { return "\033[1m\033[0m"; };
22152213

22162214
GenericGlobalHandlerTy &GHandler = Device.Plugin.getGlobalHandler();
2217-
auto GetImagePtr = [&](GlobalTy &GV) {
2215+
auto GetImagePtr = [&](GlobalTy &GV, bool Quiet = false) {
22182216
if (auto Err = GHandler.getGlobalMetadataFromImage(Device, *Image, GV)) {
2219-
REPORT("WARNING: Failed to read backtrace "
2220-
"(%s)\n",
2221-
toString(std::move(Err)).data());
2217+
if (Quiet)
2218+
consumeError(std::move(Err));
2219+
else
2220+
REPORT("WARNING: Failed to read backtrace "
2221+
"(%s)\n",
2222+
toString(std::move(Err)).data());
22222223
return false;
22232224
}
22242225
return true;
22252226
};
22262227
GlobalTy LocationsGV("__san.locations", -1);
22272228
GlobalTy LocationNamesGV("__san.location_names", -1);
2229+
GlobalTy AmbiguousCallsBitWidthGV("__san.num_ambiguous_calls", -1);
2230+
GlobalTy AmbiguousCallsLocationsGV("__san.ambiguous_calls_mapping", -1);
22282231
if (GetImagePtr(LocationsGV))
22292232
GetImagePtr(LocationNamesGV);
2233+
GetImagePtr(AmbiguousCallsBitWidthGV, /*Quiet=*/true);
2234+
GetImagePtr(AmbiguousCallsLocationsGV, /*Quiet=*/true);
22302235

22312236
fprintf(stderr, "============================================================"
22322237
"====================\n");
@@ -2236,9 +2241,12 @@ void GPUSanTy::checkAndReportError() {
22362241
fprintf(stderr, " no backtrace available\n");
22372242
return;
22382243
}
2239-
fprintf(stderr, "%lu\n", STI->CallId);
22402244
char *LocationNames = LocationNamesGV.getPtrAs<char>();
22412245
LocationEncodingTy *Locations = LocationsGV.getPtrAs<LocationEncodingTy>();
2246+
uint64_t *AmbiguousCallsBitWidth =
2247+
AmbiguousCallsBitWidthGV.getPtrAs<uint64_t>();
2248+
uint64_t *AmbiguousCallsLocations =
2249+
AmbiguousCallsLocationsGV.getPtrAs<uint64_t>();
22422250
int32_t FrameIdx = 0;
22432251
do {
22442252
LocationEncodingTy &LE = Locations[LocationId];
@@ -2247,6 +2255,13 @@ void GPUSanTy::checkAndReportError() {
22472255
&LocationNames[LE.FileNameIdx], LE.LineNo, LE.ColumnNo);
22482256
LocationId = LE.ParentIdx;
22492257
FrameIdx++;
2258+
if (LocationId < 0 && STI->CallId != 0 && AmbiguousCallsBitWidth &&
2259+
AmbiguousCallsLocations) {
2260+
uint64_t LastCallId =
2261+
STI->CallId & ((1 << *AmbiguousCallsBitWidth) - 1);
2262+
LocationId = AmbiguousCallsLocations[LastCallId - 1];
2263+
STI->CallId >>= (*AmbiguousCallsBitWidth);
2264+
}
22502265
} while (LocationId >= 0);
22512266
fputc('\n', stderr);
22522267
};

‎offload/test/sanitizer/stack_trace_multi_path_1.c‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -16,18 +16,19 @@
1616
// UNSUPPORTED: s390x-ibm-linux-gnu
1717
// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
1818

19+
#include <omp.h>
20+
1921
[[clang::optnone]] int deref(int *P) { return *P; }
2022

2123
[[gnu::noinline]] int bar(int *P) { return deref(P); }
2224
[[gnu::noinline]] int baz(int *P) { return deref(P); }
2325

2426
int main(void) {
2527

26-
#pragma omp target
28+
int *Valid = (int *)omp_target_alloc(4, omp_get_default_device());
29+
#pragma omp target is_device_ptr(Valid)
2730
{
2831
int *NullPtr = 0;
29-
int X;
30-
int *Valid = &X;
3132
// clang-format off
3233
// CHECK: ERROR: OffloadSanitizer out-of-bounds access on address 0x0000000000000000 at pc [[PC:.*]]
3334
// CHECK: WRITE of size 4 at 0x0000000000000000 thread <0, 0, 0> block <0, 0, 0> (acc 1, heap)

0 commit comments

Comments
 (0)