[PATCH] D153924: Allowing exception handling in OpenMP target regions when offloading to AMDGCN or NVPTX targets

[Anton Rydahl via Phabricator via cfe-commits]

AntonRydahl created this revision.
AntonRydahl added reviewers: jdoerfert, jhuber6.
Herald added subscribers: sunshaoce, mattd, gchakrabarti, asavonic, guansong, tpr, yaxunl, rampitec.
Herald added a project: All.
AntonRydahl requested review of this revision.
Herald added subscribers: cfe-commits, jplehr, sstefan1.
Herald added a project: clang.

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.Exceptions`. 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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D153924

Files:
  clang/include/clang/Basic/DiagnosticCommonKinds.td
  clang/lib/CodeGen/CGException.cpp
  clang/lib/Frontend/CompilerInvocation.cpp
  clang/test/OpenMP/target_throw_message.cpp
  clang/test/OpenMP/target_throw_message_fun_call.cpp
  clang/test/OpenMP/target_try_catch_message.cpp
  clang/test/OpenMP/target_try_catch_message_fun_call.cpp
  clang/test/OpenMP/target_try_catch_throw_message.cpp
  clang/test/OpenMP/target_try_catch_throw_message_no_exceptions.cpp

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D153924.535148.patch
Type: text/x-patch
Size: 15847 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230627/974c3b8e/attachment-0001.bin>