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

[franz]

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.