[Offload] Sanitize "standalone" unreachable instructions

If an unreachable is reached, the execution state is invalid. If the sanitizer is enabled, we stop and report it to the user.

Author: jdoerfert

llvm/lib/Transforms/Instrumentation/OffloadSanitizer.cpp

@@ -42,6 +42,8 @@ class OffloadSanitizerImpl final {
  bool shouldInstrumentFunction(Function &Fn);
  bool instrumentFunction(Function &Fn);
  bool instrumentTrapInstructions(SmallVectorImpl<IntrinsicInst *> &TrapCalls);
+ bool instrumentUnreachableInstructions(
+ SmallVectorImpl<UnreachableInst *> &UnreachableInsts);
 
  FunctionCallee getOrCreateFn(FunctionCallee &FC, StringRef Name, Type *RetTy,
  ArrayRef<Type *> ArgTys) {
@@ -59,6 +61,13 @@ class OffloadSanitizerImpl final {
  {/*PC*/ Int64Ty});
  }
 
+ /// void __offload_san_unreachable_info(Int64Ty);
+ FunctionCallee UnreachableInfoFn;
+ FunctionCallee getUnreachableInfoFn() {
+ return getOrCreateFn(UnreachableInfoFn, "__offload_san_unreachable_info",
+ VoidTy, {/*PC*/ Int64Ty});
+ }
+
  CallInst *createCall(IRBuilder<> &IRB, FunctionCallee Callee,
  ArrayRef<Value *> Args = std::nullopt,
  const Twine &Name = "") {
@@ -107,15 +116,34 @@ bool OffloadSanitizerImpl::instrumentTrapInstructions(
  return Changed;
 }
 
+bool OffloadSanitizerImpl::instrumentUnreachableInstructions(
+ SmallVectorImpl<UnreachableInst *> &UnreachableInsts) {
+ bool Changed = false;
+ for (auto *II : UnreachableInsts) {
+ // Skip unreachables after traps since we instrument those as well.
+ if (&II->getParent()->front() != II)
+ if (auto *CI = dyn_cast<CallInst>(II->getPrevNode()))
+ if (CI->getIntrinsicID() == Intrinsic::trap)
+ continue;
+ IRBuilder<> IRB(II);
+ createCall(IRB, getUnreachableInfoFn(), {getPC(IRB)});
+ }
+ return Changed;
+}
+
 bool OffloadSanitizerImpl::instrumentFunction(Function &Fn) {
  if (!shouldInstrumentFunction(Fn))
  return false;
 
+ SmallVector<UnreachableInst *> UnreachableInsts;
  SmallVector<IntrinsicInst *> TrapCalls;
 
  bool Changed = false;
  for (auto &I : instructions(Fn)) {
  switch (I.getOpcode()) {
+ case Instruction::Unreachable:
+ UnreachableInsts.push_back(cast<UnreachableInst>(&I));
+ break;
  case Instruction::Call: {
  auto &CI = cast<CallInst>(I);
  if (auto *II = dyn_cast<IntrinsicInst>(&CI))
@@ -129,6 +157,7 @@ bool OffloadSanitizerImpl::instrumentFunction(Function &Fn) {
  }
 
  Changed |= instrumentTrapInstructions(TrapCalls);
+ Changed |= instrumentUnreachableInstructions(UnreachableInsts);
 
  return Changed;
 }

llvm/test/Instrumentation/OffloadSanitizer/basic.ll

@@ -54,3 +54,32 @@ t:
 f:
  ret void
 }
+
+define void @test_unreachable1() {
+; CHECK-LABEL: define void @test_unreachable1() {
+; CHECK-NEXT: [[PC:%.*]] = call i64 @llvm.amdgcn.s.getpc()
+; CHECK-NEXT: call void @__offload_san_unreachable_info(i64 [[PC]])
+; CHECK-NEXT: unreachable
+;
+ unreachable
+}
+
+define void @test_unreachable2(i1 %c) {
+; CHECK-LABEL: define void @test_unreachable2(
+; CHECK-SAME: i1 [[C:%.*]]) {
+; CHECK-NEXT: [[ENTRY:.*:]]
+; CHECK-NEXT: br i1 [[C]], label %[[T:.*]], label %[[F:.*]]
+; CHECK: [[T]]:
+; CHECK-NEXT: [[PC:%.*]] = call i64 @llvm.amdgcn.s.getpc()
+; CHECK-NEXT: call void @__offload_san_unreachable_info(i64 [[PC]])
+; CHECK-NEXT: unreachable
+; CHECK: [[F]]:
+; CHECK-NEXT: ret void
+;
+entry:
+ br i1 %c, label %t ,label %f
+t:
+ unreachable
+f:
+ ret void
+}

offload/DeviceRTL/src/Sanitizer.cpp

@@ -90,6 +90,10 @@ extern "C" {
 _SAN_ENTRY_ATTRS void __offload_san_trap_info(uint64_t PC) {
  raiseExecutionError(SanitizerEnvironmentTy::TRAP, PC);
 }
+
+_SAN_ENTRY_ATTRS void __offload_san_unreachable_info(uint64_t PC) {
+ raiseExecutionError(SanitizerEnvironmentTy::UNREACHABLE, PC);
+}
 }
 
 #pragma omp end declare target

offload/include/Shared/Environment.h

@@ -111,7 +111,8 @@ struct SanitizerEnvironmentTy {
  enum ErrorCodeTy : uint8_t {
  NONE = 0,
  TRAP,
- LAST = TRAP,
+ UNREACHABLE,
+ LAST = UNREACHABLE,
  } ErrorCode;
 
  /// Flag to indicate the environment has been initialized fully.

offload/plugins-nextgen/common/include/ErrorReporting.h

@@ -281,6 +281,10 @@ class ErrorReporter {
  case SanitizerEnvironmentTy::TRAP:
  reportError("execution interrupted by hardware trap instruction");
  break;
+ case SanitizerEnvironmentTy::UNREACHABLE:
+ reportError("execution reached an \"unreachable\" state (likely caused "
+ "by undefined behavior)");
+break;
  default:
  reportError(
  "execution stopped, reason is unknown due to invalid error code");

offload/test/sanitizer/kernel_known_ub.c

@@ -0,0 +1,35 @@
+
+// clang-format off
+// RUN: %libomptarget-compile-generic -g -mllvm -amdgpu-enable-offload-sanitizer
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=SANIT
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=SANIT
+// RUN: %libomptarget-compileopt-generic -g -mllvm -amdgpu-enable-offload-sanitizer
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=SANIT
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=SANIT
+
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+// UNSUPPORTED: s390x-ibm-linux-gnu
+// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
+
+#include <omp.h>
+
+__attribute__((noinline)) void unreachable(volatile int *GoodPtr) {
+ *GoodPtr = 1;
+ __builtin_unreachable();
+}
+
+int main(void) {
+#pragma omp target
+ {
+ volatile int A = 0;
+ unreachable(&A);
+ }
+}
+// SANIT: OFFLOAD ERROR: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l27)
+// SANIT: OFFLOAD ERROR: execution reached an "unreachable" state (likely caused by undefined behavior)
+// SANIT: Triggered by thread <{{.*}},0,0> block <{{.*}},0,0> PC 0x{{.*}}