From ec6fab24505412dba1015eb845cf72b58011eb26 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Sat, 1 Aug 2026 20:52:00 -0500 Subject: [PATCH 1/3] Fix documentation deployment for tagged releases The stable docs URL (https://juliagpu.github.io/oneAPI.jl/stable) 404s because gh-pages only ever contained `dev`: Documenter creates the `stable` symlink and versioned directories only on a tag build, and no tag build has ever run. TagBot pushed tags using `secrets.GITHUB_TOKEN`, and GitHub does not trigger workflow runs for events created with the default token, so `docs.yml` never fired for v2.6.0 through v2.7.1 despite its `tags: '*'` trigger. Push tags over SSH with DOCUMENTER_KEY instead. Also add a `workflow_dispatch` trigger to the Documentation workflow so tagged docs can be re-deployed manually, and fix the canonical URL in docs/make.jl, which was copy-pasted from another project and pointed at ExaPF.jl on every page. --- .github/workflows/TagBot.yml | 3 +++ .github/workflows/docs.yml | 1 + docs/make.jl | 2 +- 3 files changed, 5 insertions(+), 1 deletion(-) diff --git a/.github/workflows/TagBot.yml b/.github/workflows/TagBot.yml index 6887673b..87d6c185 100644 --- a/.github/workflows/TagBot.yml +++ b/.github/workflows/TagBot.yml @@ -14,3 +14,6 @@ jobs: - uses: JuliaRegistries/TagBot@v1 with: token: ${{ secrets.GITHUB_TOKEN }} + # push tags over SSH so they trigger the Documentation workflow; + # tags pushed with GITHUB_TOKEN do not trigger workflow runs + ssh: ${{ secrets.DOCUMENTER_KEY }} diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index 832d8318..9043fe1e 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -9,6 +9,7 @@ on: types: [opened, synchronize, reopened] schedule: - cron: '0 0 * * 0' + workflow_dispatch: jobs: docs: diff --git a/docs/make.jl b/docs/make.jl index b5529648..fe6dd6f2 100644 --- a/docs/make.jl +++ b/docs/make.jl @@ -13,7 +13,7 @@ makedocs( sitename = "oneAPI.jl", format = Documenter.HTML( prettyurls = Base.get(ENV, "CI", nothing) == "true", - canonical = "https://exanauts.github.io/ExaPF.jl/stable/", + canonical = "https://juliagpu.github.io/oneAPI.jl/stable/", mathengine = Documenter.KaTeX(), ), modules = [oneAPI], From 15e81380e4943b35138df8acb8cf0f504cac8f9e Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Sat, 1 Aug 2026 21:06:35 -0500 Subject: [PATCH 2/3] Update README to reflect the current toolchain and feature set The README still described the pre-SPIRVIntrinsics state of the package: - SPIR-V is now produced by LLVM's SPIR-V back-end, not the Khronos SPIRV-LLVM-Translator; refresh the prose and the @device_code_llvm / @device_code_spirv output, which now shows __spirv_ControlBarrier and the new generator string. - Refresh the versioninfo() sample: NEO 26.18, libigc 2.34, SPIRV_Tools 2025.4, oneAPI_Support 0.10 (oneMKL 2025.3), plus the "Julia packages" section that versioninfo() gained. - Update the printed forms of oneArray and ZeDevice, which both changed. - Mention the oneMKL FFT support and the KernelAbstractions oneAPIBackend in the status section, and add a note on writing portable code. - Document deps/generate_interfaces.jl and the Runic formatting check, and give deps/build_local.jl its --project=deps invocation. The sparse example used oneSparseMatrixCSC; switch it to CSR and note that oneMKL's CSR-based sparse back-end leaves the triangular operations unavailable for CSC. Drop the stale "looking for a maintainer" line and the recorded test suite summary, whose counts and timings no longer match. --- README.md | 196 ++++++++++++++++++++++++++++++++++-------------------- 1 file changed, 124 insertions(+), 72 deletions(-) diff --git a/README.md b/README.md index e84e2a83..778dd221 100644 --- a/README.md +++ b/README.md @@ -28,15 +28,18 @@ Windows support is experimental. ## Status -**oneAPI.jl is looking for contributors and/or a maintainer. Reach out if you can help!** +oneAPI.jl is actively maintained and contributions are welcome — see the issue tracker for +open work. -The current version of oneAPI.jl supports most of the oneAPI Level Zero interface, has -good kernel programming capabilties, and as a demonstration of that it fully implements -the GPUArrays.jl array interfaces. This results in a full-featured GPU array type. +The current version of oneAPI.jl supports most of the oneAPI Level Zero interface, has good +kernel programming capabilities, and as a demonstration of that it fully implements the +GPUArrays.jl array interfaces. This results in a full-featured GPU array type. On top of +that, the package integrates with vendor libraries through oneMKL, covering dense BLAS and +LAPACK, sparse linear algebra, and FFTs, and it provides a KernelAbstractions.jl backend +(`oneAPIBackend`) so portable kernels run unmodified on Intel GPUs. -However, the package has not been extensively tested, and performance issues might be -present. The integration with vendor libraries like oneMKL has been extended with support -for sparse linear algebra operations. Some operations may still be unavailable or slow. +The package has not been as extensively tested as more mature back-ends, and performance +issues might be present. Some operations may still be unavailable or slow. ## Quick start @@ -66,54 +69,54 @@ julia> using oneAPI julia> oneAPI.versioninfo() Binary dependencies: -- NEO: 25.35.35096 -- libigc: 1.0.17193+0 -- gmmlib: 22.3.20+0 -- SPIRV_LLVM_Translator: 21 -- SPIRV_Tools: 2025.4.0 -- oneAPI_Support: 0.9.2 (oneMKL v2025.2.0) +- NEO: 26.18.38308+0 +- libigc: 2.34.4+1 +- gmmlib: 22.10.0+0 +- SPIRV_LLVM_Backend: 22.1.8+1 +- SPIRV_Tools: 2025.4.0+0 +- oneAPI_Support: 0.10.0+0 (oneMKL v2025.3.0) Toolchain: -- Julia: 1.11.5 -- LLVM: 16.0.6 +- Julia: 1.12.6 +- LLVM: 18.1.7 + +Julia packages: +- oneAPI.jl: 2.7.2 +- GPUArrays: 11.5.8 +- GPUCompiler: 2.1.1 +- KernelAbstractions: 0.9.42 +- LLVM: 9.11.0 +- SPIRVIntrinsics: 1.1.0 1 driver: -- 00000000-0000-0000-173d-d94201036013 (v1.3.24595, API v1.3.0) +- 00000000-0000-0000-18c7-da2e010395a4 (v1.3.38308, API v1.15.0) -2 devices: -- Intel(R) Graphics [0x56a0] -- Intel(R) HD Graphics P630 [0x591d] +1 device: +- Intel(R) Arc(TM) A750 Graphics ``` -If you have multiple compatible drivers or devices, use the `driver!` and `device!` -functions to configure which one to use in the current task: +The drivers and devices that oneAPI.jl found are listed at the end of that output. You can +also query them directly: ```julia julia> devices() -ZeDevice iterator for 2 devices: -1. Intel(R) Graphics [0x56a0] -2. Intel(R) HD Graphics P630 [0x591d] +ZeDevice iterator for 1 devices: +1. Intel(R) Arc(TM) A750 Graphics julia> device() -ZeDevice(GPU, vendor 0x8086, device 0x56a0): Intel(R) Graphics [0x56a0] - -julia> device!(2) -ZeDevice(GPU, vendor 0x8086, device 0x591d): Intel(R) HD Graphics P630 [0x591d] +ZeDevice(GPU, vendor 0x8086, device 0x56a1) ``` +If more than one compatible driver or device is listed, use the `driver!` and `device!` +functions to configure which one to use in the current task, e.g. `device!(2)` to switch to +the second device. + To ensure other functionality works as expected, you can run the test suite from the package manager REPL mode. Note that this will pull and run the test suite for [GPUArrays](https://github.com/JuliaGPU/GPUArrays.jl), which takes quite some time: ``` pkg> test oneAPI -... -Testing finished in 16 minutes, 27 seconds, 506 milliseconds - -Test Summary: | Pass Total Time - Overall | 4945 4945 - SUCCESS - Testing oneAPI tests passed ``` @@ -136,10 +139,10 @@ julia> drv = first(drivers()); julia> ctx = ZeContext(drv); julia> dev = first(devices(drv)) -ZeDevice(GPU, vendor 0x8086, device 0x1912): Intel(R) Gen9 +ZeDevice(GPU, vendor 0x8086, device 0x56a1) julia> compute_properties(dev) -(maxTotalGroupSize = 256, maxGroupSizeX = 256, maxGroupSizeY = 256, maxGroupSizeZ = 256, maxGroupCountX = 4294967295, maxGroupCountY = 4294967295, maxGroupCountZ = 4294967295, maxSharedLocalMemory = 65536, subGroupSizes = (8, 16, 32)) +(maxTotalGroupSize = 1024, maxGroupSizeX = 1024, maxGroupSizeY = 1024, maxGroupSizeZ = 1024, maxGroupCountX = 4294967295, maxGroupCountY = 4294967295, maxGroupCountZ = 4294967295, maxSharedLocalMemory = 65536, subGroupSizes = (8, 16, 32)) julia> queue = ZeCommandQueue(ctx, dev); @@ -149,8 +152,9 @@ julia> execute!(queue) do list ``` Built on top of that, are kernel programming capabilities for executing Julia code on oneAPI -accelerators. For now, we reuse OpenCL intrinsics, and compile to SPIR-V using [Khronos' -translator](https://github.com/KhronosGroup/SPIRV-LLVM-Translator): +accelerators. Device-side intrinsics are provided by +[SPIRVIntrinsics.jl](https://github.com/JuliaGPU/SPIRVIntrinsics.jl), and code is compiled to +SPIR-V using [LLVM's SPIR-V back-end](https://llvm.org/docs/SPIRVUsage.html): ```julia julia> function kernel() @@ -168,15 +172,19 @@ julia> @device_code_llvm @oneapi items=1 kernel() ``` ```llvm -; @ REPL[18]:1 within `kernel' -define dso_local spir_kernel void @_Z17julia_kernel_3053() local_unnamed_addr { -top: -; @ REPL[18]:2 within `kernel' -; ┌ @ oneAPI.jl/src/device/opencl/synchronization.jl:9 within `barrier' @ oneAPI.jl/src/device/opencl/synchronization.jl:9 -; │┌ @ oneAPI.jl/src/device/opencl/utils.jl:34 within `macro expansion' - call void @_Z7barrierj(i32 0) -; └└ -; @ REPL[18]:3 within `kernel' +; @ REPL[2]:1 within `kernel` +define spir_kernel void @_Z6kernel() local_unnamed_addr { +conversion: + br label %top + +top: ; preds = %conversion +; @ REPL[2]:2 within `kernel` +; ┌ @ SPIRVIntrinsics/src/synchronization.jl:162 within `barrier` +; │┌ @ SPIRVIntrinsics/src/synchronization.jl:154 within `work_group_barrier` +; ││┌ @ SPIRVIntrinsics/src/synchronization.jl:54 within `control_barrier` + call void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 16) +; └└└ +; @ REPL[2]:3 within `kernel` ret void } ``` @@ -187,28 +195,32 @@ julia> @device_code_spirv @oneapi items=1 kernel() ```spirv ; SPIR-V -; Version: 1.0 -; Generator: Khronos LLVM/SPIR-V Translator; 14 -; Bound: 9 +; Version: 1.4 +; Generator: LLVM LLVM SPIR-V Backend; 22 +; Bound: 22 ; Schema: 0 - OpCapability Addresses OpCapability Kernel + OpCapability Addresses %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL - OpEntryPoint Kernel %4 "_Z17julia_kernel_3067" + OpEntryPoint Kernel %_Z6kernel "_Z6kernel" + OpExecutionMode %_Z6kernel ContractionOff OpSource OpenCL_C 200000 + OpName %_Z6kernel "_Z6kernel" + OpName %conversion "conversion" OpName %top "top" + %void = OpTypeVoid + %5 = OpTypeFunction %void %uint = OpTypeInt 32 0 + %uint_16 = OpConstant %uint 16 %uint_2 = OpConstant %uint 2 - %uint_0 = OpConstant %uint 0 - %void = OpTypeVoid - %3 = OpTypeFunction %void - %4 = OpFunction %void None %3 + %_Z6kernel = OpFunction %void None %5 + %conversion = OpLabel + OpBranch %top %top = OpLabel - OpControlBarrier %uint_2 %uint_2 %uint_0 + OpControlBarrier %uint_2 %uint_2 %uint_16 OpReturn OpFunctionEnd - ``` Finally, the `oneArray` type makes it possible to use your oneAPI accelerator without the @@ -216,27 +228,50 @@ need to write custom kernels, thanks to Julia's high-level array abstractions: ```julia julia> a = oneArray(rand(Float32, 2,2)) -2×2 oneArray{Float32,2}: - 0.592979 0.996154 - 0.874364 0.232854 +2×2 oneArray{Float32, 2, oneAPI.oneL0.DeviceBuffer}: + 0.879687 0.990588 + 0.421241 0.157232 julia> a .+ 1 -2×2 oneArray{Float32,2}: - 1.59298 1.99615 - 1.87436 1.23285 +2×2 oneArray{Float32, 2, oneAPI.oneL0.DeviceBuffer}: + 1.87969 1.99059 + 1.42124 1.15723 ``` -The oneMKL integration provides extended support for linear algebra operations, including sparse -matrix operations that integrate with Julia's standard LinearAlgebra interface: +The oneMKL integration provides extended support for linear algebra operations, including +sparse matrix operations that integrate with Julia's standard LinearAlgebra interface: ```julia julia> using oneAPI, oneAPI.oneMKL, SparseArrays, LinearAlgebra -julia> A = sprand(100, 100, 0.1) -julia> dA = oneMKL.oneSparseMatrixCSC(A) -julia> x = oneArray(rand(100)) -julia> y = dA * x # Matrix-vector multiplication via LinearAlgebra + +julia> A = sprand(Float32, 100, 100, 0.1); + +julia> dA = oneMKL.oneSparseMatrixCSR(A); + +julia> x = oneArray(rand(Float32, 100)); + +julia> y = dA * x; # matrix-vector multiplication via LinearAlgebra ``` +Sparse matrices are also available in CSC (`oneSparseMatrixCSC`) and COO +(`oneSparseMatrixCOO`) formats. Note that oneMKL's sparse back-end is CSR-based, so some +operations — notably the triangular solves and multiplications — are unavailable for CSC +matrices and will throw an `ArgumentError`. + +### Writing portable code + +While oneAPI.jl provides Intel-specific functionality, it is recommended to write +backend-agnostic code where possible, so that it runs on other GPU back-ends without +modification: + +- [GPUArrays.jl](https://github.com/JuliaGPU/GPUArrays.jl) for high-level array abstractions +- [KernelAbstractions.jl](https://github.com/JuliaGPU/KernelAbstractions.jl) for kernels that + compile for CPU, CUDA, ROCm, Metal and oneAPI devices. oneAPI.jl provides the + `oneAPIBackend` back-end for this. + +Reach for `oneAPI`-specific macros (like `@oneapi`) and types (like `oneArray`) when you need +optimizations or features that the generic abstractions do not cover. + ### `Float64` support Not all oneAPI GPUs support Float64 datatypes. You can test if your GPU does using @@ -265,7 +300,24 @@ To work on oneAPI.jl, you just need to `dev` the package. In addition, you may n **build the binary support library** that's used to interface with oneMKL and other C++ vendor libraries. This library is normally provided by the oneAPI_Support_jll.jl package, however, we only guarantee to update this package when releasing oneAPI.jl. You can build -this library yourself by simply executing `deps/build_local.jl`. +this library yourself by executing: + +``` +$ julia --project=deps deps/build_local.jl +``` + +This installs a Conda environment with Intel DPC++ and MKL, builds the library, and writes a +`LocalPreferences.toml` that points oneAPI.jl at the result. + +Most of the oneMKL bindings are generated from the oneMKL C++ headers. After changing the +generator or updating oneMKL, regenerate them with: + +``` +$ julia --project=deps deps/generate_interfaces.jl +``` + +Pull requests are checked with [Runic.jl](https://github.com/fredrikekre/Runic.jl); format +your changes with `git runic main` before submitting. To facilitate development, there are other things you may want to configure: From 363e157892570b8247c8d7e038cae282d612c4b7 Mon Sep 17 00:00:00 2001 From: Michel Schanen Date: Sat, 1 Aug 2026 21:09:51 -0500 Subject: [PATCH 3/3] Correct and refresh the documentation Several documented APIs did not match the package, and the examples were verified against an Arc A750 while fixing them: - memory.md called zeMemAllocDevice(context(), device(), 1024, 1), but that raw entry point takes six arguments and the example cannot run. Use the oneL0 device_alloc/free wrappers instead. - barrier takes a mandatory flags argument. kernels.md documented it as barrier(), device.md as barrier(flags=0), and performance.md called it without arguments in five kernels. Pass oneAPI.LOCAL_MEM_FENCE. - device.md documented the indexing intrinsics as 0-based (dim=0); they are 1-based in SPIRVIntrinsics, and so are the indices they return. - onemkl.md demonstrated sparse matvec on a oneSparseMatrixCSC. Switch to CSR and describe the CSC-as-transposed-CSR representation along with the triangular operations it rules out. Its Float64 arrays also needed fp64 support the example never mentioned, so use Float32. - The FFT example loaded FFTW, which is not a dependency; AbstractFFTs is what the oneMKL DFT integration hooks into. Also record the SPIRVIntrinsics and KernelAbstractions integrations and the SPIR-V back-end in the feature list, align the OS and hardware requirements with installation.md, and mention EnableDebugBreak and ONEAPI_MEMORY_LIMIT in the troubleshooting guide. --- docs/src/device.md | 21 +++++++++++++-------- docs/src/index.md | 12 ++++++++---- docs/src/kernels.md | 7 +++++-- docs/src/memory.md | 15 +++++++++++---- docs/src/onemkl.md | 16 ++++++++++------ docs/src/troubleshooting.md | 3 +++ docs/src/usage/performance.md | 10 +++++----- 7 files changed, 55 insertions(+), 29 deletions(-) diff --git a/docs/src/device.md b/docs/src/device.md index acdb4d6e..7f708966 100644 --- a/docs/src/device.md +++ b/docs/src/device.md @@ -1,21 +1,26 @@ # Device Intrinsics -When writing custom kernels, you have access to a set of device intrinsics that map to underlying hardware instructions. +When writing custom kernels, you have access to a set of device intrinsics that map to underlying hardware instructions. These are provided by [SPIRVIntrinsics.jl](https://github.com/JuliaGPU/SPIRVIntrinsics.jl) and re-exported by oneAPI.jl. ## Indexing These functions allow you to determine the current thread's position in the execution grid. -- `get_global_id(dim=0)`: Global index of the work item. -- `get_local_id(dim=0)`: Local index of the work item within the workgroup. -- `get_group_id(dim=0)`: Index of the workgroup. -- `get_global_size(dim=0)`: Global size of the ND-range. -- `get_local_size(dim=0)`: Size of the workgroup. -- `get_num_groups(dim=0)`: Number of workgroups. +- `get_global_id(dim=1)`: Global index of the work item. +- `get_local_id(dim=1)`: Local index of the work item within the workgroup. +- `get_group_id(dim=1)`: Index of the workgroup. +- `get_global_size(dim=1)`: Global size of the ND-range. +- `get_local_size(dim=1)`: Size of the workgroup. +- `get_num_groups(dim=1)`: Number of workgroups. + +Unlike their OpenCL counterparts, both the dimension argument and the returned indices are +1-based, so the result can be used to index a Julia array directly. ## Synchronization -- `barrier(flags=0)`: Synchronizes all work items in a workgroup. +- `barrier(flags)`: Synchronizes all work items in a workgroup. The `flags` argument selects + which memory operations are fenced: `oneAPI.LOCAL_MEM_FENCE`, `oneAPI.GLOBAL_MEM_FENCE`, or + their bitwise or. There is no default; a value must be passed. ## Atomics diff --git a/docs/src/index.md b/docs/src/index.md index 988d9e92..ad959992 100644 --- a/docs/src/index.md +++ b/docs/src/index.md @@ -16,14 +16,18 @@ Direct use of `oneAPI`-specific macros (like `@oneapi`) and types (like `oneArra ## Features - **High-level Array Abstractions**: `oneArray` type fully implementing the `GPUArrays.jl` interface. -- **Kernel Programming**: Execute custom kernels written in Julia on Intel GPUs. +- **Kernel Programming**: Execute custom kernels written in Julia on Intel GPUs, compiled to + SPIR-V through LLVM's SPIR-V back-end. +- **KernelAbstractions.jl Backend**: The `oneAPIBackend` back-end runs portable kernels on + Intel GPUs unmodified. - **Level Zero Integration**: Low-level access to the Level Zero API via the `oneL0` submodule. -- **oneMKL Support**: Integration with Intel oneMKL for BLAS, LAPACK, and sparse operations. +- **oneMKL Support**: Integration with Intel oneMKL for BLAS, LAPACK, sparse and FFT operations. - **SYCL Integration**: Interoperability with SYCL (on Linux). ## Requirements - **Julia**: 1.10 or higher -- **OS**: Linux -- **Hardware**: Intel Gen9 graphics or newer (including Intel Arc A-Series) +- **OS**: Linux. Windows is experimental, via WSL2. +- **Hardware**: Intel Gen9 graphics or newer. Intel Arc GPUs (A580, A750, A770, etc) require + Linux 6.2 or newer. diff --git a/docs/src/kernels.md b/docs/src/kernels.md index b5e03b12..07e50c7a 100644 --- a/docs/src/kernels.md +++ b/docs/src/kernels.md @@ -55,7 +55,10 @@ Inside a kernel, you can use various intrinsics to interact with the hardware: - `get_global_id()`: Get the global thread ID. - `get_local_id()`: Get the local thread ID within a workgroup. - `get_group_id()`: Get the workgroup ID. -- `barrier()`: Synchronize threads within a workgroup. +- `barrier(flags)`: Synchronize threads within a workgroup. -These correspond to standard OpenCL/Level Zero intrinsics. +These are provided by [SPIRVIntrinsics.jl](https://github.com/JuliaGPU/SPIRVIntrinsics.jl) +and correspond to the standard OpenCL built-in functions. Note that the indices they return +are 1-based, so they can be used to index Julia arrays directly. See +[Device Intrinsics](device.md) for the full list. diff --git a/docs/src/memory.md b/docs/src/memory.md index 0e4ea1e9..18bb1d7c 100644 --- a/docs/src/memory.md +++ b/docs/src/memory.md @@ -17,15 +17,22 @@ oneAPI uses Unified Shared Memory, which allows for pointers that can be accessi You can perform low-level memory allocation using the `oneL0` submodule if needed, though `oneArray` handles this automatically. ```julia -using oneAPI.oneL0 +using oneAPI, oneAPI.oneL0 -# Allocate device memory -ptr = oneL0.zeMemAllocDevice(context(), device(), 1024, 1) +drv = first(drivers()) +ctx = ZeContext(drv) +dev = first(devices(drv)) + +# Allocate 1024 bytes of device memory, aligned to 8 bytes +buf = device_alloc(ctx, dev, 1024, 8) # Free memory -oneL0.zeMemFree(context(), ptr) +free(buf) ``` +`host_alloc` and `shared_alloc` allocate host and shared USM memory respectively, using the +same signature. + ## Garbage Collection Julia's garbage collector automatically manages `oneArray` objects. However, GPU memory is a limited resource. If you are running into out-of-memory errors, you might need to manually trigger garbage collection or free arrays. diff --git a/docs/src/onemkl.md b/docs/src/onemkl.md index 91f2645c..640f2a68 100644 --- a/docs/src/onemkl.md +++ b/docs/src/onemkl.md @@ -27,26 +27,30 @@ oneAPI.jl supports sparse matrix operations via oneMKL's sparse BLAS functionali using oneAPI, oneAPI.oneMKL, SparseArrays, LinearAlgebra # Create a sparse matrix on CPU -A = sprand(100, 100, 0.1) +A = sprand(Float32, 100, 100, 0.1) # Move to GPU (converts to oneMKL format) -dA = oneMKL.oneSparseMatrixCSC(A) +dA = oneMKL.oneSparseMatrixCSR(A) # Create a dense vector -x = oneArray(rand(100)) +x = oneArray(rand(Float32, 100)) # Sparse matrix-vector multiplication y = dA * x ``` -Note that `oneSparseMatrixCSC` is available for Compressed Sparse Column format, which is the standard in Julia. +Three storage formats are available: `oneSparseMatrixCSR`, `oneSparseMatrixCSC` and +`oneSparseMatrixCOO`. oneMKL's sparse back-end is CSR-based, and a `oneSparseMatrixCSC` is +therefore stored as the CSR representation of its transpose. As a consequence the triangular +operations (`sparse_trmv!`, `sparse_trsv!`, `sparse_trsm!`) cannot be expressed for CSC +matrices and throw an `ArgumentError`. Prefer CSR when you have the choice. ## FFTs -Fast Fourier Transforms are supported through `AbstractFFTs.jl` interface integration with oneMKL DFTs. +Fast Fourier Transforms are supported through `AbstractFFTs.jl` interface integration with oneMKL DFTs. oneAPI.jl depends on AbstractFFTs.jl, so no separate FFT package is required. ```julia -using oneAPI, FFTW +using oneAPI, AbstractFFTs a = oneArray(rand(ComplexF32, 1024)) diff --git a/docs/src/troubleshooting.md b/docs/src/troubleshooting.md index 9f790f65..936b435c 100644 --- a/docs/src/troubleshooting.md +++ b/docs/src/troubleshooting.md @@ -33,6 +33,8 @@ Some Intel GPUs (especially integrated graphics) lack native hardware support fo - Trigger garbage collection: `GC.gc()`. - Manually free unused arrays: `oneAPI.unsafe_free!(array)`. - Check if you are exceeding the device's memory capacity. +- When several processes share one GPU, give each a budget with `ONEAPI_MEMORY_LIMIT`; see + [Memory Management](memory.md). ## Debugging @@ -43,6 +45,7 @@ Enable the Level Zero validation layer to catch API misuse: ```bash export ZE_ENABLE_VALIDATION_LAYER=1 export ZE_ENABLE_PARAMETER_VALIDATION=1 +export EnableDebugBreak=0 # works around intel/compute-runtime#639 ``` ### Debug Mode diff --git a/docs/src/usage/performance.md b/docs/src/usage/performance.md index c7064aee..ba5e641d 100644 --- a/docs/src/usage/performance.md +++ b/docs/src/usage/performance.md @@ -116,7 +116,7 @@ function optimized_reduction!(result, input) # Load global → local (coalesced) global_id = get_global_id() @inbounds local_mem[local_id] = input[global_id] - barrier() + barrier(oneAPI.LOCAL_MEM_FENCE) # Reduce in local memory (much faster) stride = local_size ÷ 2 @@ -124,7 +124,7 @@ function optimized_reduction!(result, input) if local_id <= stride @inbounds local_mem[local_id] += local_mem[local_id + stride] end - barrier() + barrier(oneAPI.LOCAL_MEM_FENCE) stride ÷= 2 end @@ -145,9 +145,9 @@ Barriers have overhead: function wasteful_kernel!(a) i = get_local_id() a[i] += 1 - barrier() # Not needed if no data sharing + barrier(oneAPI.LOCAL_MEM_FENCE) # Not needed if no data sharing a[i] *= 2 - barrier() # Not needed + barrier(oneAPI.LOCAL_MEM_FENCE) # Not needed return end @@ -157,7 +157,7 @@ function efficient_kernel!(a, shared) # Load to shared memory shared[i] = a[i] - barrier() # Needed: ensure all loads complete + barrier(oneAPI.LOCAL_MEM_FENCE) # Needed: ensure all loads complete # Use shared data result = shared[i] + shared[i+1]