KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator

SPIR-V fails to validate with: AtomicCompareExchange: expected Pointer to point to a value of type Result Type

Open
#2,025 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
LLVM
Stars
625
Forks
279
Avg merge
3d 5h
Merged PRs (30d)
34

Description

Possibly (but i'm not 100% sure) this is identical to issue #2024.

...while trying to run OpenCL-CTS with SPIR-V offload compilation, this source:

```
#if 1
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
#endif
typedef struct BinNode {
int value;
atomic_uintptr_t pNext;
} BinNode;
__kernel void build_hash_table(__global uint* input, __global BinNode* pNodes, volatile __global atomic_uint* pNumNodes, uint numBins)
{
__global BinNode *pNew = &pNodes[ atomic_fetch_add_explicit(pNumNodes, 1u, memory_order_relaxed, memory_scope_all_svm_devices) ];
uint i = get_global_id(0);
uint b = input[i] % numBins;
pNew->value = input[i];
uintptr_t next = atomic_load_explicit(&(pNodes[b].pNext), memory_order_seq_cst, memory_scope_all_svm_devices);
do
{
atomic_store_explicit(&(pNew->pNext), next, memory_order_seq_cst, memory_scope_all_svm_devices);
} while(!atomic_compare_exchange_strong_explicit(&(pNodes[b].pNext), &next, (uintptr_t)pNew, memory_order_seq_cst, memory_order_relaxed, memory_scope_all_svm_devices));
}
```

compiled with Clang 16.0.4 && llvm-spirv branch `llvm_release_160`:

```
$CLANG --target=spir64-unknown-unknown -x cl -cl-std=CL3.0 -D__OPENCL_VERSION__=300 -D__OPENCL_C_VERSION__=300 -D__ENDIAN_LITTLE__=1 -D__IMAGE_SUPPORT__=1 -D__opencl_c_named_address_space_builtins=1 -D__opencl_c_int64=1 -D__opencl_c_atomic_order_acq_rel=1 -D__opencl_c_atomic_order_seq_cst=1 -D__opencl_c_atomic_scope_device=1 -D__opencl_c_program_scope_global_variables=1 -D__opencl_c_generic_address_space=1 -Dcl_khr_byte_addressable_store -Dcl_khr_global_int32_base_atomics -Dcl_khr_global_int32_extended_atomics -Dcl_khr_local_int32_base_atomics -Dcl_khr_local_int32_extended_atomics -Dcl_khr_3d_image_writes -D__opencl_c_3d_image_writes=1 -Dcl_khr_command_buffer -Dcl_khr_subgroups -Dcl_intel_unified_shared_memory -Dcl_khr_spir -Dcl_khr_il_program -Dcl_khr_fp64 -D__opencl_c_fp64=1 -Dcl_khr_int64_base_atomics -Dcl_khr_int64_extended_atomics -Xclang -cl-ext=-all,+__opencl_c_images,+__opencl_c_named_address_space_builtins,+__opencl_c_int64,+__opencl_c_atomic_order_acq_rel,+__opencl_c_atomic_order_seq_cst,+__opencl_c_atomic_scope_device,+__opencl_c_program_scope_global_variables,+__opencl_c_generic_address_space,+cl_khr_byte_addressable_store,+cl_khr_global_int32_base_atomics,+cl_khr_global_int32_extended_atomics,+cl_khr_local_int32_base_atomics,+cl_khr_local_int32_extended_atomics,+cl_khr_3d_image_writes,+__opencl_c_3d_image_writes,+cl_khr_command_buffer,+cl_khr_subgroups,+cl_intel_unified_shared_memory,+cl_khr_spir,+cl_khr_il_program,+cl_khr_fp64,+__opencl_c_fp64,+cl_khr_int64_base_atomics,+cl_khr_int64_extended_atomics -o test_llvm16.bc -emit-llvm -c source.cl

$TRANS --spirv-gen-kernel-arg-name-md --spirv-max-version=1.2 -o test_llvm16.spv test_llvm16.bc
```

results in a SPIR-V which fails to validate:

```
error: line 83: AtomicCompareExchange: expected Pointer to point to a value of type Result Type
%49 = OpAtomicCompareExchange %ulong %48 %uint_0 %uint_16 %uint_0 %44 %exp
```

Like with #2024, Clang-15 & llvm-spirv-15 *without* opaque pointers produce output that validates.

Contributor guide

Open the contributing guide

Research direction

Start by running the supplied OpenCL C reproducer with Clang 16.0.4 and the llvm_release_160 llvm-spirv commands. Inspect the generated test_llvm16.spv around the reported AtomicCompareExchange instruction, then validate it; done means the generated SPIR-V passes validation for this case.

Written by the indexing model from the issue text.

Assessment

Tech stack
c
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
30/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.