Create one bitcode library for AMD #15055

MartinWehking · 2024-08-13T15:54:05Z

Enable compilation of libdevice for AMD by adding AMDGCN to macro
guarded code parts in libdevice for enabling e.g. standard library
math function.

Add compilation workflow to SYCLLibdevice.cmake for AMD.

Follow the compilation mechanism for NVPTX
(56a6ae2)
and create a single bitcode library file.

Do not select builtin LLVM intrinsics for AMDGCN by default to ensure
that stdlib functions can be found when linking against libdevice.

Ensure that the clang tests check for the correctness of the new
clang driver actions and check if the driver still links the device
code against the itt device libraries when device library linkage has
been excluded.

Fix a compilation error of Intel math function libraries for MSVC
when targeting AMD. Include "device.h" before including "device_imf.hpp"
to avoid the inclusion of <type_traits>, which failed with a
redefinition of symbols error.

Enable compilation of libdevice for AMD by adding AMDGCN to macro guarded code parts in libdevice for enabling e.g. standard library math function. Add compilation workflow to SYCLLibdevice.cmake for AMD. Follow the compilation mechanism for NVPTX (3668c0a) and create a single bitcode library file. Do not select builtin LLVM intrinsics for AMDGCN by default to ensure that stdlib functions can be found when linking against libdevice. Ensure that the clang tests check for the correctness of the new clang driver actions and check if the driver still links the device code against the itt device libraries when device library linkage has been excluded. Fix a compilation error of Intel math function libraries for MSVC when targeting AMD. Include "device.h" before including "device_imf.hpp" to avoid the inclusion of <type_traits>, which failed with a redefinition of symbols error.

Check if device lib flags get treated correctly for AMDGCN and that the linking actions for the correct device libraries are generated.

steffenlarsen

Looks reasonable to me! ⭐

clang/lib/Driver/ToolChains/SYCL.cpp

libdevice/cmake/modules/SYCLLibdevice.cmake

elizabethandrews

FE changes LGTM

MartinWehking · 2024-09-23T11:33:18Z

ping @intel/dpcpp-clang-driver-reviewers

clang/test/Driver/sycl-device-lib-amdgcn.cpp

mdtoguchi

LGTM - thanks!

AlexeySachkov · 2024-09-27T09:57:14Z

@ldrumm, it seems like this PR introduced some new failures on AMD in post-commit, could you please take a look?

https://github.com/intel/llvm/actions/runs/11031105056/job/30637542157

It also looks like we have some problems with build log extraction, because logs below contain some random symbols as if we read logs from incorrect memory location.

********************
FAIL: SYCL :: Basic/launch_queries/max_num_work_groups.cpp (294 of 2185)
******************** TEST 'SYCL :: Basic/launch_queries/max_num_work_groups.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 1
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp -o /__w/llvm/llvm/build-e2e/Basic/launch_queries/Output/max_num_work_groups.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp -o /__w/llvm/llvm/build-e2e/Basic/launch_queries/Output/max_num_work_groups.cpp.tmp.out
# note: command had no output on stdout or stderr
# RUN: at line 2
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/Basic/launch_queries/Output/max_num_work_groups.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/Basic/launch_queries/Output/max_num_work_groups.cpp.tmp.out
# .---command stderr------------
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           209
# | 	Name:            hipErrorNoBinaryForGpu
# | 	Description:     no kernel image is available for execution on the device
# | 	Function:        buildProgram
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:235
# | 
# | terminate called after throwing an instance of 'sycl::_V1::exception'
# |   what():  The program was built for 1 devices
# | Build program log for 'AMD Radeon RX 6700 XT':
# | P�T�
# `-----------------------------
# error: command failed with exit status: -6

--

********************
FAIL: SYCL :: Basic/kernel_bundle/kernel_bundle_api_hip.cpp (373 of 2185)
******************** TEST 'SYCL :: Basic/kernel_bundle/kernel_bundle_api_hip.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 2
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp -fsycl-device-code-split=per_kernel -o /__w/llvm/llvm/build-e2e/Basic/kernel_bundle/Output/kernel_bundle_api_hip.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp -fsycl-device-code-split=per_kernel -o /__w/llvm/llvm/build-e2e/Basic/kernel_bundle/Output/kernel_bundle_api_hip.cpp.tmp.out
# note: command had no output on stdout or stderr
# RUN: at line 3
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/Basic/kernel_bundle/Output/kernel_bundle_api_hip.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/Basic/kernel_bundle/Output/kernel_bundle_api_hip.cpp.tmp.out
# .---command stderr------------
# | Empty list of devices for get_kernel_bundle
# | Caught: Not all devices are associated with the context or vector of devices is empty
# | Expect: Not all devices are associated with the context or vector of devices is empty
# | Mismatched contexts for join
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           209
# | 	Name:            hipErrorNoBinaryForGpu
# | 	Description:     no kernel image is available for execution on the device
# | 	Function:        buildProgram
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:235
# | 
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           UR_RESULT_ERROR_UNKNOWN
# | 	Function:        urProgramCompile
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:286
# | 
# | Caught: The program was built for 1 devices
# | Build program log for 'AMD Radeon RX 6700 XT':
# | �է�
# | Expect: Not all input bundles have the same associated context
# | kernel_bundle_api_hip.cpp.tmp.out: /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp:25: void checkException(TryBodyT, const std::string &) [TryBodyT = (lambda at /__w/llvm/llvm/llvm/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp:168:9)]: Assertion `CorrectException && "Test failed: caught exception is incorrect."' failed.
# `-----------------------------
# error: command failed with exit status: -6

--

********************
FAIL: SYCL :: GroupAlgorithm/root_group.cpp (1236 of 2185)
******************** TEST 'SYCL :: GroupAlgorithm/root_group.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 3
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/GroupAlgorithm/root_group.cpp -I . -o /__w/llvm/llvm/build-e2e/GroupAlgorithm/Output/root_group.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/GroupAlgorithm/root_group.cpp -I . -o /__w/llvm/llvm/build-e2e/GroupAlgorithm/Output/root_group.cpp.tmp.out
# note: command had no output on stdout or stderr
# RUN: at line 4
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/GroupAlgorithm/Output/root_group.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/GroupAlgorithm/Output/root_group.cpp.tmp.out
# .---command stderr------------
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           209
# | 	Name:            hipErrorNoBinaryForGpu
# | 	Description:     no kernel image is available for execution on the device
# | 	Function:        buildProgram
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:235
# | 
# | terminate called after throwing an instance of 'sycl::_V1::exception'
# |   what():  The program was built for 1 devices
# | Build program log for 'AMD Radeon RX 6700 XT':
# | �#e�(�
# `-----------------------------
# error: command failed with exit status: -6

--

********************
FAIL: SYCL :: KernelAndProgram/kernel-bundle-find-run.cpp (1[59](https://github.com/intel/llvm/actions/runs/11031105056/job/30637542157#step:22:60)7 of 2185)
******************** TEST 'SYCL :: KernelAndProgram/kernel-bundle-find-run.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 1
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/KernelAndProgram/kernel-bundle-find-run.cpp -o /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-find-run.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/KernelAndProgram/kernel-bundle-find-run.cpp -o /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-find-run.cpp.tmp.out
# note: command had no output on stdout or stderr
# RUN: at line 2
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-find-run.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-find-run.cpp.tmp.out
# .---command stdout------------
# | sycl_kernel done
# `-----------------------------
# .---command stderr------------
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           209
# | 	Name:            hipErrorNoBinaryForGpu
# | 	Description:     no kernel image is available for execution on the device
# | 	Function:        buildProgram
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:235
# | 
# | terminate called after throwing an instance of 'sycl::_V1::exception'
# |   what():  The program was built for 1 devices
# | Build program log for 'AMD Radeon RX [67](https://github.com/intel/llvm/actions/runs/11031105056/job/30637542157#step:22:68)00 XT':
# | �G��
# `-----------------------------
# error: command failed with exit status: -6

--

********************
FAIL: SYCL :: KernelAndProgram/kernel-bundle-get-kernel.cpp (1599 of 2185)
******************** TEST 'SYCL :: KernelAndProgram/kernel-bundle-get-kernel.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 1
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/KernelAndProgram/kernel-bundle-get-kernel.cpp -o /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-get-kernel.cpp.tmp.out
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/KernelAndProgram/kernel-bundle-get-kernel.cpp -o /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-get-kernel.cpp.tmp.out
# note: command had no output on stdout or stderr
# RUN: at line 2
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-get-kernel.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/KernelAndProgram/Output/kernel-bundle-get-kernel.cpp.tmp.out
# .---command stderr------------
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           209
# | 	Name:            hipErrorNoBinaryForGpu
# | 	Description:     no kernel image is available for execution on the device
# | 	Function:        buildProgram
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:235
# | 
# | terminate called after throwing an instance of 'sycl::_V1::exception'
# |   what():  The program was built for 1 devices
# | Build program log for 'AMD Radeon RX 6[70](https://github.com/intel/llvm/actions/runs/11031105056/job/30637542157#step:22:71)0 XT':
# | 
# `-----------------------------
# error: command failed with exit status: -6

--

********************
FAIL: SYCL :: SpecConstants/2020/host_apis.cpp (1903 of 2185)
******************** TEST 'SYCL :: SpecConstants/2020/host_apis.cpp' FAILED ********************
Exit Code: -6

Command Output (stdout):
--
# RUN: at line 1
/__w/llvm/llvm/toolchain/bin//clang++  -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa  /__w/llvm/llvm/llvm/sycl/test-e2e/SpecConstants/2020/host_apis.cpp -Wno-error=unused-command-line-argument -o /__w/llvm/llvm/build-e2e/SpecConstants/2020/Output/host_apis.cpp.tmp.out -fsycl-dead-args-optimization
# executed command: /__w/llvm/llvm/toolchain/bin//clang++ -Werror -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl -fsycl-targets=amdgcn-amd-amdhsa /__w/llvm/llvm/llvm/sycl/test-e2e/SpecConstants/2020/host_apis.cpp -Wno-error=unused-command-line-argument -o /__w/llvm/llvm/build-e2e/SpecConstants/2020/Output/host_apis.cpp.tmp.out -fsycl-dead-args-optimization
# .---command stderr------------
# | clang++: warning: argument unused during compilation: '-fsycl-dead-args-optimization' [-Wunused-command-line-argument]
# `-----------------------------
# RUN: at line 2
env ONEAPI_DEVICE_SELECTOR=hip:gpu  /__w/llvm/llvm/build-e2e/SpecConstants/2020/Output/host_apis.cpp.tmp.out
# executed command: env ONEAPI_DEVICE_SELECTOR=hip:gpu /__w/llvm/llvm/build-e2e/SpecConstants/2020/Output/host_apis.cpp.tmp.out
# .---command stderr------------
# | <HIP>[ERROR]: 
# | UR HIP ERROR:
# | 	Value:           209
# | 	Name:            hipErrorNoBinaryForGpu
# | 	Description:     no kernel image is available for execution on the device
# | 	Function:        buildProgram
# | 	Source Location: /__w/llvm/llvm/build/_deps/unified-runtime-src/source/adapters/hip/program.cpp:235
# | 
# | terminate called after throwing an instance of 'sycl::_V1::exception'
# |   what():  The program was built for 1 devices
# | Build program log for 'AMD Radeon RX 6700 XT':
# | @
# `-----------------------------
# error: command failed with exit status: -6

--

sarnex · 2024-09-30T14:36:47Z

@frasercrmck @MartinWehking Can someone take a look at Alexey's comment about this change breaking HIP CI? Thanks.

frasercrmck · 2024-09-30T14:38:07Z

@frasercrmck @MartinWehking Can someone take a look at Alexey's comment about this change breaking HIP CI? Thanks.

Yep we're taking a look 👍

npmiller · 2024-09-30T17:37:13Z

This patch should fix all of the failures:

[SYCL][HIP] Add --only-needed flag to llvm-link for AMDGCN #15558

We'll need to follow up on this to figure out what's going on with the build log, but it should be enough to unblock the CI

MartinWehking changed the title ~~Amd libdev~~ Create one bitcode library for NVPTX Aug 13, 2024

MartinWehking changed the title ~~Create one bitcode library for NVPTX~~ Create one bitcode library for AMD Aug 13, 2024

MartinWehking had a problem deploying to WindowsCILock August 13, 2024 16:01 — with GitHub Actions Failure

MartinWehking mentioned this pull request Aug 13, 2024

[SYCL][NVPTX] Create one bitcode library for NVPTX #15048

Merged

MartinWehking temporarily deployed to WindowsCILock August 13, 2024 17:02 — with GitHub Actions Inactive

MartinWehking force-pushed the amd_libdev branch from d7effee to e57f743 Compare September 13, 2024 10:59

MartinWehking temporarily deployed to WindowsCILock September 13, 2024 11:01 — with GitHub Actions Inactive

MartinWehking temporarily deployed to WindowsCILock September 13, 2024 12:14 — with GitHub Actions Inactive

MartinWehking had a problem deploying to WindowsCILock September 16, 2024 11:54 — with GitHub Actions Failure

MartinWehking temporarily deployed to WindowsCILock September 16, 2024 13:13 — with GitHub Actions Inactive

Martin Wehking added 3 commits September 17, 2024 11:22

Add devicelib linking test

71a64ed

Check if device lib flags get treated correctly for AMDGCN and that the linking actions for the correct device libraries are generated.

Fix unused argument error for devicelib test case

43dfbee

MartinWehking force-pushed the amd_libdev branch from b5ae788 to 43dfbee Compare September 17, 2024 14:46

MartinWehking had a problem deploying to WindowsCILock September 17, 2024 14:47 — with GitHub Actions Failure

MartinWehking temporarily deployed to WindowsCILock September 17, 2024 15:29 — with GitHub Actions Inactive

MartinWehking marked this pull request as ready for review September 20, 2024 12:31

MartinWehking requested review from a team as code owners September 20, 2024 12:31

MartinWehking requested a review from steffenlarsen September 20, 2024 12:31

steffenlarsen approved these changes Sep 20, 2024

View reviewed changes

frasercrmck approved these changes Sep 20, 2024

View reviewed changes

clang/lib/Driver/ToolChains/SYCL.cpp Show resolved Hide resolved

libdevice/cmake/modules/SYCLLibdevice.cmake Outdated Show resolved Hide resolved

elizabethandrews approved these changes Sep 20, 2024

View reviewed changes

smanna12 approved these changes Sep 20, 2024

View reviewed changes

Change naming to a proper backend target

6ceefc6

MartinWehking temporarily deployed to WindowsCILock September 23, 2024 11:28 — with GitHub Actions Inactive

MartinWehking temporarily deployed to WindowsCILock September 23, 2024 12:02 — with GitHub Actions Inactive

mdtoguchi reviewed Sep 23, 2024

View reviewed changes

clang/test/Driver/sycl-device-lib-amdgcn.cpp Outdated Show resolved Hide resolved

Remove not needed flag + add dummy input

a7666f6

MartinWehking temporarily deployed to WindowsCILock September 24, 2024 10:58 — with GitHub Actions Inactive

MartinWehking temporarily deployed to WindowsCILock September 24, 2024 11:56 — with GitHub Actions Inactive

MartinWehking requested a review from mdtoguchi September 24, 2024 15:00

mdtoguchi approved these changes Sep 24, 2024

View reviewed changes

ldrumm merged commit 2f0abc6 into intel:sycl Sep 25, 2024
12 checks passed

MartinWehking mentioned this pull request Sep 25, 2024

Create one bitcode library for AMD and NVPTX #13930

Closed

JackAKirk mentioned this pull request Sep 26, 2024

Not possible to use some math functions with the CUDA and HIP backends #5326

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Create one bitcode library for AMD #15055

Create one bitcode library for AMD #15055

MartinWehking commented Aug 13, 2024 •

edited

Loading

steffenlarsen left a comment

elizabethandrews left a comment

MartinWehking commented Sep 23, 2024

mdtoguchi left a comment

AlexeySachkov commented Sep 27, 2024

sarnex commented Sep 30, 2024

frasercrmck commented Sep 30, 2024

npmiller commented Sep 30, 2024

Create one bitcode library for AMD #15055

Create one bitcode library for AMD #15055

Conversation

MartinWehking commented Aug 13, 2024 • edited Loading

steffenlarsen left a comment

Choose a reason for hiding this comment

elizabethandrews left a comment

Choose a reason for hiding this comment

MartinWehking commented Sep 23, 2024

mdtoguchi left a comment

Choose a reason for hiding this comment

AlexeySachkov commented Sep 27, 2024

sarnex commented Sep 30, 2024

frasercrmck commented Sep 30, 2024

npmiller commented Sep 30, 2024

MartinWehking commented Aug 13, 2024 •

edited

Loading