Atomic operations on local memory doesn't work on integrated GPU

[AlexanderKalistratov]

Reproducer:

# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation
#
# SPDX-License-Identifier: Apache-2.0

import dpnp as np

import numba_dpex as ndpx


@ndpx.kernel
def atomic_reduction(a):
    local_a_0 = ndpx.local.array(1, dtype=a.dtype)
    lid = ndpx.get_local_id(0)

    if lid == 0:
        local_a_0[0] = 0

    ndpx.barrier(ndpx.LOCAL_MEM_FENCE)

    idx = ndpx.get_global_id(0)
    ndpx.atomic.add(local_a_0, 0, a[idx])

    ndpx.barrier(ndpx.LOCAL_MEM_FENCE)

    ndpx.atomic.add(a, 0, a[idx])


def main():
    N = 10
    a = np.ones(N, dtype='float32')

    print("Using device ...")
    print(a.device)

    atomic_reduction[ndpx.Range(N)](a)
    print("Reduction sum =", a[0])

    print("Done...")


if __name__ == "__main__":
    main()

Result:

Using device ...
Device(level_zero:gpu:0)
error: 1: Conflicting SPIR-V versions: 1.4 (input modules 1 through 1) vs 1.0 (input module 2).
Traceback (most recent call last):
  File "/home/akalistr/repo/./atomic_op.py", line 42, in <module>
    main()
  File "/home/akalistr/repo/./atomic_op.py", line 35, in main
    atomic_reduction[ndpx.Range(N)](a)
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/core/kernel_interface/dispatcher.py", line 455, in __call__
    ) = self._compile_and_cache(
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/core/kernel_interface/dispatcher.py", line 141, in _compile_and_cache
    kernel.compile(
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/core/kernel_interface/spirv_kernel.py", line 161, in compile
    self._device_driver_ir_module = spirv_generator.llvm_to_spirv(
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/spirv_generator.py", line 250, in llvm_to_spirv
    return mod.finalize()
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/spirv_generator.py", line 193, in finalize
    self._cmd.link(spirv_path, binary_paths)
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/spirv_generator.py", line 108, in link
    check_call(["spirv-link", *flags, "-o", opath, *binaries])
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/spirv_generator.py", line 27, in check_call
    return _real_check_call(*args, **kwargs)
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/subprocess.py", line 369, in check_call
    raise CalledProcessError(retcode, cmd)
subprocess.CalledProcessError: Command '['spirv-link', '--allow-partial-linkage', '-o', '/tmp/tmpkewpm7n7/2-linked-spirv', '/tmp/tmpkewpm7n7/1-generated-spirv', '/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/ocl/atomics/atomic_ops.spir']' returned non-zero exit status 1.
Exception ignored in: <function Module.__del__ at 0x7f76bf085900>
Traceback (most recent call last):
  File "/home/akalistr/miniconda3/envs/dpbench-dev/lib/python3.10/site-packages/numba_dpex/spirv_generator.py", line 128, in __del__
    os.unlink(afile)
FileNotFoundError: [Errno 2] No such file or directory: '/tmp/tmpkewpm7n7/2-linked-spirv'

Expected result:

Using device ...
Device(level_zero:gpu:0)
Reduction sum = 11.
Done...

NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 fix the issue on 0.21.dev1 but doesn't work on latest main.

[AlexanderKalistratov]

On CPU doesn't work either.
NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 doesn't work on numba-dpex 0.21

[fcharras]

I'm encountering the same error with numba_dpex==0.21.1 (latest release) for some of my kernels, I confirm that NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 does fix the issue on gpu.

[fcharras]

Apparently this has been fixed by 2032e47 (by reverting the spirv bumps)

[ZzEeKkAa]

Fixed in #1103