Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

AOMP 17.0.3 crashes at compile time with conjunction in if clause #601

Closed
tom91136 opened this issue Aug 6, 2023 · 7 comments
Closed

AOMP 17.0.3 crashes at compile time with conjunction in if clause #601

tom91136 opened this issue Aug 6, 2023 · 7 comments
Assignees
Labels
bug Something isn't working

Comments

@tom91136
Copy link

tom91136 commented Aug 6, 2023

The following crashes the compiler:

#include <stdio.h>
int main(void) {
  bool a = true;
  bool b = true;
  int * p = nullptr;
  #pragma omp target data if(a && b) map(to: p[0])
  {
      printf("%p\n", p);
  }
  return 0;
}
/usr/lib/aomp_17.0-3/bin/clang++  -std=c++17 -O3 -Wall -march=native -fopenmp --offload-arch=gfx908  -fopenmp=libomp  boom.cpp
Basic Block in function 'main' does not have terminator!
label %land.end
fatal error: error in backend: Broken module found, compilation aborted!
clang++: error: clang frontend command failed with exit code 70 (use -v to see invocation)
AOMP_STANDALONE_17.0-3 clang version 17.0.0 (https://github.com/radeonopencompute/llvm-project f959ea5d8d1e5aef4b6d06727a9698316d3d33cd)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/lib/aomp_17.0-3/bin
clang++: note: diagnostic msg: Error generating preprocessed source(s).

Replacing the conjunction (i.e replace with if(a)) compiles successfully.

@gregrodgers Probably not the right place but I remember chatting with you during SC22 about AOMP and its improved reduction performance, as shown in BabelStream, is that available in any of the release already?

@dhruvachak
Copy link
Contributor

@tom91136 Thanks for the bug report.

Regarding your question about reduction, yes, it is available in the recent ROCm and AOMP releases. You can try it out using the clang compiler option -fopenmp-target-fast. Documentation about the Xteam-Reduction feature is at https://rocm.docs.amd.com/en/latest/reference/openmp/openmp.html#specialized-kernels

@tom91136
Copy link
Author

Also reproducible on AOMP18:

/usr/lib/aomp_18.0-0/bin/clang++  -std=c++17 -O3 -Wall -march=native -fopenmp --offload-arch=gfx908  -fopenmp=libomp  aomp_bad.cpp
Basic Block in function 'main' does not have terminator!
label %land.end
fatal error: error in backend: Broken module found, compilation aborted!
clang++: error: clang frontend command failed with exit code 70 (use -v to see invocation)
AOMP_STANDALONE_18.0-0 clang version 18.0.0 (https://github.com/radeonopencompute/llvm-project def7057717b5098f6a9f773fc6e7b2a7f59cdd50)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /usr/lib/aomp_18.0-0/bin
clang++: note: diagnostic msg: 
********************

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang++: note: diagnostic msg: /tmp/aomp_bad-gfx908-0593e3.cpp
clang++: note: diagnostic msg: /tmp/aomp_bad-84120f.cpp
clang++: note: diagnostic msg: /tmp/aomp_bad-gfx908-0593e3.sh
clang++: note: diagnostic msg: 

********************

@gregrodgers
Copy link
Contributor

I know it should not crash the compiler but I believe this is invalid code. You cannot put an if clause on target data. If you remove data, it almost works. You cannot send p[0] to device because it is a nullptr to nothing. The following test cases works.

#include <stdio.h>
#include <omp.h>
int main(void) {
int is_host, rc = 0;
int * p = &is_host ;

bool a = true;
bool b = true;
#pragma omp target if(a && b) map(to: p) map(from:is_host)
{
is_host=omp_is_initial_device();
printf("if(t&&t) threads:%d isInitial:%d p:%p\n", omp_get_num_threads(),is_host,(void*) p);
}
if(is_host) rc++;

a = true;
b = false;
#pragma omp target if(a && b) map(to: p) map(from:is_host)
{
is_host=omp_is_initial_device();
printf("if(t&&f) threads:%d isInitial:%d p:%p\n", omp_get_num_threads(),is_host,(void*) p);
}
if(!is_host) rc++;
a = false;
b = false;
#pragma omp target if(a && b) map(to: p) map(from:is_host)
{
is_host=omp_is_initial_device();
printf("if(f&&f) threads:%d isInitial:%d p:%p\n", omp_get_num_threads(),is_host,(void*) p);
}
if(!is_host) rc++;
return rc;
}

@tom91136
Copy link
Author

Yes, the snippet is just something that is short to trigger the issue, it was reduced from a TeaLeaf kernel.
I'm reading https://www.openmp.org/spec-html/5.0/openmpsu57.html and it seems to suggest an if clause is allowed for target data?

@ronlieb
Copy link
Contributor

ronlieb commented Sep 13, 2023

fail in build of trunk llvm as well

@ronlieb
Copy link
Contributor

ronlieb commented Sep 13, 2023

/home/rlieberm/rocm/trunk_1.0//bin/clang++ -O2 -fopenmp --offload-arch=gfx90a -D__OFFLOAD_ARCH_gfx90a__ issue-27.cpp -o issue-27
Basic Block in function 'main' does not have terminator!
label %land.end
fatal error: error in backend: Broken module found, compilation aborted!

@gregrodgers gregrodgers added the bug Something isn't working label Sep 26, 2023
@ddpagan ddpagan self-assigned this Oct 24, 2023
@ronlieb
Copy link
Contributor

ronlieb commented Nov 21, 2023

Dave fixed it, should be in 18.0-1

@ronlieb ronlieb closed this as completed Nov 21, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working
Projects
None yet
Development

No branches or pull requests

5 participants