Authored by vzakhari on Mar 29 2021, 4:13 PM.
Tags
• Restricted Project
• Restricted Project
• Restricted Project

# Details

Reviewers
 ABataev sdmitriev grokos alexander-shaposhnikov jdoerfert
Commits
Summary

This patch is a piece of D99360.

The new notes use a null-terminated "LLVMOMPOFFLOAD" note name. There are currently three types of notes:

VERSION: a string (not null-terminated) representing the ELF offload image structure. The current version '1.0' does not put any restrictions on the structure of the image. If we ever need to come up with a common structure for ELF offload images (e.g. to be able to analyze the images in libomptarget in some standard way), then we will introduce new versions.
PRODUCER: a vendor specific name of the producing toolchain. Upstream LLVM uses "LLVM" (not null-terminated).
PRODUCER_VERSION: a vendor specific version of the producing toolchain. Upstream LLVM uses LLVM_VERSION_STRING with optional <space> LLVM_REVISION.

All three notes are not mandatory currently.

# Unit TestsFailedView All

TimeTest
Script: -- : 'RUN: at line 6'; clang-offload-wrapper --help | /mnt/disks/ssd0/agent/llvm-project/build/bin/FileCheck /mnt/disks/ssd0/agent/llvm-project/clang/test/Driver/clang-offload-wrapper.c --check-prefix CHECK-HELP
40 msx64 debian > LLVM.tools/dsymutil/X86::label2.test
Script: -- : 'RUN: at line 14'; /mnt/disks/ssd0/agent/llvm-project/build/bin/dsymutil -oso-prepend-path /mnt/disks/ssd0/agent/llvm-project/llvm/test/tools/dsymutil/X86/../Inputs /mnt/disks/ssd0/agent/llvm-project/llvm/test/tools/dsymutil/X86/../Inputs/private/tmp/label/label.out -o /mnt/disks/ssd0/agent/llvm-project/build/test/tools/dsymutil/X86/Output/label2.test.tmp.dSYM
Script: -- : 'RUN: at line 6'; clang-offload-wrapper --help | c:\ws\w16-1\llvm-project\premerge-checks\build\bin\filecheck.exe C:\ws\w16-1\llvm-project\premerge-checks\clang\test\Driver\clang-offload-wrapper.c --check-prefix CHECK-HELP
120 msx64 windows > LLVM.tools/dsymutil/X86::label2.test
Script: -- : 'RUN: at line 14'; c:\ws\w16-1\llvm-project\premerge-checks\build\bin\dsymutil.exe -oso-prepend-path C:\ws\w16-1\llvm-project\premerge-checks\llvm\test\tools\dsymutil\X86/../Inputs C:\ws\w16-1\llvm-project\premerge-checks\llvm\test\tools\dsymutil\X86/../Inputs/private/tmp/label/label.out -o C:\ws\w16-1\llvm-project\premerge-checks\build\test\tools\dsymutil\X86\Output\label2.test.tmp.dSYM

### Event Timeline

vzakhari created this revision.Mar 29 2021, 4:13 PM
vzakhari requested review of this revision.Mar 29 2021, 4:13 PM
Herald added projects: Restricted Project, Restricted Project.
Herald added subscribers: llvm-commits, cfe-commits, sstefan1.
vzakhari updated this revision to Diff 334282.Mar 30 2021, 3:22 PM

I am not sure what is wrong with the uploaded ELF files. The test works locally. Let's try it with a thinner YAML template.

vzakhari updated this revision to Diff 334292.Mar 30 2021, 3:59 PM

Updated revision fixes BUILD_SHARED_LIBS build.

Herald added a subscriber: mgorny. Mar 30 2021, 3:59 PM
grokos accepted this revision.Apr 6 2021, 7:19 PM

Change looks good, so it's accepted on my end. I'll let the other reviewers have a look and post their comments. Please do not commit until we have reached an agreement for all 4 patches together (D99551, D99552, D99553, D99612).

This revision is now accepted and ready to land.Apr 6 2021, 7:19 PM
vzakhari updated this revision to Diff 364193.Aug 4 2021, 11:10 AM
This revision was landed with ongoing or failed builds.Aug 16 2021, 1:20 PM
This revision was automatically updated to reflect the committed changes.

#include <omp.h>

int main() {
int isHost;
#pragma omp target map(from:isHost)
{ isHost = omp_is_initial_device(); }
return isHost;
}

Compiling and running with clang -fopenmp -fopenmp-targets=nvptx64 gives me the following results.

CUDA error: Error returned from cuModuleLoadDataEx
CUDA error: device kernel image is invalid
Libomptarget error: Unable to generate entries table for device id 0.
Libomptarget error: Failed to init globals on device 0
Libomptarget error: Run with LIBOMPTARGET_INFO=4 to dump host-target pointer mappings.
Libomptarget fatal error 1: failure of target construct while offloading is mandatory

Reverting 93d08acaacec951dbb302f77eeae51974985b6b2 fixes it.

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

Unfortunately, I did not receive it. Can you please check if the mail was blocked on your side? Can you please try to archive it with a password and send it again?
I would like to prepare a patch to put this functionality under a switch that is off by default. I think we want to have this working eventually (and it actually works with x86_64 offload), so it will be easier to just flip a switch, when I figure out what CUDA API does not like about the modified ELF image. Does it sound appropriate to you?

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

Unfortunately, I did not receive it. Can you please check if the mail was blocked on your side? Can you please try to archive it with a password and send it again?
I would like to prepare a patch to put this functionality under a switch that is off by default. I think we want to have this working eventually (and it actually works with x86_64 offload), so it will be easier to just flip a switch, when I figure out what CUDA API does not like about the modified ELF image. Does it sound appropriate to you?

I put it in a .tar.gz file, I think your domain is blocking it. Anything else I can use?

@jhuber6, thank you for reporting this. I do not have a properly setup CUDA system currently. Can you please invoke clang-offload-wrapper with -save-temps and send the temporary files to me?

I sent it to you, any luck or should we revert this upstream for the time being.

Unfortunately, I did not receive it. Can you please check if the mail was blocked on your side? Can you please try to archive it with a password and send it again?
I would like to prepare a patch to put this functionality under a switch that is off by default. I think we want to have this working eventually (and it actually works with x86_64 offload), so it will be easier to just flip a switch, when I figure out what CUDA API does not like about the modified ELF image. Does it sound appropriate to you?

I put it in a .tar.gz file, I think your domain is blocking it. Anything else I can use?