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.
While I have almost not modified any code, I have made tests that verify that we still do not allow compilation of exceptions, unless the user explicitly enables -fcxx-exceptions or -fexceptions.
Please let me know what you think of this patch and if the warnings could be improved.
Example
With this patch, the compilation of the following example
{C++}
#include <iostream>
#pragma omp declare target
int gaussian_sum(int a,int b){
if ((a + b) % 2 == 0) {throw -1;};
return (a+b) * ((a+b)/2);
}
#pragma omp end declare target
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)Issues
The patch make 11 of the tests from clang/test/OpenMP fail. It seems to be related only with the values of Opts.Exceptions and Opts.CXXExceptions. I have tested that this change alone breaks the aforementioned tests. It would
be nice if somebody with a better knowledge of Clang and NVPTX would help me understand to what degree I have done something wrong, and when, if in any cases, it would be allowed to modify the tests.
Check the style of other messages, they have a single sentence and start with a lower case letter.
Also, the explanation doesn't make sense for users. They don't know about traps, basic blocks, etc.
Maybe:
target '%0' does not support exception handling; 'throw' is assumed to be never reached
target '%0' does not support exception handling; 'catch' block is ignored
and nothing for try.
Finally, these need a group so users can also disable them.