Skip to content
New issue

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

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

Already on GitHub? Sign in to your account

[SYCL][E2E] Ensuring lowering of llvm.bitreverse is functionally correct #12774

Merged
merged 18 commits into from
Mar 21, 2024

Conversation

LU-JOHN
Copy link
Contributor

@LU-JOHN LU-JOHN commented Feb 20, 2024

Ensure that lowering of llvm.bitreverse* intrinsics by llvm-spirv is functionally correct.

@LU-JOHN LU-JOHN requested a review from a team as a code owner February 20, 2024 22:43
@LU-JOHN LU-JOHN marked this pull request as draft February 20, 2024 22:43
Copy link
Contributor

@sarnex sarnex left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I will leave in-depth review of the test logic to others.

sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
@@ -0,0 +1,189 @@
// Test that llvm.bitreverse is lowered correctly by llvm-spirv

// RUN: %{build} -o %t.O2.out -O2
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we really need to pass O2 here?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Optimization is necessary to ensure that functions are recognized as having bitreverse functionality and optimized in LLVM IR to llvm.bitreverse.* intrinsics.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

-O2 would not be necessary if we could use __builtin_elementwise_bitreverse for all cases, but there is an issue with this builtin for 8-bit and 16-bit types.

sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
@LU-JOHN LU-JOHN changed the title [NFC] Ensuring lowering of llvm.bitreverse is functionally correct [SYCL][E2E] Ensuring lowering of llvm.bitreverse is functionally correct Feb 26, 2024
@LU-JOHN
Copy link
Contributor Author

LU-JOHN commented Mar 6, 2024

There are a few remaining questions/issues with this testcase:

  1. This testcase is a draft because it is waiting for the llvm-spirv changes to be pulled in.
  2. This testcase cannot use __builtin_elementwise_bitreverse on 8-bit and 16-bit data. This testcase has a workaround to avoid this issue: Unexpected codegen for __builtin_elementwise_bitreverse on unsigned char and short  llvm/llvm-project#84047
  3. This testcase also builds without explicitly disabling SPV_KHR_bitinstructions (i.e. the straightforward compile options). This fails to execute because of an unsupported capability. We recently enabled SPV_KHR_bitinstructions in the driver: [SYCL] Enable SPV_KHR_bit_instructions for SYCL code #12754 . Is this safe?
# .---command stderr------------
# | terminate called after throwing an instance of 'sycl::_V1::compile_program_error'
# |   what():  The program was built for 1 devices
# | Build program log for 'Intel(R) FPGA Emulation Device':
# | Compilation started
# | Unsupported SPIR-V module
# | SPIRV module requires unsupported capability 6025
# | Compilation failed
# |  -11 (PI_ERROR_BUILD_PROGRAM_FAILURE)
# `-----------------------------
# error: command failed with exit status: -6

@LU-JOHN LU-JOHN force-pushed the test_llvm_bitreverse_lowering branch from 3581e87 to e9f0c61 Compare March 8, 2024 15:54
@LU-JOHN LU-JOHN marked this pull request as ready for review March 8, 2024 15:55
@LU-JOHN
Copy link
Contributor Author

LU-JOHN commented Mar 8, 2024

llvm-spirv changes to lower llvm.bitreverse.* have been pulled in. Thus, test is no longer a draft and ready for review.

Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
…explicitly disabling SPV_KHR_bit_instructions passes

Signed-off-by: Lu, John <john.lu@intel.com>
…vm.bitreverse is not supported

Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/t.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/t.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/t.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/t.cpp Outdated Show resolved Hide resolved
sycl/test-e2e/LLVMIntrinsicLowering/t.cpp Outdated Show resolved Hide resolved
Signed-off-by: Lu, John <john.lu@intel.com>
Signed-off-by: Lu, John <john.lu@intel.com>
@steffenlarsen steffenlarsen merged commit 597619c into intel:sycl Mar 21, 2024
11 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

6 participants