-
Notifications
You must be signed in to change notification settings - Fork 10.8k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[OpenMP] Allow exceptions in target regions when offloading to GPUs
The motivation for this patch is that many code bases use exception handling. As GPUs are not expected to support exception handling in the near future, we can experiment with compiling the code for GPU targets anyway. This will allow us to run the code, as long as no exception is thrown. The overall idea is very simple: - If a throw expression is compiled to AMDGCN or NVPTX, it is replaced with a trap during code generation. - If a try/catch statement is compiled to AMDGCN or AMDHSA, we ganerate code for the try statement as if it were a basic block. With this patch, the compilation of the following example ```{C++} int gaussian_sum(int a,int b){ if ((a + b) % 2 == 0) {throw -1;}; return (a+b) * ((a+b)/2); } int main(void) { int gauss = 0; #pragma omp target map(from:gauss) { try { gauss = gaussian_sum(1,100); } catch (int e){ gauss = e; } } std::cout << "GaussianSum(1,100)="<<gauss<<std::endl; #pragma omp target map(from:gauss) { try { gauss = gaussian_sum(1,101); } catch (int e){ gauss = e; } } std::cout << "GaussianSum(1,101)="<<gauss<<std::endl; return (gauss > 1) ? 0 : 1; } ``` with offloading to `gfx906` results in ```{bash} ./bin/target_try_minimal_fail GaussianSum(1,100)=5050 AMDGPU fatal error 1: Received error in queue 0x155555506000: HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception. zsh: abort (core dumped) ``` Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D153924
- Loading branch information
1 parent
f6259d9
commit 0cfc2db
Showing
17 changed files
with
371 additions
and
13 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,46 @@ | ||
/** | ||
* The first four lines test that a warning is produced when enabling | ||
* -Wopenmp-target-exception no matter what combination of -fexceptions and | ||
* -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the | ||
* target region but emit a warning instead. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
|
||
/** | ||
* The following four lines test that no warning is emitted when providing | ||
* -Wno-openmp-target-exception no matter the combination of -fexceptions and | ||
* -fcxx-exceptions. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
|
||
/** | ||
* Finally we should test that we only ignore exceptions in the OpenMP | ||
* offloading tool-chain | ||
*/ | ||
|
||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o - | ||
|
||
// noexceptions-error@37 {{cannot use 'try' with exceptions disabled}} | ||
// noexceptions-error@38 {{cannot use 'throw' with exceptions disabled}} | ||
|
||
#pragma omp declare target | ||
int foo(void) { | ||
int error = -1; | ||
try { // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'catch' block is ignored}} | ||
throw 404; // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'throw' is assumed to be never reached}} | ||
} | ||
catch (int e){ | ||
error = e; | ||
} | ||
return error; | ||
} | ||
#pragma omp end declare target | ||
// without-no-diagnostics |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,38 @@ | ||
/** | ||
* The first four lines test that a warning is produced when enabling | ||
* -Wopenmp-target-exception no matter what combination of -fexceptions and | ||
* -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the | ||
* target region but emit a warning instead. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
|
||
/** | ||
* The following four lines test that no warning is emitted when providing | ||
* -Wno-openmp-target-exception no matter the combination of -fexceptions and | ||
* -fcxx-exceptions. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
|
||
/** | ||
* Finally we should test that we only ignore exceptions in the OpenMP | ||
* offloading tool-chain | ||
*/ | ||
|
||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o - | ||
|
||
// noexceptions-error@35 {{cannot use 'throw' with exceptions disabled}} | ||
|
||
#pragma omp declare target | ||
void foo(void) { | ||
throw 404; // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'throw' is assumed to be never reached}} | ||
} | ||
#pragma omp end declare target | ||
// without-no-diagnostics |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,11 @@ | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=DEVICE %s | ||
// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=HOST %s | ||
// DEVICE: s_trap | ||
// DEVICE-NOT: __cxa_throw | ||
// HOST: __cxa_throw | ||
// HOST-NOT: s_trap | ||
#pragma omp declare target | ||
void foo(void) { | ||
throw 404; | ||
} | ||
#pragma omp end declare target |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,45 @@ | ||
/** | ||
* The first four lines test that a warning is produced when enabling | ||
* -Wopenmp-target-exception no matter what combination of -fexceptions and | ||
* -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the | ||
* target region but emit a warning instead. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
|
||
/** | ||
* The following four lines test that no warning is emitted when providing | ||
* -Wno-openmp-target-exception no matter the combination of -fexceptions and | ||
* -fcxx-exceptions. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
|
||
/** | ||
* Finally we should test that we only ignore exceptions in the OpenMP | ||
* offloading tool-chain | ||
*/ | ||
|
||
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o - | ||
|
||
// noexceptions-error@36 {{cannot use 'try' with exceptions disabled}} | ||
|
||
#pragma omp declare target | ||
int foo(void) { | ||
int error = -1; | ||
try { // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'catch' block is ignored}} | ||
error = 1; | ||
} | ||
catch (int e){ | ||
error = e; | ||
} | ||
return error; | ||
} | ||
#pragma omp end declare target | ||
// without-no-diagnostics |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,46 @@ | ||
/** | ||
* The first four lines test that a warning is produced when enabling | ||
* -Wopenmp-target-exception no matter what combination of -fexceptions and | ||
* -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the | ||
* target region but emit a warning instead. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
|
||
/** | ||
* The following four lines test that no warning is emitted when providing | ||
* -Wno-openmp-target-exception no matter the combination of -fexceptions and | ||
* -fcxx-exceptions. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
|
||
/** | ||
* Finally we should test that we only ignore exceptions in the OpenMP | ||
* offloading tool-chain | ||
*/ | ||
|
||
// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o - | ||
|
||
// noexceptions-error@37 {{cannot use 'try' with exceptions disabled}} | ||
// noexceptions-error@38 {{cannot use 'throw' with exceptions disabled}} | ||
|
||
#pragma omp declare target | ||
int foo(void) { | ||
int error = -1; | ||
try { // with-warning {{target 'nvptx64' does not support exception handling; 'catch' block is ignored}} | ||
throw 404; // with-warning {{target 'nvptx64' does not support exception handling; 'throw' is assumed to be never reached}} | ||
} | ||
catch (int e){ | ||
error = e; | ||
} | ||
return error; | ||
} | ||
#pragma omp end declare target | ||
// without-no-diagnostics |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,38 @@ | ||
/** | ||
* The first four lines test that a warning is produced when enabling | ||
* -Wopenmp-target-exception no matter what combination of -fexceptions and | ||
* -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the | ||
* target region but emit a warning instead. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze | ||
|
||
/** | ||
* The following four lines test that no warning is emitted when providing | ||
* -Wno-openmp-target-exception no matter the combination of -fexceptions and | ||
* -fcxx-exceptions. | ||
*/ | ||
|
||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze | ||
|
||
/** | ||
* Finally we should test that we only ignore exceptions in the OpenMP | ||
* offloading tool-chain | ||
*/ | ||
|
||
// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o - | ||
|
||
// noexceptions-error@35 {{cannot use 'throw' with exceptions disabled}} | ||
|
||
#pragma omp declare target | ||
void foo(void) { | ||
throw 404; // with-warning {{target 'nvptx64' does not support exception handling; 'throw' is assumed to be never reached}} | ||
} | ||
#pragma omp end declare target | ||
// without-no-diagnostics |
Oops, something went wrong.