3c9988f85d
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 NVPTX, we generate code for the try statement as if it were a basic block.
With this patch, the compilation of the following example
```
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
```
./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
104 lines
3.0 KiB
C++
104 lines
3.0 KiB
C++
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown \
|
|
// RUN: -verify=host -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc \
|
|
// RUN: %s -o %t-ppc-host.bc -fexceptions -fcxx-exceptions
|
|
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown \
|
|
// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s \
|
|
// RUN: -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - \
|
|
// RUN: -fexceptions -fcxx-exceptions -ferror-limit 100
|
|
|
|
#ifndef HEADER
|
|
#define HEADER
|
|
|
|
template <typename T>
|
|
class TemplateClass {
|
|
T a;
|
|
public:
|
|
TemplateClass() { throw 1;}
|
|
T f_method() const { return a; }
|
|
};
|
|
|
|
int foo();
|
|
|
|
int baz1();
|
|
|
|
int baz2();
|
|
|
|
int baz4() { return 5; }
|
|
|
|
template <typename T>
|
|
T FA() {
|
|
TemplateClass<T> s;
|
|
return s.f_method();
|
|
}
|
|
|
|
#pragma omp declare target
|
|
struct S {
|
|
int a;
|
|
S(int a) : a(a) { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
|
|
};
|
|
|
|
int foo() { return 0; }
|
|
int b = 15;
|
|
int d;
|
|
#pragma omp end declare target
|
|
int c;
|
|
|
|
int bar() { return 1 + foo() + bar() + baz1() + baz2(); } // expected-note {{called by 'bar'}}
|
|
|
|
int maini1() {
|
|
int a;
|
|
static long aa = 32;
|
|
try {
|
|
#pragma omp target map(tofrom \
|
|
: a, b)
|
|
{
|
|
// expected-note@+1 {{called by 'maini1'}}
|
|
S s(a);
|
|
static long aaa = 23;
|
|
a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); // expected-note{{called by 'maini1'}}
|
|
if (!a)
|
|
throw "Error"; // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
|
|
}
|
|
} catch(...) {
|
|
}
|
|
return baz4();
|
|
}
|
|
|
|
int baz3() { return 2 + baz2(); }
|
|
int baz2() {
|
|
#pragma omp target
|
|
try { // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'catch' block is ignored}}
|
|
++c;
|
|
} catch (...) {
|
|
}
|
|
return 2 + baz3();
|
|
}
|
|
|
|
int baz1() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
|
|
|
|
int foobar1();
|
|
int foobar2();
|
|
|
|
int (*A)() = &foobar1;
|
|
#pragma omp declare target
|
|
int (*B)() = &foobar2;
|
|
#pragma omp end declare target
|
|
|
|
int foobar1() { throw 1; }
|
|
int foobar2() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
|
|
|
|
|
|
int foobar3();
|
|
int (*C)() = &foobar3; // expected-warning {{declaration is not declared in any declare target region}}
|
|
// host-warning@-1 {{declaration is not declared in any declare target region}}
|
|
#pragma omp declare target
|
|
int (*D)() = C; // expected-note {{used here}}
|
|
// host-note@-1 {{used here}}
|
|
#pragma omp end declare target
|
|
int foobar3() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
|
|
|
|
// Check no infinite recursion in deferred diagnostic emitter.
|
|
long E = (long)&E;
|
|
|
|
#endif // HEADER
|