SYCL :: OptionalKernelFeatures/is_compatible.cpp LIT test fails on HIP

[AlexeySachkov]

Describe the bug

The test case crashes within llvm::AMDGPUAsmPrinter::emitEndOfAsmFile with assertion Success && "Malformed HSA Metadata".
See https://github.com/intel/llvm/actions/runs/4038705536/jobs/6943046676 for full logs, for example.

Also copying full log here in case it expires from GH Actions:

FAIL: SYCL :: OptionalKernelFeatures/is_compatible.cpp (956 of 1183)
******************** TEST 'SYCL :: OptionalKernelFeatures/is_compatible.cpp' FAILED ********************
Script:
--
: 'RUN: at line 2';    /__w/llvm/llvm/toolchain/bin/clang++     -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa -O0 /__w/llvm/llvm/llvm_test_suite/SYCL/OptionalKernelFeatures/is_compatible.cpp -o /__w/llvm/llvm/build/SYCL/OptionalKernelFeatures/Output/is_compatible.cpp.tmp.out
: 'RUN: at line 3';   true /__w/llvm/llvm/build/SYCL/OptionalKernelFeatures/Output/is_compatible.cpp.tmp.out
: 'RUN: at line 4';    env ONEAPI_DEVICE_SELECTOR=ext_oneapi_hip:gpu  /__w/llvm/llvm/build/SYCL/OptionalKernelFeatures/Output/is_compatible.cpp.tmp.out
: 'RUN: at line 5';   true /__w/llvm/llvm/build/SYCL/OptionalKernelFeatures/Output/is_compatible.cpp.tmp.out
--
Exit Code: 1

Command Output (stdout):
--
$ ":" "RUN: at line 2"
note: command had no output on stdout or stderr
$ "/__w/llvm/llvm/toolchain/bin/clang++" "-Xsycl-target-backend=amdgcn-amd-amdhsa" "--offload-arch=gfx1031" "-fsycl" "-fsycl-targets=amdgcn-amd-amdhsa" "-O0" "/__w/llvm/llvm/llvm_test_suite/SYCL/OptionalKernelFeatures/is_compatible.cpp" "-o" "/__w/llvm/llvm/build/SYCL/OptionalKernelFeatures/Output/is_compatible.cpp.tmp.out"
# command stderr:
lld: /__w/llvm/llvm/src/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp:173: virtual void llvm::AMDGPUAsmPrinter::emitEndOfAsmFile(llvm::Module&): Assertion `Success && "Malformed HSA Metadata"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: /__w/llvm/llvm/toolchain/bin/lld -flavor gnu -m elf64_amdgpu --no-undefined -shared -plugin-opt=-amdgpu-internalize-symbols -plugin-opt=mcpu=gfx1031 -plugin-opt=O0 -plugin-opt=no-opaque-pointers -o /tmp/lit-tmp-z95x2lnm/is_compatible-gfx1031-49a176.out /tmp/lit-tmp-z95x2lnm/is_compatible-gfx1031-d0242b.o
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it):
0  lld             0x000055bb431c3a1f llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 63
1  lld             0x000055bb431c13e4
2  libpthread.so.0 0x00007f473a250420
3  libc.so.6       0x00007f4739ced00b gsignal + 203
4  libc.so.6       0x00007f4739ccc[859](https://github.com/intel/llvm/actions/runs/4038705536/jobs/6943046676#step:8:882) abort + 299
5  libc.so.6       0x00007f4739ccc729
6  libc.so.6       0x00007f4739cddfd6
7  lld             0x000055bb43aa0cb0
8  lld             0x000055bb43e8de58 llvm::AsmPrinter::doFinalization(llvm::Module&) + 6440
9  lld             0x000055bb46132305 llvm::FPPassManager::doFinalization(llvm::Module&) + 117
10 lld             0x000055bb4613eac6 llvm::legacy::PassManagerImpl::run(llvm::Module&) + 1206
11 lld             0x000055bb444bea6d
12 lld             0x000055bb444bf01d llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) + 701
13 lld             0x000055bb444b1cb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) + 2901
14 lld             0x000055bb444b2305 llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) + 757
15 lld             0x000055bb43388840 lld::elf::BitcodeCompiler::compile() + 384
16 lld             0x000055bb432f529f void lld::elf::LinkerDriver::compileBitcodeFiles<llvm::object::ELFType<(llvm::support::endianness)1, true>>(bool) + 207
17 lld             0x000055bb4330b934 lld::elf::LinkerDriver::link(llvm::opt::InputArgList&) + 7492
18 lld             0x000055bb4330dd2e lld::elf::LinkerDriver::linkerMain(llvm::ArrayRef<char const*>) + 4302
19 lld             0x000055bb4330f41a lld::elf::link(llvm::ArrayRef<char const*>, llvm::raw_ostream&, llvm::raw_ostream&, bool, bool) + 4426
20 lld             0x000055bb43135704
21 lld             0x000055bb43135b56 lld_main(int, char**) + 422
22 libc.so.6       0x00007f4739cce083 __libc_start_main + 243
23 lld             0x000055bb4313308e _start + 46
llvm-foreach: Aborted (core dumped)
lld: /__w/llvm/llvm/src/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp:173: virtual void llvm::AMDGPUAsmPrinter::emitEndOfAsmFile(llvm::Module&): Assertion `Success && "Malformed HSA Metadata"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: /__w/llvm/llvm/toolchain/bin/lld -flavor gnu -m elf64_amdgpu --no-undefined -shared -plugin-opt=-amdgpu-internalize-symbols -plugin-opt=mcpu=gfx1031 -plugin-opt=O0 -plugin-opt=no-opaque-pointers -o /tmp/lit-tmp-z95x2lnm/is_compatible-gfx1031-2d6778.out /tmp/lit-tmp-z95x2lnm/is_compatible-gfx1031-6b24b3.o
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it):
0  lld             0x0000555baf344a1f llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) + 63
1  lld             0x0000555baf3423e4
2  libpthread.so.0 0x00007fb798036420
3  libc.so.6       0x00007fb797ad300b gsignal + 203
4  libc.so.6       0x00007fb797ab2859 abort + 299
5  libc.so.6       0x00007fb797ab2729
6  libc.so.6       0x00007fb797ac3fd6
7  lld             0x0000555bafc21cb0
8  lld             0x0000555bb000ee58 llvm::AsmPrinter::doFinalization(llvm::Module&) + 6440
9  lld             0x0000555bb22b3305 llvm::FPPassManager::doFinalization(llvm::Module&) + 117
10 lld             0x0000555bb22bfac6 llvm::legacy::PassManagerImpl::run(llvm::Module&) + 1206
11 lld             0x0000555bb063fa6d
12 lld             0x0000555bb064001d llvm::lto::backend(llvm::lto::Config const&, std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) + 701
13 lld             0x0000555bb0632cb5 llvm::lto::LTO::runRegularLTO(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) + 2[901](https://github.com/intel/llvm/actions/runs/4038705536/jobs/6943046676#step:8:924)
14 lld             0x0000555bb0633305 llvm::lto::LTO::run(std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::function<llvm::Expected<std::function<llvm::Expected<std::unique_ptr<llvm::CachedFileStream, std::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) + 757
15 lld             0x0000555baf509840 lld::elf::BitcodeCompiler::compile() + 384
16 lld             0x0000555baf47629f void lld::elf::LinkerDriver::compileBitcodeFiles<llvm::object::ELFType<(llvm::support::endianness)1, true>>(bool) + 207
17 lld             0x0000555baf48c934 lld::elf::LinkerDriver::link(llvm::opt::InputArgList&) + 7492
18 lld             0x0000555baf48ed2e lld::elf::LinkerDriver::linkerMain(llvm::ArrayRef<char const*>) + 4302
19 lld             0x0000555baf4[904](https://github.com/intel/llvm/actions/runs/4038705536/jobs/6943046676#step:8:927)1a lld::elf::link(llvm::ArrayRef<char const*>, llvm::raw_ostream&, llvm::raw_ostream&, bool, bool) + 4426
20 lld             0x0000555baf2b6704
21 lld             0x0000555baf2b6b56 lld_main(int, char**) + 422
22 libc.so.6       0x00007fb797ab4083 __libc_start_main + 243
23 lld             0x0000555baf2b408e _start + 46
llvm-foreach: Aborted (core dumped)
clang-16: error: amdgcn-link command failed with exit code 254 (use -v to see invocation)
clang version 16.0.0 (https://github.com/intel/llvm 39f4c53186a4e375ac956978f14fc29a62945d31)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /__w/llvm/llvm/toolchain/bin
clang-16: note: diagnostic msg: Error generating preprocessed source(s).

error: command failed with exit status: 1

--

[AlexeySachkov]

The test will be disabled on HIP by intel/llvm-test-suite#1556

[zjin-lcf]

@steffenlarsen

Do you know what changes are needed to support sycl::reqd_work_group_size in the HIP backend ? Thanks.

[zjin-lcf]

@AlexeySachkov

I think Compatible = false for the following check. Before the check, Compatible is either true or false. Hence, the check does not change the value of Compatible.

if (Dev.get_infosycl::info::device::max_work_group_size() > INT_MAX) {
Compatible &= true;
}

[steffenlarsen]

Do you know what changes are needed to support sycl::reqd_work_group_size in the HIP backend ? Thanks.

There are a couple of parts to this. For it to be fully supported, we need this issue fixed and we need to make the runtime throw when the specified work-group size doesn't adhere to the sycl::reqd_work_group_size instead of relying on the backends reporting it upon launch.

I think Compatible = false for the following check. Before the check, Compatible is either true or false. Hence, the check does not change the value of Compatible.

if (Dev.get_infosycl::info::device::max_work_group_size() > INT_MAX) { Compatible &= true; }

Though I agree this looks a little off (tag @dm-vodopyanov) it is not the cause of this failure. I suspect the AMDGPU ASM printer is unable to handle reqd_work_group_size metadata that does not have 3 arguments, which we are now able to generate.

@AerialMantis | @jchlanda - Based on llvm/docs/AMDGPUUsage.rst it seems like there is such a relation (MetadataStreamerYamlV2::emitKernelAttrs I suspect). It likely has the same problem with "work_group_size_hint".

[zjin-lcf]

As long as the Clang compiler generates a 3D required work-group sizes when the sycl::reqd_work_group_size is 1D or 2D, then the fact that the AMDGPU ASM printer couldn't handle 3-argument metadata does not matter.

[steffenlarsen]

As long as the Clang compiler generates a 3D required work-group sizes when the sycl::reqd_work_group_size is 1D or 2D, then the fact that the AMDGPU ASM printer couldn't handle 3-argument metadata does not matter.

But it doesn't anymore. It now generates between 1 and 3 arguments, based on how many were passed. This is needed to correctly identify the dimensionality of the attribute as padding would remove that information.

[zjin-lcf]

I am not familiar with LLVM and don't know about padding and loss of dimensionality, but find a section in CodeGen. I assume that the values of YDimVal and ZDimVal are NULL when sycl::reqd_work_group_size is 1D. When they are false, Builder.getInt(1) is pushed to AttrMDArgs. Thanks.

  if (const SYCLReqdWorkGroupSizeAttr *A =
          FD->getAttr<SYCLReqdWorkGroupSizeAttr>()) {
    llvm::Optional<llvm::APSInt> XDimVal = A->getXDimVal();
    llvm::Optional<llvm::APSInt> YDimVal = A->getYDimVal();
    llvm::Optional<llvm::APSInt> ZDimVal = A->getZDimVal();
    llvm::SmallVector<llvm::Metadata *, 3> AttrMDArgs;

    // On SYCL target the dimensions are reversed if present.
    if (ZDimVal)
      AttrMDArgs.push_back(
          llvm::ConstantAsMetadata::get(Builder.getInt(*ZDimVal)));
    if (YDimVal)
      AttrMDArgs.push_back(
          llvm::ConstantAsMetadata::get(Builder.getInt(*YDimVal)));
    AttrMDArgs.push_back(
        llvm::ConstantAsMetadata::get(Builder.getInt(*XDimVal)));

    Fn->setMetadata("reqd_work_group_size",
                    llvm::MDNode::get(Context, AttrMDArgs));

I

[steffenlarsen]

Originally we did pad with 1's, as you mention, but it caused mistakes when we needed to determine if the original was 1D, 2D or 3D. This was changed as of #7450 so now the attribute (in SYCL) only generates metadata with the arguments specified in the attribute.