[Driver][SYCL][FPGA] Improve AOCX behaviors with -fsycl-link=image #10499
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…image Use of -fsycl-link=image allows for FPGA users to create a full device binary and store it in an archive to be used later. This archive can then be picked up and linked in. The information stored in the archive was not sufficient, causing the device binary to not be 'visible' when linked in. Improve the situation here by saving all of the information provided by the wrapping step (including symbols) in the archive, but also be sure that this entire final binary is saved as it would be if we were performing a full E2E compilation.
mdtoguchi
changed the title
- clang format - improved behaviors from -fsycl-link=early type images - some test updates
Needed to properly handle run_launcher.
The main goal of this refactoring is to reduce amount of includes to the context.hpp header. Two new includes are added: 1. sycl/include/sycl/detail/impl_utils.hpp is intended to replace sycl/include/sycl/detail/common.hpp when only interopability with "impl" class is needed. common.hpp include is overloaded with a bunch of other functionality. 2. sycl/include/sycl/async_handler.hpp separates async_handler definiton from exception.hpp.
This adds references to unused function parameters to avoid build errors.
added a "__builtin_intel_sycl_ptr_annotation(ptr, PROPERTY1, VALUE1, PROPERTY2, VALUE2 ...)" builtin function for FPGA SYCL ptr annotation using SYCL extension properties. Updated the "annotated_ptr.hpp" to use "__builtin_intel_sycl_ptr_annotation" to pass annotation instead of using "add_sycl_ir_member_annotation". This helps to insert the llvm.ptr.annotation at the correct position Update the compileTimeProperties pass to consume "sycl-alignment" and thus do not propagate this annotation to SPIRV translator --------------------------------------------------------------------- After using "__builtin_intel_sycl_ptr_annotation" , IR generated are changed from: ``` %ptr.to.ptr = GEP %ptr.to.ptr.annotated = llvm.ptr.annotation %ptr.to.ptr %ptr = load %ptr.to.ptr.annotated %4 = load %ptr ``` to ``` %ptr.to.ptr = GEP %ptr = load %ptr.to.ptr %ptr.annotated = llvm.ptr.annotation %ptr %4 = load %ptr.annotated ``` Thus, SYCLLowerIR pass are expecting different IR forms and thus these two changes are merged together.
This adds references to unused function parameters to avoid build errors. --------- Co-authored-by: Vyacheslav Klochkov <[email protected]>
…tel#10456) This test requires an autogenerated file (DeviceConfigFile.inc) to be present at the time of its compilation, and it was possible in some scenarios that the autogenerated file wasn't generated before the compilation of the test, resulting in an error. Moving the test to `sycl/test` should prevent that from happening. --------- Signed-off-by: Maronas, Marcos <[email protected]>
…ull (intel#10277) Signed-off-by: jinge90 <[email protected]>
* Remove temporary experimental task sycl_linux_gen12_exp.yml * Add possibility to trigger E2E testing on a nightly image manually So far it needs a branch inside intel/llvm but I plan on extending it to PRs later.
…#10455) In the sycl-post-link tool, most of the changes towards moving the passes to use the new pass manager were added as part of this PR by Konstantin Bobrovsky: intel#6835 There was one site where legacy pass manager was used. This change moves that to use the new pass manager. It is sufficient to check for regression in existing tests. There is no need to add a new test. Thanks Sincerely --------- Signed-off-by: Arvind Sudarsanam <[email protected]>
The -fsycl-device-obj option replaces the -fsycl-use-bitcode option. The -fsycl-use-bitcode option is deprecated as well. Update the documentation to reflect this fact, also update any testing using the older option with a replacement of the new option. The use of -fsycl-use-bitcode should also emit the proper deprecation message when used on the command line.
Signed-off-by: Arvind Sudarsanam <[email protected]>
We've removed support for E2E tests on nightly compiler/image in pre-commit, so these aren't used anymore.
enabled Author: Joshua Cranmer <[email protected]> Signed-off-by: Sidorov, Dmitry <[email protected]>
Most of the changes are adding -emit-opaque-pointers=0 lines to test code. The code generally works in the forward translation at this point, although there is still substantial work that needs to be done to finish porting the tests. This test is not adapted multiple_user_semantic_on_struct.ll Author: Joshua Cranmer <[email protected]> Signed-off-by: Sidorov, Dmitry <[email protected]>
Remove `opencl.queue_t.ll` as it would be identical to `spirv.Queue.ll` after migrating. Author: Sven van Haastregt <[email protected]> Signed-off-by: Sidorov, Dmitry <[email protected]>
Signed-off-by: Sidorov, Dmitry <[email protected]>
Made the bindless image sampler member names snake case to match SYCL convention.
Reorganize annotated_ptr/arg folder sycl/ext ----oneapi/ --------experimental --------------annotated_ptr --------------------annotated_ptr.hpp --------------annotated_arg --------------------annotated_arg.hpp --------------common_annotated_properties --------------------properties.hpp ----intel/ --------experimental --------------fpga_annotated_properties.hpp --------------fpga_kernel_properties.hpp --------- Co-authored-by: Dmitry Vodopyanov <[email protected]>
… - USM and SLM (intel#10041) This patch replaces some uses of GenX intrinsics with regular LLVM IR. This change allowed to remove most of restrictions that previously accompanied slm_block_load/store() and block_load/store() accepting USM ptr. GPU RT/driver can lower 'load' instructions to one or several flat-loads and/or gather instructions depending on the target device, alignment and length. Same for 'store' instructions. --------- Signed-off-by: Vyacheslav N Klochkov <[email protected]>
# Fix Mipmap Tests Fixing the computation of expected output values so that the tests also work with input sizes that are not powers of 2. --------- Co-authored-by: Dmitry Vodopyanov <[email protected]>
This pull request introduces a new stand alone library, `SYCLcompat`: a simplified wrapper on top of SYCL, aiming to make it more accessible to developers familiar with other heterogeneous programming models. SYCLcompat has two primary goals: - Improve the adoption of SYCL. This library is designed to provide a familiar programming interface that resembles other popular heterogeneous programming models. By reducing the learning curve, it enables developers to leverage SYCL's power and features more easily. - Source-to-Source Translation Support. SYCLcompat is also designed to facilitate automatic source-to-source translation from other heterogeneous programming models to SYCL and offer a more standardized and consistent programming interface. This feature can significantly streamline the migration and integration of existing codebases into the SYCL ecosystem. The first commit of this PR includes the proposed library README, providing explanation of its motivation, public interface, usage guidelines, and code examples. A set of PRs will follow, including subsets of the current implementation including their tests. We are open to any suggestions, concerns, or improvements you may have, so please, let us know if you have any. Edit: Updated from extension to stand alone library. intel#9976 `dims.hpp` and `defs.hpp` headers --------- Co-authored-by: Gordon Brown <[email protected]> Co-authored-by: Joe Todd <[email protected]> Co-authored-by: Pietro Ghiglio <[email protected]> Co-authored-by: Ruyman Reyes <[email protected]> Co-authored-by: Steffen Larsen <[email protected]> Co-authored-by: aelovikov-intel <[email protected]> Co-authored-by: Sami Hatna <[email protected]> Co-authored-by: Joe Todd <[email protected]> Co-authored-by: Alexey Bader <[email protected]>
The recent sycl::vec changes (intel#9492) broke they unary operations. This PR fixes them and adds some testing to avoid that in the future.
…invocation (intel#10495) This PR adds support to multiple SYCL targets alongside `native_cpu` in the same compiler invocation (e.g. `clang++ -fsycl -fsycl-targets=native_cpu,spir64 input.cpp`). In order to implement this we had to make changes to multiple components, here is a quick overview: * Driver: changes in the Driver allow to correctly parse all the targets passed to `-fsycl-targets` (before we were just looking for `native_cpu`, ignoring the others). The Driver now also calls `sycl-post-link` and `clang-offload-wrapper`, performing a compilation flow more similar to the one used for other targets. * Sema: since the kernel name needs to be the same for all the SYCL targets, the change to the kernel name in Sema has been removed, and replaced with an LLVM Pass that gets run when lowering the device module (`llvm/lib/SYCLLowerIR/RenameKernelSYCLNativeCPU.cpp`). * Runtime: The definition for `_pi_program` in the Native CPU Plug-In now supports multiple kernels in one program, and the `__SYCL_PI_DEVICE_BINARY_TARGET_NATIVE_CPU` binary type has been added in order to identify kernels compiled for Native CPU. * clang-offload-wrapper: for Native CPU, the offload-wrapper doesn't bundle the device code in the host module, but instead produces an array containing function declarations that are resolved by the linker, see `sycl/doc/design/SYCLNativeCPU.md` for more information. --------- Co-authored-by: Uwe Dolinsky <[email protected]> Co-authored-by: Alexey Sachkov <[email protected]> Co-authored-by: Steffen Larsen <[email protected]>
This is an implementation of SYCLcompat, the stand alone library proposed in intel#9646. Please, find the documentation there. SYCLcompat has two primary goals: - Improve the adoption of SYCL. This library is designed to provide a familiar programming interface that resembles other popular heterogeneous programming models. By reducing the learning curve, it enables developers to leverage SYCL's power and features more easily. - Source-to-Source Translation Support. SYCLcompat is also designed to facilitate automatic source-to-source translation from other heterogeneous programming models to SYCL and offer a more standardized and consistent programming interface. This feature can significantly streamline the migration and integration of existing codebases into the SYCL ecosystem The PR also includes tests. We were advised to include the tests as part of `sycl/unittests`, so we added a different mechanism to build them using `clang`. As we stated in the docs PR, we are open to any suggestions, concerns, or improvements you may have, so please, let us know if you have any. --------- Co-authored-by: Gordon Brown <[email protected]> Co-authored-by: Joe Todd <[email protected]> Co-authored-by: Pietro Ghiglio <[email protected]> Co-authored-by: Ruyman Reyes <[email protected]> Co-authored-by: tomflinda <[email protected]>
Fix static analyzer concern about dereferencing null value.
…10650) This PR changes `SYCLPropagateAspectsPass` to propagate aspects that come from `sycl_declared_aspects` even if they are excluded. The reason for this change is because a test like `no-fp64-optimization-declared-aspects.cpp` added in this PR would failed before with higher optimization level because - on the first aspect propagation pass, `fp64` is not propagated (to allow for trivial uses of `float x = 1.5` to optimized out) - the call to the function marked with `device_has(fp64)` is inlined on higher optimizations - that function does not actually use `double` in its body which means no usage of double ends up in the optimized function, leading the second aspect propagation pass to not attach `fp64` to its used aspects metadata. --------- Co-authored-by: Alexey Sachkov <[email protected]> Co-authored-by: Marcos Maronas <[email protected]>
…l#10740) Scheduler::checkLeavesCompletion checks status of all leaves of the buffer to see whether we can destroy that sycl::buffer. There are many scenarios when alloca commands are leaves, these commands don't have associated event and currently they are always incorrectly considered "in progress" because of that preventing buffers to be destroyed timely and deferring their destruction till the point of program termination. Skip alloca commands to fix that.
…ghtly (intel#10745) Also changes daily build upload (via github release) to only require successful build/LIT and not E2E tests. In addition to that, I switched CUDA E2E from using AWS runner to using our self-hosted one. The load of this workflow is low so our single runner can handle that. --------- Co-authored-by: Steffen Larsen <[email protected]>
intel#10743) This eliminates the usage of matrix generator in post-commit, following a similar change done in pre-commit earlier.
The underlying issue in the GPU RT seems to have been fixed, verified locally. Closes intel#8103.
…intel#10716) This patch updates sycl_ext_oneapi_device_architecture extension with the following: * add missing architecture: * intel_gpu_ehl and its sibling - intel_gpu_jsl * add missing aliases to existing architectures: * intel_gpu_bxt (alias to intel_gpu_apl) * intel_gpu_dg2_g10 (alias to intel_gpu_acm_g10) * intel_gpu_dg2_g11 (alias to intel_gpu_acm_g11) * intel_gpu_dg2_g12 (alias to intel_gpu_acm_g12) * remove unnecessary implementation details for intel_gpu_rpl_s from the driver and make it dependent on intel_gpu_adl_s as they are the same architectures
This PR adds C++ wrappers in intel math header for type conversion functions provided in SYCL libdevice. The function names align with CUDA correspondence but we don't use "__" prefix in function name. For example, CUDA __double2float_ru maps to sycl::ext::intel::math::double2float_ru. --------- Signed-off-by: jinge90 <[email protected]>
…0620) When a kernel with a large (unsupported) work-group size is submitted, the backend would fail upon trying build the program, and throw the wrong exception (compile_program_error with `errc::build`, but it should just be a `sycl::exception` with `errc::kernel_not_supported` according to [5.8.1. Kernel attributes](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.attributes)). https://github.com/intel/llvm/blob/b65969014f001f9730349a5caad5c2b85c9bc378/sycl/source/detail/program_manager/program_manager.cpp#L749 This PR adds a check before building to validate the `reqd_work_group_size`. Additionally, another bug in `sycl-post-link` was fixed. The bug occurred when a kernel was decorated with a multi-dimensional required work-group size and two dimensions had the same value. Due to how `sycl-post-link` worked, the attached metadata on the device image would include fewer dimensions than specified. (e.g. If the decorated with `reqd_work_group_size(16, 16)`, then the device image would appear as if decorated by `reqd_work_group_size(16)`.)
Fixed the native_cpu plugin only being built if level_zero was enabled as well
Signed-off-by: Sidorov, Dmitry <[email protected]>
This commit changes the builtin headers from defining all SYCL builtins to using a number of generated header files containing these builtin functions. --------- Signed-off-by: Larsen, Steffen <[email protected]>
…ntel#10767) Linux had been updated in intel#10720 already.
L0 already provides an interface to access PCI information without needing to set an environment variable to access SYSMAN interfaces. --------- Signed-off-by: Jaime Arteaga <[email protected]>
To accomodate the ability to pass around the needed symbols and properties that are used when wrapping the final binaries for any AOCR or AOCX type archives, update the way this information is passed around. Multiple binaries are added to the archive, including the full wrapped .bc file that is gathered and passed along to any subsequent wrapping sequence that occurs.
|
Bad rebase - closing. New PR here: #11592 |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Use of -fsycl-link=image allows for FPGA users to create a full device binary and store it in an archive to be used later. This archive can then be picked up and linked in. The information stored in the archive was not sufficient, causing the device binary to not be 'visible' when linked in.
Improve the situation here by saving all of the information provided by the wrapping step (including symbols) in the archive, but also be sure that this entire final binary is saved as it would be if we were performing a full E2E compilation.