Author: Shilei Tian Date: 2021-02-05T20:47:20-08:00 New Revision: 716eef9ad5b367e5cbcc22c8ac53395f9bdbe7a5
URL: https://github.com/llvm/llvm-project/commit/716eef9ad5b367e5cbcc22c8ac53395f9bdbe7a5 DIFF: https://github.com/llvm/llvm-project/commit/716eef9ad5b367e5cbcc22c8ac53395f9bdbe7a5.diff LOG: [OpenMP][libomptarget] Fixed an issue that device sync is skipped if the kernel doesn't have any argument Currently if there is not kernel argument, device synchronization will be skipped. This can lead to two issues: 1. If there is any device error, it will not be captured; 2. The target region might end before the kernel is done, which is not spec conformant. The test added in this patch only runs on NVPTX platform, although it will not be executed by Phab at all. It also requires `not` which is not available on most systems. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D96067 (cherry picked from commit b68a6b09e60a24733b923a0fc282746a855852da) Added: openmp/libomptarget/test/offloading/assert.cpp Modified: openmp/libomptarget/src/omptarget.cpp Removed: ################################################################################ diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 90966d25fb26..e4b7b18bc70b 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -451,6 +451,17 @@ struct DeallocTgtPtrInfo { : HstPtrBegin(HstPtr), DataSize(Size), ForceDelete(ForceDelete), HasCloseModifier(HasCloseModifier) {} }; + +/// Synchronize device +static int syncDevice(DeviceTy &Device, __tgt_async_info *AsyncInfo) { + assert(AsyncInfo && AsyncInfo->Queue && "Invalid AsyncInfo"); + if (Device.synchronize(AsyncInfo) != OFFLOAD_SUCCESS) { + REPORT("Failed to synchronize device.\n"); + return OFFLOAD_FAIL; + } + + return OFFLOAD_SUCCESS; +} } // namespace /// Internal function to undo the mapping and retrieve the data from the device. @@ -631,11 +642,9 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum, // AsyncInfo->Queue will not be nullptr, so again, we don't need to // synchronize. if (AsyncInfo && AsyncInfo->Queue) { - Ret = Device.synchronize(AsyncInfo); - if (Ret != OFFLOAD_SUCCESS) { - REPORT("Failed to synchronize device.\n"); + Ret = syncDevice(Device, AsyncInfo); + if (Ret != OFFLOAD_SUCCESS) return OFFLOAD_FAIL; - } } // Deallocate target pointer @@ -1307,6 +1316,11 @@ int target(ident_t *loc, int64_t DeviceId, void *HostPtr, int32_t ArgNum, REPORT("Failed to process data after launching the kernel.\n"); return OFFLOAD_FAIL; } + } else if (AsyncInfo.Queue) { + // If ArgNum is zero, but AsyncInfo.Queue is valid, then the kernel doesn't + // hava any argument, and the device supports async operations, so we need a + // sync at this point. + return syncDevice(Device, &AsyncInfo); } return OFFLOAD_SUCCESS; diff --git a/openmp/libomptarget/test/offloading/assert.cpp b/openmp/libomptarget/test/offloading/assert.cpp new file mode 100644 index 000000000000..00112dd92cc6 --- /dev/null +++ b/openmp/libomptarget/test/offloading/assert.cpp @@ -0,0 +1,8 @@ +// RUN: %libomptarget-compilexx-nvptx64-nvidia-cuda && %libomptarget-run-fail-nvptx64-nvidia-cuda + +int main(int argc, char *argv[]) { +#pragma omp target + { __builtin_trap(); } + + return 0; +} _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
