[OpenMP][DeviceRTL][AMDGPU] Support code object version 5

[saiislam]

Summary

Update DeviceRTL and the AMDGPU plugin to use code
object version 5. Default is code object version 4.

DeviceRTL uses rocm-device-libs instead of directly calling
amdgcn builtins for the functions which are affected by
cov5.

AMDGPU plugin queries the ELF for code object version
and then prepares various implicitargs accordingly.

[jhuber6]

Maybe we should wait until D138389 lands and we can update both, otherwise we'd need a second patch.

[jhuber6]

I'm not fully up-to-date, what's the main difference and advantage of the new code object version? What do all the new implicit arguments do.

clang/lib/Driver/ToolChains/AMDGPU.cpp

953

Unrelated?

openmp/libomptarget/DeviceRTL/include/Interface.h

169

This should probably use variants to match the rest of the style, also if you intend to read these outside of the library you'll need to put them in the exports file and set their visibility.

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

19

What if this isn't defined? We should be able to use the OpenMP library without the AMD device libraries. Should it be extern weak?

openmp/libomptarget/DeviceRTL/src/State.cpp

73

Variants

openmp/libomptarget/plugins/amdgpu/impl/get_elf_mach_gfx_name.h

15

Unrelated, but is there any particular reason these aren't defined in the hsa_amd_ext.h?

[yaxunl]

clang/lib/Driver/ToolChains/Clang.cpp

7085

Any reason you need the original args? This will bypass the driver translation, which should not in normal cases.

7086

clang -cc1 needs this to be default value false to emit code object version module flag

[jdoerfert]

Could we elaborate on the benefits, please. Now we support two versions?

Why is this helpful:

DeviceRTL uses rocm-device-libs instead of directly calling amdgcn builtins for the functions which are affected by cov5.

openmp/libomptarget/DeviceRTL/src/State.cpp

81

Why do we need the "external..." stuff anyway?

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

2233

What is this all about?

[JonChesterfield]

I am reluctant to add the dependency edge to rocm device libs to openmp's GPU runtime.

We currently require that library for libm, which I'm also not thrilled about, but at least you can currently build and run openmp programs (that don't use libm, like much of our tests) without it.

My problem with device libs is that it usually doesn't build with trunk. It follows a rolling dev model tied to rocm clang and when upstream does something that takes a long time to apply, device libs doesn't build until rocm catches up. I've literally never managed to compile any branch of device libs with trunk clang without modifying the source, generally to delete directories that don't look necessary for libm.

Further, selecting an ABI based on runtime code found in a library which is hopefully constant folded is a weird layering choice. The compiler knows what ABI it is emitting code for, and that's how it picks files from device libs to effect that choice, but it would make far more sense to me for the compiler back end to set this stuff up itself.

Also, if we handle ABI in the back end, then we don't get the inevitable problem of rocm device libs and trunk clang having totally different ideas of what the ABI is as they drift in and out of sync.

[tianshilei1992]

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

19

It should be put into AMD's declare variant.

[saiislam]

Thanks everyone for your review and comments!
I am going to address all of them in a series of smaller patches starting with D140784.

[saiislam]

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

50

If we still don't want to depend on rocm-device-libs then we will have to do something like (haven't tried this code yet):

uint32_t getNumHardwareThreadsInBlock() {
   if (__oclc_ABI_version < 500) {
      return __builtin_amdgcn_workgroup_size_x();
   } else {
      void *implicitArgPtr = __builtin_amdgcn_implicitarg_ptr();
      return (ushort)implicitArgPtr[6];
}

80

uint32_t getNumberOfBlocks() {
   if (__oclc_ABI_version < 500) {
      return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
   } else {
      void *implicitArgPtr = __builtin_amdgcn_implicitarg_ptr();
      return (uint)implicitArgPtr[0];
}

[saiislam]

clang/lib/Driver/ToolChains/Clang.cpp

7085

We need derived args to look for mcode-object-version. I have created a separate review for this change. Please have a look at D142022

[arsenm]

In D139730#3991628, @JonChesterfield wrote:

We currently require that library for libm, which I'm also not thrilled about, but at least you can currently build and run openmp programs (that don't use libm, like much of our tests) without it.

The ABI isn't defined in terms of what device-libs does. It's fixed offsets off of pointers accessible through amdgcn intrinsics. You can also just directly emit the same IR, these functions aren't complicated

[jhuber6]

In D139730#4418630, @arsenm wrote:

In D139730#3991628, @JonChesterfield wrote:

We currently require that library for libm, which I'm also not thrilled about, but at least you can currently build and run openmp programs (that don't use libm, like much of our tests) without it.

The ABI isn't defined in terms of what device-libs does. It's fixed offsets off of pointers accessible through amdgcn intrinsics. You can also just directly emit the same IR, these functions aren't complicated

This is the suggestion I've talked with @saiislam about. I think we should just copy the magic intrinsics that are being queried here. I'm assuming we don't need to bother with supporting both v4 and v5 so we can just make the switch all at once.

[saiislam]

Another attempt at cov5 support by using CodeGen for buitlin_amdgpu_workgroup_size.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

17146 ↗

(On Diff #547297)

this must always pass

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

3037 ↗

(On Diff #547297)

This isn't doing anything?

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #547297)

This is getting duplicated a few places, should it move to a support header?

I don't love the existing APIs for this, I think a struct definition makes more sense

[jhuber6]

Could you explain briefly what the approach here is? I'm confused as to what's actually changed and how we're handling this difference. I thought if this was just the definition of some builtin function we could just rely on the backend to figure it out. Why do we need to know the code object version inside the device RTL?

clang/lib/CodeGen/CGBuiltin.cpp

17140 ↗

(On Diff #547297)

Could you explain the function of this in a comment? Are we emitting generic code if unspecified?

17187–17188 ↗

(On Diff #547297)

nit.

17194 ↗

(On Diff #547297)

Leftover debugging?

clang/lib/Driver/ToolChain.cpp

1360 ↗

(On Diff #547297)

Shouldn't we be able to put this under the OPT_m_group below?

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

1752 ↗

(On Diff #547297)

Leftoever?

2542 ↗

(On Diff #547297)

Why do we need this? The current method shouldn't need to change if all we're doing is allocating memory of greater size.

3036 ↗

(On Diff #547297)

So we're required to emit some new arguments? I don't have any idea what'schanged between this COV4 and COV5 stuff.

[jhuber6]

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #547297)

The other user here is my custom loader, @JonChesterfield has talked about wanting a common HSA helper header for awhile now.

I agree that the struct definition is much better. Being able to simply allocate this size and then zero fill it is much cleaner.

[arsenm]

In D139730#4561540, @jhuber6 wrote:

Could you explain briefly what the approach here is? I'm confused as to what's actually changed and how we're handling this difference. I thought if this was just the definition of some builtin function we could just rely on the backend to figure it out. Why do we need to know the code object version inside the device RTL?

The builtin is called in the device rtl, so the device RTL needs to contain both implementations. The "backend figuring it out" is dead code elimination

[jhuber6]

In D139730#4561573, @arsenm wrote:

In D139730#4561540, @jhuber6 wrote:

Could you explain briefly what the approach here is? I'm confused as to what's actually changed and how we're handling this difference. I thought if this was just the definition of some builtin function we could just rely on the backend to figure it out. Why do we need to know the code object version inside the device RTL?

The build is called in the device rtl, so the device RTL needs to contain both implementations. The "backend figuring it out" is dead code elimination

Okay, do we expect to re-use this interface anywhere? If it's just for OpenMP then we should probably copy the approach taken for __omp_rtl_debug_kind, which is a global created on the GPU by CGOpenMPRuntimeGPU's constructor and does more or less the same thing.

[arsenm]

In D139730#4561575, @jhuber6 wrote:

In D139730#4561573, @arsenm wrote:

In D139730#4561540, @jhuber6 wrote:

Could you explain briefly what the approach here is? I'm confused as to what's actually changed and how we're handling this difference. I thought if this was just the definition of some builtin function we could just rely on the backend to figure it out. Why do we need to know the code object version inside the device RTL?

The build is called in the device rtl, so the device RTL needs to contain both implementations. The "backend figuring it out" is dead code elimination

Okay, do we expect to re-use this interface anywhere? If it's just for OpenMP then we should probably copy the approach taken for __omp_rtl_debug_kind, which is a global created on the GPU by CGOpenMPRuntimeGPU's constructor and does more or less the same thing.

device libs replicates the same scheme using its own copy of an equivalent variable. Trying to merge those two together

[arsenm]

In D139730#4561619, @arsenm wrote:

In D139730#4561575, @jhuber6 wrote:

In D139730#4561573, @arsenm wrote:

In D139730#4561540, @jhuber6 wrote:

Could you explain briefly what the approach here is? I'm confused as to what's actually changed and how we're handling this difference. I thought if this was just the definition of some builtin function we could just rely on the backend to figure it out. Why do we need to know the code object version inside the device RTL?

The build is called in the device rtl, so the device RTL needs to contain both implementations. The "backend figuring it out" is dead code elimination

Okay, do we expect to re-use this interface anywhere? If it's just for OpenMP then we should probably copy the approach taken for __omp_rtl_debug_kind, which is a global created on the GPU by CGOpenMPRuntimeGPU's constructor and does more or less the same thing.

device libs replicates the same scheme using its own copy of an equivalent variable. Trying to merge those two together

Although I guess that doesn't really need the builtin changes?

[saiislam]

Removed unused cov5 implicitargs fields.
Added comments about EmitAMDGPUWorkGroupSize and ABI-agnostica code emission.
Adressed reviewers' comments.

[saiislam]

In D139730#4561622, @arsenm wrote:

In D139730#4561619, @arsenm wrote:

In D139730#4561575, @jhuber6 wrote:

In D139730#4561573, @arsenm wrote:

In D139730#4561540, @jhuber6 wrote:

Could you explain briefly what the approach here is? I'm confused as to what's actually changed and how we're handling this difference. I thought if this was just the definition of some builtin function we could just rely on the backend to figure it out. Why do we need to know the code object version inside the device RTL?

The build is called in the device rtl, so the device RTL needs to contain both implementations. The "backend figuring it out" is dead code elimination

Okay, do we expect to re-use this interface anywhere? If it's just for OpenMP then we should probably copy the approach taken for __omp_rtl_debug_kind, which is a global created on the GPU by CGOpenMPRuntimeGPU's constructor and does more or less the same thing.

device libs replicates the same scheme using its own copy of an equivalent variable. Trying to merge those two together

Although I guess that doesn't really need the builtin changes?

This builtin was already aware about cov4 and cov5. All this patch is changing is making it aware about a possibility where both needs to be present.
It is already used by device-libs, deviceRTL, and libc-gpu.
Also, encapsulating ABI related changes in implementation of the builtin allows other runtime developers to be agnostic to these lower level changes.

clang/lib/CodeGen/CGBuiltin.cpp

17187–17188 ↗

(On Diff #547297)

There are a couple of common lines after the inner if-else, in the outer else section.

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

1752 ↗

(On Diff #547297)

No, it is not a left over.
One of the fields in cov5 implicitikernarg is heap_v1 ptr. It should point to a 128KB zero-initialized block of coarse-grained memory on each device before launching the kernel. This code was working a while ago, but right now it is failing most likely due to some latest change in devicertl memory handling mechanism.
I need to debug it with this patch, otherwise it will cause all target region code calling device-malloc to fail.
I will try to fix it before the next revision.

2542 ↗

(On Diff #547297)

PreAllocatedDeviceMemoryPool is the pointer which stores the intermediate value before it is written to heap_v1_ptr field of cov5 implicitkernarg.

3036 ↗

(On Diff #547297)

In cov5, we need to set certain fields of the implicit kernel arguments before launching the kernel.
Please see AMDHSA Code Object V5 Kernel Argument Metadata Map Additions and Changes for more details.

Only NumBlocks, NumThreads(XYZ), GridDims, and Heap_V1_ptr are relevant for us, so I have simplified code further.

3037 ↗

(On Diff #547297)

Earlier we used to set hostcall_buffer here, but not anymore.
I have left the message in DP just for debug help.

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #547297)

Defining a struct for whole 256 byte of implicitargs in cov5 was becoming a little difficult due to different sizes of various fields (2, 4, 6, 8, 48, 72 bytes) along with multiple reserved fields in between. It made sense for cov4 because it only had 7 fields of 8 bytes each, where we needed only 4th field in OpenMP runtime (for hostcall_buffer).

Offset based lookups like the following allows handling/exposing only required fields across generations of ABI.

[jhuber6]

clang/lib/CodeGen/CGBuiltin.cpp

17187–17188 ↗

(On Diff #547297)

You should be able to factor out

LD = CGF.Builder.CreateLoad(
    Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2)));

from both by making each assign the Result to a value.

17138 ↗

(On Diff #547751)

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

2550–2551 ↗

(On Diff #547751)

This and below isn't correct. You can't discard an llvm::Error value like this without either doing consumeError(std::move(Err)) or toString(std::move(Err)). However, you don't need to consume these in the first place, they already contain the error message from the callee and should just be forwarded.

1752 ↗

(On Diff #547297)

Do we really need that? We only use a fraction of the existing implicit arguments. My understanding is that most of these are more for runtime handling for HIP and OpenCL while we would most likely want our own solution. I'm assuming that the 128KB is not required for anything we use?

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #547297)

If we don't use it, just put it as unused. It's really hard to read as-is and it makes it more difficult to just zero fill.

[yaxunl]

need a lit test for the codegen of the clang builtin for cov 4/5/none and a lit test to show the branching code generated with cov none can be optimized away when linked with cov4 or cov5.

clang/lib/CodeGen/Targets/AMDGPU.cpp

383 ↗

(On Diff #547751)

I am not sure weak_odr linkage will work when code object version is none. This will cause conflict when a module emitted with cov none is linked with a module emitted with cov4 or cov5. Also, when all modules are emitted with cov none, we end up with a linked module with cov none and the work group size code will not work.

Probably we need to emit llvm.amdgcn.abi.version with external linkage for cov none.

Another issue is that llvm.amdgcn.abi.version is not internalized. It is always loaded from memory even though it is in constant address space. This will cause bad performance. Considering device libs may use clang builtin for workgroup size. The performance impact may be significant. To avoid performance degradation, we need to internalize it as early as possible in the optimization pipeline.

[yaxunl]

I would suggest separating the clang/llvm part into a separate review.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

17143–17145 ↗

(On Diff #547751)

Move down to define and initialize

17163–17165 ↗

(On Diff #547751)

You could write all of this in terms of selects and avoid introducing all these blocks

clang/lib/CodeGen/Targets/AMDGPU.cpp

358 ↗

(On Diff #547751)

Don't need this?

[saiislam]

Updated the patch as per reviewers comments.

clang/lib/CodeGen/CGBuiltin.cpp

17143–17145 ↗

(On Diff #547751)

There are multiple uses of the same identifier. Defining them four times looks odd.

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

2550–2551 ↗

(On Diff #547751)

Removed the logic for preallocatedheap.

1752 ↗

(On Diff #547297)

I have removed the preallocatedheap work from this patch.

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #547297)

I have reduced the fields to bare minimum required for OpenMP.

[jhuber6]

Some nits. I'm assuming we're getting the code object in the backend now? We'll need to make sure that -Wl,--amdhsa-code-object-version is passed to the clang invocation inside of the clang-linker-wrapper to handle -save-temps mode.

clang/lib/CodeGen/CGBuiltin.cpp

17053 ↗

(On Diff #551266)

clang/lib/Driver/ToolChain.cpp

1363 ↗

(On Diff #551266)

Random whitespace.

clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu

54 ↗

(On Diff #551266)

Need newline

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

3016 ↗

(On Diff #551266)

Don't think this needs to be a debug message, same below

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #551266)

I'm still not a fan of replacing the struct. The mnemonic of having a struct is much more user friendly.

ImplicitArgsTy Args{};
std::memset(&Args, sizeof(ImplicitArgsTy), 0);
...

If we don't use something, just make it some random bytes, e.g.

struct ImplicitArgsTy {
  uint64_t OffsetX;
  uint8_t Unused[64]; // 64 byte offset.
};

[saiislam]

Changed ImplitArgs implementation using struct.

[saiislam]

In D139730#4597504, @jhuber6 wrote:

Some nits. I'm assuming we're getting the code object in the backend now? We'll need to make sure that -Wl,--amdhsa-code-object-version is passed to the clang invocation inside of the clang-linker-wrapper to handle -save-temps mode.

Clang-linker-wrapper was not passing -mllvm option to the clang backend.

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

36 ↗

(On Diff #551266)

Replaced.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

17057 ↗

(On Diff #551266)

Spell out to DispatchPtr?

clang/lib/CodeGen/CodeGenModule.cpp

1206–1208 ↗

(On Diff #551266)

These could be one combined hook? this isn't really different from metadata

clang/lib/CodeGen/Targets/AMDGPU.cpp

369–386 ↗

(On Diff #551266)

You moved GetOrCreateLLVMGlobal but don't use it?

The lamdba is unnecessary for a single local use

clang/lib/Driver/ToolChain.cpp

1369 ↗

(On Diff #551266)

Capitalize

1372 ↗

(On Diff #551266)

Don't understand why this is necessary

[arsenm]

clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu

40–43 ↗

(On Diff #551597)

test all the builtins?

[saiislam]

clang/lib/CodeGen/Targets/AMDGPU.cpp

383 ↗

(On Diff #547751)

I tried external linkage but it didn't work. Only weak_odr is working fine.

369–386 ↗

(On Diff #551266)

I am using GetOrCreateLLVMGlobal in CGBuiltin.cpp while emitting code for amdgpu_worgroup_size.

369–386 ↗

(On Diff #551266)

I was hoping that this patch will pave way for D130096, so that it can generate rest of the control constants using the same lambda.
I can remove this and simplify the code if you want.

clang/lib/Driver/ToolChain.cpp

1372 ↗

(On Diff #551266)

This function creates a derived argument list for OpenMP target specific flags.
mcode-object-version remains unset for device compilation step if we don't pass it here.

[saiislam]

Adressed reviewer's comments.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

17067 ↗

(On Diff #552085)

Capitalization is weird, IsCOV5?

17082–17083 ↗

(On Diff #552085)

CreateConstInBoundsGEP1_64

17100 ↗

(On Diff #552085)

CreateConstInBoundsGEP1_64

clang/lib/CodeGen/Targets/AMDGPU.cpp

364 ↗

(On Diff #552085)

Single use lamdba, just make this the function body

381 ↗

(On Diff #552085)

No real point setting the alignment

[saiislam]

Used CreateConstInBoundsGEP1_32 for emitting GEP statements. Changed lambda function to simple fucntion body for defining the global variable.

[arsenm]

Codegen parts LGTM, questions with the driver parts

clang/lib/Driver/ToolChain.cpp

1368 ↗

(On Diff #552344)

Typos

1369 ↗

(On Diff #552344)

clang/lib/Driver/ToolChains/Clang.cpp

8368–8369

so device rtl is linked once as a normal library?

8372–8373

Why do you need this? The code object version is supposed to come from a module flag. We should be getting rid of the command line argument for it

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

406–410 ↗

(On Diff #552344)

Shouldn't need this?

417 ↗

(On Diff #552344)

Commented out code

[yaxunl]

clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu

12 ↗

(On Diff #552344)

need to test using clang -cc1 with -O3 and -mlink-builtin-bitcode to link the device lib and verify the load of llvm.amdgcn.abi.version being eliminated after optimization.

I think currently it cannot do that since llvm.amdgcn.abi.version is not internalized by the internalization pass. This can cause some significant perf drops since loading is expensive. Need to tweak the function controlling what variables can be internalized for amdgpu so that this variable gets internalized, or having a generic way to tell that function which variables should be internalized, e.g. by adding a metadata amdgcn.internalize

[saiislam]

Updated test case to check internalization of newly inserted global variable.

clang/lib/Driver/ToolChains/Clang.cpp

8368–8369

No, this is command generation for clang-linker-wrapper. Since, devicertl is compiled only to get bitcode file (-c), it is never called.

8372–8373

During command generation for clang-linker-wrapper, it is required to check user's provided mcode-object-version=X so that amdhsa-code-object-version=X can be passed to the clang/lto backend.

getAmdhsaCodeObjectVersion() and getHsaAbiVersion() both still use the above command line argument to override user's choice of COV, instead of the module flag.

clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu

12 ↗

(On Diff #552344)

load of llvm.amdgcn.abi.version is being eliminated with cc1, -O3, and mlink-builtin-bitcode of device lib.

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

406–410 ↗

(On Diff #552344)

It is required so that when clang pass (not the lto backend) is called from clang-linker-wrapper due to -save-temps, user provided COV is correctly propagated.

[jhuber6]

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

49 ↗

(On Diff #552085)

We should probably be using sizeof now that it's back to being a struct and keep the old struct definition.

[saiislam]

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

49 ↗

(On Diff #552085)

AMDGPU plugin doesn't use any implicitarg for COV4, but it does so for COV5. So, we are not keeping two separate structures for implicitargs of COV4 and COV5.
If we use sizeof then it will always return 256 corresponding to COV5 (even for cov4, which should be 56). That's why we need this function.

[jhuber6]

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

49 ↗

(On Diff #552085)

Yeah, I guess for COV4 the only thing that mattered was the size so that we could make sure it's all set to zero. We shouldn't use the enum value. It should be sizeof(ImplicitArgsTy) for COV5 and either hard-code it in the function for V4 or make a dummy struct.

[saiislam]

Changed getImplicitArgsSize to use sizeof.

[jhuber6]

Just a few more nits. I think it's looking fine but I haven't tested it. Anyone else?

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

406 ↗

(On Diff #553413)

415–417 ↗

(On Diff #553413)

No braces around a single line if.

openmp/libomptarget/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h

54 ↗

(On Diff #553413)

We return uint16_t here? These are sizes.

[saiislam]

Minor fixes addressing reviewer's comment.

[jhuber6]

I think it's fine now given that it's passing tests. Others feel free to comment.

[yaxunl]

LGTM. Thanks

clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu

12 ↗

(On Diff #552344)

It seems being eliminated by IPSCCP. It makes sense since it is constant weak_odr without externally_initialized. Either changing it to weak or adding externally_initialized will keep the load. Normal __constant__ var in device code may be changed by host code, therefore they are emitted with externally_initialized and do not have the load eliminated.

[saiislam]

clang/test/CodeGenCUDA/amdgpu-code-object-version-linking.cu

12 ↗

(On Diff #552344)

Thank you @yaxunl !
I have added these observations as comments in the code at load emit and global emit locations.

clang/lib/Driver/Driver.cpp

for (StringRef Val : OpenMPTriples) {
           auto &DeviceTC =
               ToolChains[TT.str() + "/" + HostTC->getTriple().normalize()];
           if (!DeviceTC) {
             if (TT.isNVPTX())
               DeviceTC = std::make_unique<toolchains::CudaToolChain>(
                   *this, TT, *HostTC, C.getInputArgs());
             else if (TT.isAMDGCN())
               DeviceTC = std::make_unique<toolchains::AMDGPUOpenMPToolChain>(
-                  *this, TT, *HostTC, C.getInputArgs());
+                  *this, TT, *HostTC, C.getInputArgs(), C.getArgs());
             else
               assert(DeviceTC && "Device toolchain not defined.");
           }
 
           TC = DeviceTC.get();
         } else
           TC = &getToolChain(C.getInputArgs(), TT);
         C.addOffloadDeviceToolChain(TC, Action::OFK_OpenMP);

clang/lib/Driver/ToolChains/AMDGPU.cpp

ROCMToolChain::getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs,
                                        const std::string &GPUArch,
                                        bool isOpenMP) const {
   auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch);
   const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
 
   StringRef LibDeviceFile = RocmInstallation.getLibDeviceFile(CanonArch);
   auto ABIVer = DeviceLibABIVersion::fromCodeObjectVersion(
       getAMDGPUCodeObjectVersion(getDriver(), DriverArgs));
+
   if (!RocmInstallation.checkCommonBitcodeLibs(CanonArch, LibDeviceFile,
                                                ABIVer))
     return {};
 
   // If --hip-device-lib is not set, add the default bitcode libraries.
   // TODO: There are way too many flags that change this. Do we need to check
   // them all?
   bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero,

clang/lib/Driver/ToolChains/AMDGPUOpenMP.h

 }
 
 namespace toolchains {
 
 class LLVM_LIBRARY_VISIBILITY AMDGPUOpenMPToolChain final
     : public ROCMToolChain {
 public:
   AMDGPUOpenMPToolChain(const Driver &D, const llvm::Triple &Triple,
-                        const ToolChain &HostTC,
-                        const llvm::opt::ArgList &Args);
+                        const ToolChain &HostTC, const llvm::opt::ArgList &Args,
+                        const llvm::opt::DerivedArgList &DerivedArgs);
 
   const llvm::Triple *getAuxTriple() const override {
     return &HostTC.getTriple();
   }
 
   llvm::opt::DerivedArgList *
   TranslateArgs(const llvm::opt::DerivedArgList &Args, StringRef BoundArch,
                 Action::OffloadKind DeviceOffloadKind) const override;
public:
   VersionTuple
   computeMSVCVersion(const Driver *D,
                      const llvm::opt::ArgList &Args) const override;
 
   llvm::SmallVector<BitCodeLibraryInfo, 12>
   getDeviceLibs(const llvm::opt::ArgList &Args) const override;
 
   const ToolChain &HostTC;
+  const llvm::opt::DerivedArgList &DerivedArgs;
 };
 
 } // end namespace toolchains
 } // end namespace driver
 } // end namespace clang
 
 #endif // LLVM_CLANG_LIB_DRIVER_TOOLCHAINS_AMDGPUOPENMP_H

clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp

static bool checkSystemForAMDGPU(const ArgList &Args, const AMDGPUToolChain &TC,
 
   return true;
 }
 } // namespace
 
 AMDGPUOpenMPToolChain::AMDGPUOpenMPToolChain(const Driver &D,
                                              const llvm::Triple &Triple,
                                              const ToolChain &HostTC,
-                                             const ArgList &Args)
-    : ROCMToolChain(D, Triple, Args), HostTC(HostTC) {
+                                             const ArgList &Args,
+                                             const DerivedArgList &DerivedArgs)
+    : ROCMToolChain(D, Triple, Args), HostTC(HostTC), DerivedArgs(DerivedArgs) {
   // Lookup binaries into the driver directory, this is used to
   // discover the clang-offload-bundler executable.
   getProgramPaths().push_back(getDriver().Dir);
 }
 
 void AMDGPUOpenMPToolChain::addClangTargetOptions(
     const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
     Action::OffloadKind DeviceOffloadingKind) const {
if (!RocmInstallation.hasDeviceLibrary()) {
     getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 0;
     return {};
   }
 
   StringRef GpuArch = getProcessorFromTargetID(
       getTriple(), Args.getLastArgValue(options::OPT_march_EQ));
 
   SmallVector<BitCodeLibraryInfo, 12> BCLibs;
-  for (auto BCLib : getCommonDeviceLibNames(Args, GpuArch.str(),
+  for (auto BCLib : getCommonDeviceLibNames(DerivedArgs, GpuArch.str(),
                                             /*IsOpenMP=*/true))
     BCLibs.emplace_back(BCLib);
 
   return BCLibs;
 }

clang/lib/Driver/ToolChains/Clang.cpp

if (IsOpenMPDevice) {
     CmdArgs.push_back("-fopenmp-is-device");
     if (OpenMPDeviceInput) {
       CmdArgs.push_back("-fopenmp-host-ir-file-path");
       CmdArgs.push_back(Args.MakeArgString(OpenMPDeviceInput->getFilename()));
     }
   }
 
   if (Triple.isAMDGPU()) {
-    handleAMDGPUCodeObjectVersionOptions(D, Args, CmdArgs);
+    handleAMDGPUCodeObjectVersionOptions(D, C.getArgs(), CmdArgs,
+                                         /*IsCC1As=*/true);
 
     Args.addOptInFlag(CmdArgs, options::OPT_munsafe_fp_atomics,
                       options::OPT_mno_unsafe_fp_atomics);
   }
 
   // For all the host OpenMP offloading compile jobs we need to pass the targets
   // information using -fopenmp-targets= option.
   if (JA.isHostOffloading(Action::OFK_OpenMP)) {
void ClangAs::ConstructJob(Compilation &C, const JobAction &JA,
   Arg *A;
   if (getDebugFissionKind(D, Args, A) == DwarfFissionKind::Split &&
       T.isOSBinFormatELF()) {
     CmdArgs.push_back("-split-dwarf-output");
     CmdArgs.push_back(SplitDebugName(JA, Args, Input, Output));
   }
 
   if (Triple.isAMDGPU())
-    handleAMDGPUCodeObjectVersionOptions(D, Args, CmdArgs, /*IsCC1As=*/true);
+    handleAMDGPUCodeObjectVersionOptions(D, C.getArgs(), CmdArgs,
+                                         /*IsCC1As=*/true);
 
   assert(Input.isFilename() && "Invalid input.");
   CmdArgs.push_back(Input.getFilename());
 
   const char *Exec = getToolChain().getDriver().getClangProgramPath();
   if (D.CC1Main && !D.CCGenDiagnostics) {
     // Invoke cc1as directly in this process.
     C.addCommand(std::make_unique<CC1Command>(JA, *this,
void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
   if (Args.hasArg(options::OPT_v))
     CmdArgs.push_back("--wrapper-verbose");
 
   if (const Arg *A = Args.getLastArg(options::OPT_g_Group)) {
     if (!A->getOption().matches(options::OPT_g0))
       CmdArgs.push_back("--device-debug");
   }
 
   for (const auto &A : Args.getAllArgValues(options::OPT_Xcuda_ptxas))
     CmdArgs.push_back(Args.MakeArgString("--ptxas-args=" + A));
 
   // Forward remarks passes to the LLVM backend in the wrapper.
   if (const Arg *A = Args.getLastArg(options::OPT_Rpass_EQ))
     CmdArgs.push_back(Args.MakeArgString(Twine("--offload-opt=-pass-remarks=") +
                                          A->getValue()));
   if (const Arg *A = Args.getLastArg(options::OPT_Rpass_missed_EQ))
     CmdArgs.push_back(Args.MakeArgString(
         Twine("--offload-opt=-pass-remarks-missed=") + A->getValue()));
   if (const Arg *A = Args.getLastArg(options::OPT_Rpass_analysis_EQ))
     CmdArgs.push_back(Args.MakeArgString(
         Twine("--offload-opt=-pass-remarks-analysis=") + A->getValue()));
   if (Args.getLastArg(options::OPT_save_temps_EQ))

openmp/libomptarget/DeviceRTL/include/Interface.h

 
 /// Misc
 ///
 ///{
 double omp_get_wtick(void);
 
 double omp_get_wtime(void);
 ///}
+
+#ifdef __AMDGCN__
+size_t external_get_local_size(uint32_t dim);
+size_t external_get_num_groups(uint32_t dim);
+#endif
 }
 
 extern "C" {
 /// Allocate \p Bytes in "shareable" memory and return the address. Needs to be
 /// called balanced with __kmpc_free_shared like a stack (push/pop). Can be
 /// called by any thread, allocation happens *per thread*.
 void *__kmpc_alloc_shared(uint64_t Bytes);
 

openmp/libomptarget/DeviceRTL/src/Mapping.cpp

 
 #include "Mapping.h"
 #include "Interface.h"
 #include "State.h"
 #include "Types.h"
 #include "Utils.h"
 
 #pragma omp begin declare target device_type(nohost)
-
+extern const uint16_t __oclc_ABI_version;
 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
 using namespace _OMP;
 
 namespace _OMP {
 namespace impl {
 
 // Forward declarations defined to be defined for AMDGCN and NVPTX.

 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 const llvm::omp::GV &getGridValue() {
   return llvm::omp::getAMDGPUGridValues<__AMDGCN_WAVEFRONT_SIZE>();
 }
 
-uint32_t getNumHardwareThreadsInBlock() {
-  return __builtin_amdgcn_workgroup_size_x();
-}
+uint32_t getNumHardwareThreadsInBlock() { return external_get_local_size(0); }
 
 LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
 
 LaneMaskTy lanemaskLT() {
   uint32_t Lane = mapping::getThreadIdInWarp();
   int64_t Ballot = mapping::activemask();
   uint64_t Mask = ((uint64_t)1 << Lane) - (uint64_t)1;
   return Mask & Ballot;

 }
 
 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 
 uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); }
 
 uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); }
 
-uint32_t getNumberOfBlocks() {
-  return __builtin_amdgcn_grid_size_x() / __builtin_amdgcn_workgroup_size_x();
-}
+uint32_t getNumberOfBlocks() { return external_get_num_groups(0); }
 
 uint32_t getWarpId() {
   return impl::getThreadIdInBlock() / mapping::getWarpSize();
 }
 
 uint32_t getNumberOfWarpsInBlock() {
   return mapping::getBlockSize() / mapping::getWarpSize();
 }

openmp/libomptarget/DeviceRTL/src/State.cpp

 
 ///}
 
 /// AMDGCN implementations of the shuffle sync idiom.
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
+extern "C" size_t __ockl_get_local_size(uint32_t dim);
+extern "C" size_t __ockl_get_num_groups(uint32_t dim);
+
 extern "C" {
 void *malloc(uint64_t Size) {
   // TODO: Use some preallocated space for dynamic malloc.
   return nullptr;
 }
 
 void free(void *Ptr) {}
 }
 
 #pragma omp end declare variant
 ///}
 
+extern "C" {
+#ifdef __AMDGCN__
+size_t external_get_local_size(uint32_t dim) {
+  return __ockl_get_local_size(dim);
+}
+size_t external_get_num_groups(uint32_t dim) {
+  return __ockl_get_num_groups(dim);
+}
+#endif
+} // extern "C"
+
 /// A "smart" stack in shared memory.
 ///
 /// The stack exposes a malloc/free interface but works like a stack internally.
 /// In fact, it is a separate stack *per warp*. That means, each warp must push
 /// and pop symmetrically or this breaks, badly. The implementation will (aim
 /// to) detect non-lock-step warps and fallback to malloc/free. The same will
 /// happen if a warp runs out of memory. The master warp in generic memory is
 /// special and is given more memory than the rest.

openmp/libomptarget/plugins/amdgpu/impl/get_elf_mach_gfx_name.h

 //===--- amdgpu/impl/get_elf_mach_gfx_name.h ---------------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 #ifndef GET_ELF_MACH_GFX_NAME_H_INCLUDED
 #define GET_ELF_MACH_GFX_NAME_H_INCLUDED
 
 #include <stdint.h>
 
 const char *get_elf_mach_gfx_name(uint32_t EFlags);
 
+enum IMPLICITARGS : uint16_t {
+  COV4_SIZE = 56,
+  COV4_HOSTCALL_PTR_OFFSET = 24,
+  HOSTCALL_PTR_SIZE = 8,
+
+  COV5_SIZE = 256,
+
+  COV5_BLOCK_COUNT_X_OFFSET = 0,
+  COV5_BLOCK_COUNT_X_SIZE = 4,
+
+  COV5_BLOCK_COUNT_Y_OFFSET = 4,
+  COV5_BLOCK_COUNT_Y_SIZE = 4,
+
+  COV5_BLOCK_COUNT_Z_OFFSET = 8,
+  COV5_BLOCK_COUNT_Z_SIZE = 4,
+
+  COV5_GROUP_SIZE_X_OFFSET = 12,
+  COV5_GROUP_SIZE_X_SIZE = 2,
+
+  COV5_GROUP_SIZE_Y_OFFSET = 14,
+  COV5_GROUP_SIZE_Y_SIZE = 2,
+
+  COV5_GROUP_SIZE_Z_OFFSET = 16,
+  COV5_GROUP_SIZE_Z_SIZE = 2,
+
+  COV5_REMAINDER_X_OFFSET = 18,
+  COV5_REMAINDER_X_SIZE = 2,
+
+  COV5_REMAINDER_Y_OFFSET = 20,
+  COV5_REMAINDER_Y_SIZE = 2,
+
+  COV5_REMAINDER_Z_OFFSET = 22,
+  COV5_REMAINDER_Z_SIZE = 2,
+
+  COV5_GRID_DIMS_OFFSET = 64,
+  COV5_GRID_DIMS_SIZE = 2,
+
+  COV5_HOSTCALL_PTR_OFFSET = 80,
+
+  COV5_HEAPV1_PTR_OFFSET = 96,
+  COV5_HEAPV1_PTR_SIZE = 8
+};
+
+const uint16_t implicitArgsSize(uint16_t Version);
+
 #endif

openmp/libomptarget/plugins/amdgpu/impl/get_elf_mach_gfx_name.cpp

 //===--- amdgpu/impl/get_elf_mach_gfx_name.cpp -------------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 #include "get_elf_mach_gfx_name.h"
 
 // This header conflicts with the system elf.h (macros vs enums of the same
 // identifier) and contains more up to date values for the enum checked here.
 // rtl.cpp uses the system elf.h.
 #include "llvm/BinaryFormat/ELF.h"
+using namespace llvm::ELF;
 
 const char *get_elf_mach_gfx_name(uint32_t EFlags) {
   using namespace llvm::ELF;
   uint32_t Gfx = (EFlags & EF_AMDGPU_MACH);
   switch (Gfx) {
   case EF_AMDGPU_MACH_AMDGCN_GFX801:
     return "gfx801";
   case EF_AMDGPU_MACH_AMDGCN_GFX802:
const char *get_elf_mach_gfx_name(uint32_t EFlags) {
   case EF_AMDGPU_MACH_AMDGCN_GFX1102:
     return "gfx1102";
   case EF_AMDGPU_MACH_AMDGCN_GFX1103:
     return "gfx1103";
   default:
     return "--unknown gfx";
   }
 }
+
+const uint16_t implicitArgsSize(uint16_t Version) {
+  return Version < ELFABIVERSION_AMDGPU_HSA_V5 ? IMPLICITARGS::COV4_SIZE
+                                               : IMPLICITARGS::COV5_SIZE;
+}

openmp/libomptarget/plugins/amdgpu/impl/internal.h

 #ifndef TARGET_NAME
 #error "Missing TARGET_NAME macro"
 #endif
 #define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
 #include "Debug.h"
 
 #define MAX_NUM_KERNELS (1024 * 16)
 
-typedef struct impl_implicit_args_s {
-  uint64_t offset_x;
-  uint64_t offset_y;
-  uint64_t offset_z;
-  uint64_t hostcall_ptr;
-  uint64_t unused0;
-  uint64_t unused1;
-  uint64_t unused2;
-} impl_implicit_args_t;
-static_assert(sizeof(impl_implicit_args_t) == 56, "");
-
 // ---------------------- Kernel Start -------------
 typedef struct atl_kernel_info_s {
   uint64_t kernel_object;
   uint32_t group_segment_size;
   uint32_t private_segment_size;
   uint32_t sgpr_count;
   uint32_t vgpr_count;
   uint32_t sgpr_spill_count;

openmp/libomptarget/plugins/amdgpu/impl/system.cpp

enum class ValueKind {
     HiddenGlobalOffsetZ,
     HiddenNone,
     HiddenPrintfBuffer,
     HiddenDefaultQueue,
     HiddenCompletionAction,
     HiddenMultiGridSyncArg,
     HiddenHostcallBuffer,
     HiddenHeapV1,
+    HiddenBlockCountX,
+    HiddenBlockCountY,
+    HiddenBlockCountZ,
+    HiddenGroupSizeX,
+    HiddenGroupSizeY,
+    HiddenGroupSizeZ,
+    HiddenRemainderX,
+    HiddenRemainderY,
+    HiddenRemainderZ,
+    HiddenGridDims,
+    HiddenQueuePtr,
     Unknown
   };
 
   KernelArgMD()
       : name_(std::string()), size_(0), offset_(0),
         valueKind_(ValueKind::Unknown) {}
 
   // fields
static const std::map<std::string, KernelArgMD::ValueKind> ArgValueKind = {
     {"hidden_none", KernelArgMD::ValueKind::HiddenNone},
     {"hidden_printf_buffer", KernelArgMD::ValueKind::HiddenPrintfBuffer},
     {"hidden_default_queue", KernelArgMD::ValueKind::HiddenDefaultQueue},
     {"hidden_completion_action",
      KernelArgMD::ValueKind::HiddenCompletionAction},
     {"hidden_multigrid_sync_arg",
      KernelArgMD::ValueKind::HiddenMultiGridSyncArg},
     {"hidden_hostcall_buffer", KernelArgMD::ValueKind::HiddenHostcallBuffer},
-    {"hidden_heap_v1", KernelArgMD::ValueKind::HiddenHeapV1}};
+    {"hidden_heap_v1", KernelArgMD::ValueKind::HiddenHeapV1},
+    {"hidden_block_count_x", KernelArgMD::ValueKind::HiddenBlockCountX},
+    {"hidden_block_count_y", KernelArgMD::ValueKind::HiddenBlockCountY},
+    {"hidden_block_count_z", KernelArgMD::ValueKind::HiddenBlockCountZ},
+    {"hidden_group_size_x", KernelArgMD::ValueKind::HiddenGroupSizeX},
+    {"hidden_group_size_y", KernelArgMD::ValueKind::HiddenGroupSizeY},
+    {"hidden_group_size_z", KernelArgMD::ValueKind::HiddenGroupSizeZ},
+    {"hidden_remainder_x", KernelArgMD::ValueKind::HiddenRemainderX},
+    {"hidden_remainder_y", KernelArgMD::ValueKind::HiddenRemainderY},
+    {"hidden_remainder_z", KernelArgMD::ValueKind::HiddenRemainderZ},
+    {"hidden_grid_dims", KernelArgMD::ValueKind::HiddenGridDims},
+    {"hidden_queue_ptr", KernelArgMD::ValueKind::HiddenQueuePtr},
+};
 
 namespace core {
 
 hsa_status_t callbackEvent(const hsa_amd_event_t *event, void *data) {
   if (event->event_type == HSA_AMD_GPU_MEMORY_FAULT_EVENT) {
     hsa_amd_gpu_memory_fault_info_t memory_fault = event->memory_fault;
     // memory_fault.agent
     // memory_fault.virtual_address
static bool isImplicit(KernelArgMD::ValueKind value_kind) {
   case KernelArgMD::ValueKind::HiddenGlobalOffsetZ:
   case KernelArgMD::ValueKind::HiddenNone:
   case KernelArgMD::ValueKind::HiddenPrintfBuffer:
   case KernelArgMD::ValueKind::HiddenDefaultQueue:
   case KernelArgMD::ValueKind::HiddenCompletionAction:
   case KernelArgMD::ValueKind::HiddenMultiGridSyncArg:
   case KernelArgMD::ValueKind::HiddenHostcallBuffer:
   case KernelArgMD::ValueKind::HiddenHeapV1:
+  case KernelArgMD::ValueKind::HiddenBlockCountX:
+  case KernelArgMD::ValueKind::HiddenBlockCountY:
+  case KernelArgMD::ValueKind::HiddenBlockCountZ:
+  case KernelArgMD::ValueKind::HiddenGroupSizeX:
+  case KernelArgMD::ValueKind::HiddenGroupSizeY:
+  case KernelArgMD::ValueKind::HiddenGroupSizeZ:
+  case KernelArgMD::ValueKind::HiddenRemainderX:
+  case KernelArgMD::ValueKind::HiddenRemainderY:
+  case KernelArgMD::ValueKind::HiddenRemainderZ:
+  case KernelArgMD::ValueKind::HiddenGridDims:
+  case KernelArgMD::ValueKind::HiddenQueuePtr:
     return true;
   default:
     return false;
   }
 }
 
 static std::pair<const unsigned char *, const unsigned char *>
 findMetadata(const ELFObjectFile<ELF64LE> &ELFObj) {
if (kernel_segment_size > 0) {
           printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
                  "iterate args map in kernel args metadata");
           return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
         }
         // v3 has offset field and not align field
         size_t new_offset = lcArg.offset_;
         size_t padding = new_offset - offset;
         offset = new_offset;
-        DP("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(), lcArg.size_,
-           lcArg.offset_);
         offset += lcArg.size_;
 
         // check if the arg is a hidden/implicit arg
         // this logic assumes that all hidden args are 8-byte aligned
         if (!isImplicit(lcArg.valueKind_)) {
           info.explicit_argument_count++;
           kernel_explicit_args_size += lcArg.size_;
+          DP("Explicit Kernel Arg[%lu] \"%s\" (%u, %u)\n", i,
+             lcArg.name_.c_str(), lcArg.size_, lcArg.offset_);
         } else {
           info.implicit_argument_count++;
           hasHiddenArgs = true;
+          DP("Implicit Kernel Arg[%lu] \"%s\" (%u, %u)\n", i,
+             lcArg.name_.c_str(), lcArg.size_, lcArg.offset_);
         }
         kernel_explicit_args_size += padding;
       }
     }
 
     // TODO: Probably don't want this arithmetic
     info.kernel_segment_size =
-        (hasHiddenArgs ? kernel_explicit_args_size : kernel_segment_size);
+        (!hasHiddenArgs ? kernel_explicit_args_size : kernel_segment_size);
     DP("[%s: kernarg seg size] (%lu --> %u)\n", kernelName.c_str(),
        kernel_segment_size, info.kernel_segment_size);
 
     // kernel received, now add it to the kernel info table
     KernelInfoTable[kernelName] = info;
   }
 
   return HSA_STATUS_SUCCESS;

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

 struct KernelArgPool {
 private:
   static pthread_mutex_t Mutex;
 
 public:
   uint32_t KernargSegmentSize;
   void *KernargRegion = nullptr;
   std::queue<int> FreeKernargSegments;
+  uint16_t CodeObjectVersion;
 
   uint32_t kernargSizeIncludingImplicit() {
-    return KernargSegmentSize + sizeof(impl_implicit_args_t);
+    return KernargSegmentSize + implicitArgsSize(CodeObjectVersion);
   }
 
   ~KernelArgPool() {
     if (KernargRegion) {
       auto R = hsa_amd_memory_pool_free(KernargRegion);
       if (R != HSA_STATUS_SUCCESS) {
         DP("hsa_amd_memory_pool_free failed: %s\n", get_error_string(R));
       }
     }
   }
 
   // Can't really copy or move a mutex
   KernelArgPool() = default;
   KernelArgPool(const KernelArgPool &) = delete;
   KernelArgPool(KernelArgPool &&) = delete;
 
-  KernelArgPool(uint32_t KernargSegmentSize, hsa_amd_memory_pool_t &MemoryPool)
-      : KernargSegmentSize(KernargSegmentSize) {
+  KernelArgPool(uint32_t KernargSegmentSize, hsa_amd_memory_pool_t &MemoryPool,
+                uint16_t CodeObjectVersion)
+      : KernargSegmentSize(KernargSegmentSize),
+        CodeObjectVersion(CodeObjectVersion) {
 
     // impl uses one pool per kernel for all gpus, with a fixed upper size
     // preserving that exact scheme here, including the queue<int>
 
     hsa_status_t Err = hsa_amd_memory_pool_allocate(
         MemoryPool, kernargSizeIncludingImplicit() * MAX_NUM_KERNELS, 0,
         &KernargRegion);
 
struct KernelTy {
   int16_t ConstWGSize;
   int32_t DeviceId;
   void *CallStackAddr = nullptr;
   const char *Name;
 
   KernelTy(llvm::omp::OMPTgtExecModeFlags ExecutionMode, int16_t ConstWgSize,
            int32_t DeviceId, void *CallStackAddr, const char *Name,
            uint32_t KernargSegmentSize,
-           hsa_amd_memory_pool_t &KernArgMemoryPool)
+           hsa_amd_memory_pool_t &KernArgMemoryPool, uint16_t CodeObjectVersion)
       : ExecutionMode(ExecutionMode), ConstWGSize(ConstWgSize),
         DeviceId(DeviceId), CallStackAddr(CallStackAddr), Name(Name) {
     DP("Construct kernelinfo: ExecMode %d\n", ExecutionMode);
 
     std::string N(Name);
     if (KernelArgPoolMap.find(N) == KernelArgPoolMap.end()) {
-      KernelArgPoolMap.insert(
-          std::make_pair(N, std::unique_ptr<KernelArgPool>(new KernelArgPool(
-                                KernargSegmentSize, KernArgMemoryPool))));
+      KernelArgPoolMap.insert(std::make_pair(
+          N, std::unique_ptr<KernelArgPool>(new KernelArgPool(
+                 KernargSegmentSize, KernArgMemoryPool, CodeObjectVersion))));
     }
   }
 };
 
 /// List that contains all the kernels.
 /// FIXME: we may need this to be per device and per library.
 std::list<KernelTy> KernelsList;
 
public:
 
   // Device properties
   std::vector<int> ComputeUnits;
   std::vector<int> GroupsPerDevice;
   std::vector<int> ThreadsPerGroup;
   std::vector<int> WarpSize;
   std::vector<std::string> GPUName;
   std::vector<std::string> TargetID;
+  uint16_t CodeObjectVersion;
 
   // OpenMP properties
   std::vector<int> NumTeams;
   std::vector<int> NumThreads;
 
   // OpenMP Environment properties
   EnvironmentVariables Env;
 
   // OpenMP Requires Flags
   int64_t RequiresFlags;
 
   // Resource pools
   SignalPoolT FreeSignalPool;
+  std::vector<void *> PreallocatedDeviceHeap;
 
   bool HostcallRequired = false;
 
   std::vector<hsa_executable_t> HSAExecutables;
 
   std::vector<std::map<std::string, atl_kernel_info_t>> KernelInfoTable;
   std::vector<std::map<std::string, atl_symbol_info_t>> SymbolInfoTable;
 
public:
   }
 
   // Clear entries table for a device
   void clearOffloadEntriesTable(int DeviceId) {
     assert(DeviceId < (int32_t)FuncGblEntries.size() &&
            "Unexpected device id!");
     FuncGblEntries[DeviceId].emplace_back();
     FuncOrGblEntryTy &E = FuncGblEntries[DeviceId].back();
-    // KernelArgPoolMap.clear();
     E.Entries.clear();
     E.Table.EntriesBegin = E.Table.EntriesEnd = 0;
   }
 
   hsa_status_t addDeviceMemoryPool(hsa_amd_memory_pool_t MemoryPool,
                                    unsigned int DeviceId) {
     assert(DeviceId < DeviceFineGrainedMemoryPools.size() && "Error here.");
     uint32_t GlobalFlags = 0;
RTLDeviceInfoTy() {
     WarpSize.resize(NumberOfDevices);
     NumTeams.resize(NumberOfDevices);
     NumThreads.resize(NumberOfDevices);
     DeviceStateStore.resize(NumberOfDevices);
     KernelInfoTable.resize(NumberOfDevices);
     SymbolInfoTable.resize(NumberOfDevices);
     DeviceCoarseGrainedMemoryPools.resize(NumberOfDevices);
     DeviceFineGrainedMemoryPools.resize(NumberOfDevices);
+    PreallocatedDeviceHeap.resize(NumberOfDevices);
 
     Err = setupDevicePools(HSAAgents);
     if (Err != HSA_STATUS_SUCCESS) {
       DP("Setup for Device Memory Pools failed\n");
       return;
     }
 
     Err = setupHostMemoryPools(CPUAgents);
static uint64_t acquireAvailablePacketId(hsa_queue_t *Queue) {
   bool Full = true;
   while (Full) {
     Full =
         PacketId >= (Queue->size + hsa_queue_load_read_index_scacquire(Queue));
   }
   return PacketId;
 }
 
+const uint16_t getCodeObjectVersionFromELF(__tgt_device_image *Image) {
+  char *ImageBegin = (char *)Image->ImageStart;
+  size_t ImageSize = (char *)Image->ImageEnd - ImageBegin;
+
+  StringRef Buffer = StringRef(ImageBegin, ImageSize);
+  auto ElfOrErr = ObjectFile::createELFObjectFile(MemoryBufferRef(Buffer, ""),
+                                                  /*InitContent=*/false);
+  if (!ElfOrErr) {
+    REPORT("Failed to load ELF: %s\n", toString(ElfOrErr.takeError()).c_str());
+    return 1;
+  }
+
+  if (const auto *ELFObj = dyn_cast<ELF64LEObjectFile>(ElfOrErr->get())) {
+    auto Header = ELFObj->getELFFile().getHeader();
+    uint16_t Version = (uint8_t)(Header.e_ident[EI_ABIVERSION]);
+    DP("ELFABIVERSION Version: %u\n", Version);
+    return Version;
+  }
+  return 0;
+}
+
 int32_t runRegionLocked(int32_t DeviceId, void *TgtEntryPtr, void **TgtArgs,
                         ptrdiff_t *TgtOffsets, int32_t ArgNum, int32_t NumTeams,
                         int32_t ThreadLimit, uint64_t LoopTripcount) {
   // Set the context we are using
   // update thread limit content in gpu memory if un-initialized or specified
   // from host
 
   DP("Run target team region thread_limit %d\n", ThreadLimit);
int32_t runRegionLocked(int32_t DeviceId, void *TgtEntryPtr, void **TgtArgs,
   const uint32_t GroupSegmentSize =
       KernelInfoEntry.group_segment_size + DeviceInfo().Env.DynamicMemSize;
   const uint32_t SgprCount = KernelInfoEntry.sgpr_count;
   const uint32_t VgprCount = KernelInfoEntry.vgpr_count;
   const uint32_t SgprSpillCount = KernelInfoEntry.sgpr_spill_count;
   const uint32_t VgprSpillCount = KernelInfoEntry.vgpr_spill_count;
 
   assert(ArgNum == (int)KernelInfoEntry.explicit_argument_count);
+  int NumGroups = 0;
+  uint16_t ThreadsPerGroup = 0;
 
   /*
    * Set limit based on ThreadsPerGroup and GroupsPerDevice
    */
   LaunchVals LV =
       getLaunchVals(DeviceInfo().WarpSize[DeviceId], DeviceInfo().Env,
                     KernelInfo->ConstWGSize, KernelInfo->ExecutionMode,
                     NumTeams,      // From run_region arg
                     ThreadLimit,   // From run_region arg
                     LoopTripcount, // From run_region arg
                     DeviceInfo().NumTeams[KernelInfo->DeviceId]);
   const int GridSize = LV.GridSize;
   const int WorkgroupSize = LV.WorkgroupSize;
 
   if (print_kernel_trace >= LAUNCH) {
-    int NumGroups = GridSize / WorkgroupSize;
+    NumGroups = GridSize / WorkgroupSize;
     // enum modes are SPMD, GENERIC, NONE 0,1,2
     // if doing rtl timing, print to stderr, unless stdout requested.
     bool TraceToStdout = print_kernel_trace & (RTL_TO_STDOUT | RTL_TIMING);
     fprintf(TraceToStdout ? stdout : stderr,
             "DEVID:%2d SGN:%1d ConstWGSize:%-4d args:%2d teamsXthrds:(%4dX%4d) "
             "reqd:(%4dX%4d) lds_usage:%uB sgpr_count:%u vgpr_count:%u "
             "sgpr_spill_count:%u vgpr_spill_count:%u tripcount:%lu n:%s\n",
             DeviceId, KernelInfo->ExecutionMode, KernelInfo->ConstWGSize,
             ArgNum, NumGroups, WorkgroupSize, NumTeams, ThreadLimit,
             GroupSegmentSize, SgprCount, VgprCount, SgprSpillCount,
             VgprSpillCount, LoopTripcount, KernelInfo->Name);
   }
 
   // Run on the device.
   {
     hsa_queue_t *Queue = DeviceInfo().HSAQueueSchedulers[DeviceId].next();
     if (!Queue) {
       return OFFLOAD_FAIL;
     }
     uint64_t PacketId = acquireAvailablePacketId(Queue);
 
+    uint16_t CodeObjectVersion = DeviceInfo().CodeObjectVersion;
     const uint32_t Mask = Queue->size - 1; // size is a power of 2
     hsa_kernel_dispatch_packet_t *Packet =
         (hsa_kernel_dispatch_packet_t *)Queue->base_address + (PacketId & Mask);
 
     // packet->header is written last
     Packet->setup = UINT16_C(1) << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
     Packet->workgroup_size_x = WorkgroupSize;
     Packet->workgroup_size_y = 1;
if (!ArgPool) {
         return OFFLOAD_FAIL;
       }
 
       // Copy explicit arguments
       for (int I = 0; I < ArgNum; I++) {
         memcpy((char *)KernArg + sizeof(void *) * I, Args[I], sizeof(void *));
       }
 
-      // Initialize implicit arguments. TODO: Which of these can be dropped
-      impl_implicit_args_t *ImplArgs = reinterpret_cast<impl_implicit_args_t *>(
-          static_cast<char *>(KernArg) + ArgPool->KernargSegmentSize);
-      memset(ImplArgs, 0,
-             sizeof(impl_implicit_args_t)); // may not be necessary
-      ImplArgs->offset_x = 0;
-      ImplArgs->offset_y = 0;
-      ImplArgs->offset_z = 0;
+      uint8_t *ImplArgs =
+          static_cast<uint8_t *>(KernArg) + sizeof(void *) * ArgNum;
+      memset(ImplArgs, 0, implicitArgsSize(CodeObjectVersion));
+
+      uint64_t Buffer = 0;
 
       // assign a hostcall buffer for the selected Q
       if (__atomic_load_n(&DeviceInfo().HostcallRequired, __ATOMIC_ACQUIRE)) {
         // hostrpc_assign_buffer is not thread safe, and this function is
         // under a multiple reader lock, not a writer lock.
         static pthread_mutex_t HostcallInitLock = PTHREAD_MUTEX_INITIALIZER;
         pthread_mutex_lock(&HostcallInitLock);
-        uint64_t Buffer = hostrpc_assign_buffer(
-            DeviceInfo().HSAAgents[DeviceId], Queue, DeviceId);
+        Buffer = hostrpc_assign_buffer(DeviceInfo().HSAAgents[DeviceId], Queue,
+                                       DeviceId);
         pthread_mutex_unlock(&HostcallInitLock);
         if (!Buffer) {
           DP("hostrpc_assign_buffer failed, gpu would dereference null and "
              "error\n");
           return OFFLOAD_FAIL;
         }
+      }
 
-        DP("Implicit argument count: %d\n",
-           KernelInfoEntry.implicit_argument_count);
-        if (KernelInfoEntry.implicit_argument_count >= 4) {
-          // Initialise pointer for implicit_argument_count != 0 ABI
-          // Guess that the right implicit argument is at offset 24 after
-          // the explicit arguments. In the future, should be able to read
-          // the offset from msgpack. Clang is not annotating it at present.
-          uint64_t Offset =
-              sizeof(void *) * (KernelInfoEntry.explicit_argument_count + 3);
-          if ((Offset + 8) > ArgPool->kernargSizeIncludingImplicit()) {
-            DP("Bad offset of hostcall: %lu, exceeds kernarg size w/ implicit "
-               "args: %d\n",
-               Offset + 8, ArgPool->kernargSizeIncludingImplicit());
-          } else {
-            memcpy(static_cast<char *>(KernArg) + Offset, &Buffer, 8);
-          }
-        }
 
-        // initialise pointer for implicit_argument_count == 0 ABI
-        ImplArgs->hostcall_ptr = Buffer;
+      if (CodeObjectVersion < llvm::ELF::ELFABIVERSION_AMDGPU_HSA_V5) {
+        DP("Setting Hostcall buffer for COV4\n");
+        memcpy(&ImplArgs[IMPLICITARGS::COV4_HOSTCALL_PTR_OFFSET], &Buffer,
+               IMPLICITARGS::HOSTCALL_PTR_SIZE);
+      } else {
+        DP("Setting fields of ImplicitArgs for COV5\n");
+        uint16_t Remainder = 0;
+        uint16_t GridDims = 1;
+        uint32_t NumGroupsYZ = 1;
+        uint16_t ThreadsPerGroupYZ = 0;
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_BLOCK_COUNT_X_OFFSET], &NumGroups,
+               IMPLICITARGS::COV5_BLOCK_COUNT_X_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_BLOCK_COUNT_Y_OFFSET], &NumGroupsYZ,
+               IMPLICITARGS::COV5_BLOCK_COUNT_Y_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_BLOCK_COUNT_Z_OFFSET], &NumGroupsYZ,
+               IMPLICITARGS::COV5_BLOCK_COUNT_Z_SIZE);
+
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_GROUP_SIZE_X_OFFSET],
+               &ThreadsPerGroup, IMPLICITARGS::COV5_GROUP_SIZE_X_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_GROUP_SIZE_Y_OFFSET],
+               &ThreadsPerGroupYZ, IMPLICITARGS::COV5_GROUP_SIZE_Y_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_GROUP_SIZE_Z_OFFSET],
+               &ThreadsPerGroupYZ, IMPLICITARGS::COV5_GROUP_SIZE_Z_SIZE);
+
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_REMAINDER_X_OFFSET], &Remainder,
+               IMPLICITARGS::COV5_REMAINDER_X_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_REMAINDER_Y_OFFSET], &Remainder,
+               IMPLICITARGS::COV5_REMAINDER_Y_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_REMAINDER_Z_OFFSET], &Remainder,
+               IMPLICITARGS::COV5_REMAINDER_Z_SIZE);
+
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_GRID_DIMS_OFFSET], &GridDims,
+               IMPLICITARGS::COV5_GRID_DIMS_SIZE);
+
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_HOSTCALL_PTR_OFFSET], &Buffer,
+               IMPLICITARGS::HOSTCALL_PTR_SIZE);
+        memcpy(&ImplArgs[IMPLICITARGS::COV5_HEAPV1_PTR_OFFSET],
+               &(DeviceInfo().PreallocatedDeviceHeap[DeviceId]),
+               IMPLICITARGS::COV5_HEAPV1_PTR_SIZE);
       }
 
       Packet->kernarg_address = KernArg;
     }
 
     hsa_signal_t S = DeviceInfo().FreeSignalPool.pop();
     if (S.handle == 0) {
       DP("Failed to get signal instance\n");
DP("Device %d: wavefront size %d, total threads %d x %d = %d\n", DeviceId,
      DeviceInfo().WarpSize[DeviceId], DeviceInfo().ThreadsPerGroup[DeviceId],
      DeviceInfo().GroupsPerDevice[DeviceId],
      DeviceInfo().GroupsPerDevice[DeviceId] *
          DeviceInfo().ThreadsPerGroup[DeviceId]);
 
   return OFFLOAD_SUCCESS;
 }
 
+static void preAllocateHeapMemoryForCov5() {
+  void *DevPtr;
+  for (int I = 0; I < DeviceInfo().NumberOfDevices; I++) {
+    DevPtr = nullptr;
+    size_t PreAllocSize = 131072; // 128KB per device
+
+    hsa_amd_memory_pool_t MemoryPool =
+        DeviceInfo().DeviceCoarseGrainedMemoryPools[I];
+    hsa_status_t Err =
+        hsa_amd_memory_pool_allocate(MemoryPool, PreAllocSize, 0, &DevPtr);
+    if (Err != HSA_STATUS_SUCCESS) {
+      DP("Error allocating preallocated heap device memory: %s\n",
+         get_error_string(Err));
+    }
+
+    Err = hsa_amd_agents_allow_access(1, &DeviceInfo().HSAAgents[I], NULL,
+                                      DevPtr);
+    if (Err != HSA_STATUS_SUCCESS) {
+      DP("hsa allow_access_to_all_gpu_agents failed: %s\n",
+         get_error_string(Err));
+    }
+
+    uint64_t Rounded =
+        sizeof(uint32_t) * ((PreAllocSize + 3) / sizeof(uint32_t));
+    Err = hsa_amd_memory_fill(DevPtr, 0, Rounded / sizeof(uint32_t));
+    if (Err != HSA_STATUS_SUCCESS) {
+      DP("Error zero-initializing preallocated heap device memory:%s\n",
+         get_error_string(Err));
+    }
+
+    DeviceInfo().PreallocatedDeviceHeap[I] = DevPtr;
+  }
+}
+
 static __tgt_target_table *
 __tgt_rtl_load_binary_locked(int32_t DeviceId, __tgt_device_image *Image);
 
 __tgt_target_table *__tgt_rtl_load_binary(int32_t DeviceId,
                                           __tgt_device_image *Image) {
   DeviceInfo().LoadRunLock.lock();
   __tgt_target_table *Res = __tgt_rtl_load_binary_locked(DeviceId, Image);
   DeviceInfo().LoadRunLock.unlock();
__tgt_target_table *__tgt_rtl_load_binary_locked(int32_t DeviceId,
   DeviceInfo().clearOffloadEntriesTable(DeviceId);
 
   // We do not need to set the ELF version because the caller of this function
   // had to do that to decide the right runtime to use
 
   if (!elfMachineIdIsAmdgcn(Image))
     return NULL;
 
+  DeviceInfo().CodeObjectVersion = getCodeObjectVersionFromELF(Image);
+  if (DeviceInfo().CodeObjectVersion >=
+      llvm::ELF::ELFABIVERSION_AMDGPU_HSA_V5) {
+    preAllocateHeapMemoryForCov5();
+  }
+
   {
     auto Env =
         DeviceEnvironment(DeviceId, DeviceInfo().NumberOfDevices,
                           DeviceInfo().Env.DynamicMemSize, Image, ImgSize);
 
     auto &KernelInfo = DeviceInfo().KernelInfoTable[DeviceId];
     auto &SymbolInfo = DeviceInfo().SymbolInfoTable[DeviceId];
     hsa_status_t Err = moduleRegisterFromMemoryToPlace(
if (Err == HSA_STATUS_SUCCESS) {
          "value "
          "GENERIC (1)\n",
          ExecModeName);
     }
     check("Loading computation property", Err);
 
     KernelsList.push_back(KernelTy(ExecModeVal, WGSizeVal, DeviceId,
                                    CallStackAddr, E->name, KernargSegmentSize,
-                                   DeviceInfo().KernArgPool));
+                                   DeviceInfo().KernArgPool,
+                                   DeviceInfo().CodeObjectVersion));
     __tgt_offload_entry Entry = *E;
     Entry.addr = (void *)&KernelsList.back();
     DeviceInfo().addOffloadEntry(DeviceId, Entry);
     DP("Entry point %ld maps to %s\n", E - HostBegin, E->name);
   }
 
   return DeviceInfo().getOffloadEntriesTable(DeviceId);
 }