KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator
SPIR-V fails to validate with: AtomicCompareExchange: expected Pointer to point to a value of type Result Type
- 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
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