Skip to content

Offload sanitizer traps fix #13

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

Open
wants to merge 2 commits into
base: pr/offload_sanitizer_traps
Choose a base branch
from
Open

Offload sanitizer traps fix #13

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6955d57
Fix build by bringing back terminateWarp
EthanLuisMcDonough Nov 18, 2024
f6d27a6
Fix tests
EthanLuisMcDonough Nov 18, 2024
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
3 changes: 3 additions & 0 deletions offload/DeviceRTL/include/DeviceUtils.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,9 @@ int64_t shuffleDown(uint64_t Mask, int64_t Var, uint32_t Delta, int32_t Width);

uint64_t ballotSync(uint64_t Mask, int32_t Pred);

/// Terminate the execution of this warp.
void terminateWarp();

/// Return \p LowBits and \p HighBits packed into a single 64 bit value.
uint64_t pack(uint32_t LowBits, uint32_t HighBits);

Expand Down
7 changes: 7 additions & 0 deletions offload/DeviceRTL/src/DeviceUtils.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,7 @@ int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
int32_t Width);

uint64_t ballotSync(uint64_t Mask, int32_t Pred);
void terminateWarp();

/// AMDGCN Implementation
///
Expand All @@ -62,6 +63,8 @@ uint64_t ballotSync(uint64_t Mask, int32_t Pred) {
return Mask & __builtin_amdgcn_ballot_w64(Pred);
}

void terminateWarp() { __builtin_amdgcn_endpgm(); }

bool isSharedMemPtr(const void *Ptr) {
return __builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void *)Ptr);
Expand Down Expand Up @@ -89,6 +92,8 @@ uint64_t ballotSync(uint64_t Mask, int32_t Pred) {
return __nvvm_vote_ballot_sync(static_cast<uint32_t>(Mask), Pred);
}

void terminateWarp() { __nvvm_exit(); }

bool isSharedMemPtr(const void *Ptr) { return __nvvm_isspacep_shared(Ptr); }

#pragma omp end declare variant
Expand Down Expand Up @@ -126,6 +131,8 @@ uint64_t utils::ballotSync(uint64_t Mask, int32_t Pred) {
return impl::ballotSync(Mask, Pred);
}

void utils::terminateWarp() { return impl::terminateWarp(); }

bool utils::isSharedMemPtr(void *Ptr) { return impl::isSharedMemPtr(Ptr); }

extern "C" {
Expand Down
3 changes: 1 addition & 2 deletions offload/DeviceRTL/src/Sanitizer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,6 @@
#include "Mapping.h"
#include "Shared/Environment.h"
#include "Synchronization.h"
#include "gpuintrin.h"

using namespace ompx;

Expand Down Expand Up @@ -69,7 +68,7 @@ void raiseExecutionError(SanitizerEnvironmentTy::ErrorCodeTy ErrorCode,
// If no thread of this warp has the lock, end execution gracefully.
bool AnyThreadHasLock = utils::ballotSync(lanes::All, HasLock);
if (!AnyThreadHasLock)
__gpu_exit();
utils::terminateWarp();

// One thread will set the location information and signal that the rest of
// the wapr that the actual trap can be executed now.
Expand Down
4 changes: 2 additions & 2 deletions offload/test/jit/type_punning.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,8 +13,8 @@
// Ensure that there is only the kernel function left, not any outlined
// parallel regions.
//
// CHECK: define
// CHECK-NOT: define
// CHECK: define {{.*}}offloading
// CHECK-NOT: define {{.*}}offloading

#include <omp.h>
#include <stdio.h>
Expand Down
14 changes: 5 additions & 9 deletions offload/test/sanitizer/kernel_trap_all.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,15 +3,9 @@
// 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=CHECK,TRACE,DEBUG
// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
// clang-format on

// 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
// REQUIRES: amdgpu

#include <omp.h>

Expand All @@ -23,9 +17,11 @@ int main(void) {
__builtin_trap();
}
}
// CHECK: OFFLOAD ERROR: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l20)
// clang-format off
// CHECK: OFFLOAD ERROR: Kernel {{.*}} (__omp_offloading_{{.*}}_main_{{.*}})
// CHECK: OFFLOAD ERROR: execution interrupted by hardware trap instruction
// CHECK: Triggered by thread <{{[0-9]*}},0,0> block <{{[0-9]*}},0,0> PC 0x{{.*}}
// TRACE: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_trap_all.c:
// clang-format on