Fix mangling for atomic builtins used with SPV_KHR_untyped_pointers (#…

…2771) This change allows to preserve the correct builtin mangling in reverse translation. All the existing tests for atomics (except atomic flag instructions which are not covered by the extension) were updated to verify we get the same mangling with and without extension enabled.
KhronosGroup · Oct 22, 2024 · 5660237 · 5660237
1 parent 9d2926d
commit 5660237
Show file tree

Hide file tree

Showing 18 changed files with 134 additions and 12 deletions.
diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h
@@ -557,6 +557,19 @@ inline unsigned findFirstPtr(const Container &Args) {
   return PtArg - Args.begin();
 }
 
+// Utility function to check if a type is a TypedPointerType
+inline bool isTypedPointerType(llvm::Type *Ty) {
+  return llvm::isa<llvm::TypedPointerType>(Ty);
+}
+
+template <typename Container>
+inline unsigned findFirstPtrType(const Container &Args) {
+  auto PtArg = std::find_if(Args.begin(), Args.end(), [](Type *T) {
+    return T->isPointerTy() || isTypedPointerType(T);
+  });
+  return PtArg - Args.begin();
+}
+
 bool isSupportedTriple(Triple T);
 void removeFnAttr(CallInst *Call, Attribute::AttrKind Attr);
 void addFnAttr(CallInst *Call, Attribute::AttrKind Attr);

diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp
@@ -3369,6 +3369,21 @@ Instruction *SPIRVToLLVM::transBuiltinFromInst(const std::string &FuncName,
   transOCLBuiltinFromInstPreproc(BI, RetTy, Ops);
   std::vector<Type *> ArgTys =
       transTypeVector(SPIRVInstruction::getOperandTypes(Ops), true);
+
+  // Special handling for "truly" untyped pointers to preserve correct
+  // builtin mangling of atomic operations.
+  auto Ptr = findFirstPtrType(ArgTys);
+  if (Ptr < ArgTys.size() &&
+      BI->getValueType(Ops[Ptr]->getId())->isTypeUntypedPointerKHR()) {
+    if (isAtomicOpCodeUntypedPtrSupported(BI->getOpCode())) {
+      auto *AI = static_cast<SPIRVAtomicInstBase *>(BI);
+      ArgTys[Ptr] = TypedPointerType::get(
+          transType(AI->getSemanticType()),
+          SPIRSPIRVAddrSpaceMap::rmap(
+              BI->getValueType(Ops[Ptr]->getId())->getPointerStorageClass()));
+    }
+  }
+
   for (auto &I : ArgTys) {
     if (isa<FunctionType>(I)) {
       I = TypedPointerType::get(I, SPIRAS_Private);

diff --git a/lib/SPIRV/libSPIRV/SPIRVInstruction.h b/lib/SPIRV/libSPIRV/SPIRVInstruction.h
@@ -2953,6 +2953,24 @@ class SPIRVAtomicInstBase : public SPIRVInstTemplateBase {
       assert(this->getModule()->getSPIRVVersion() < VersionNumber::SPIRV_1_4 &&
              "OpAtomicCompareExchangeWeak is removed starting from SPIR-V 1.4");
   }
+
+  // This method is needed for correct translation of atomic instructions when
+  // SPV_KHR_untyped_pointers is enabled.
+  // The interpreted data type for untyped pointers is specified by the Result
+  // Type if it exists, or from the type of the object being stored in other
+  // case.
+  SPIRVType *getSemanticType() {
+    switch (OpCode) {
+    case OpAtomicStore:
+      // Get type of Value operand
+      return getOperand(3)->getType();
+    default: {
+      if (hasType())
+        return getType();
+      return nullptr;
+    }
+    }
+  }
 };
 
 class SPIRVAtomicStoreInst : public SPIRVAtomicInstBase {

diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCode.h b/lib/SPIRV/libSPIRV/SPIRVOpCode.h
@@ -69,6 +69,13 @@ inline bool isAtomicOpCode(Op OpCode) {
          OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear ||
          isFPAtomicOpCode(OpCode);
 }
+inline bool isAtomicOpCodeUntypedPtrSupported(Op OpCode) {
+  static_assert(OpAtomicLoad < OpAtomicXor, "");
+  return ((unsigned)OpCode >= OpAtomicLoad &&
+          (unsigned)OpCode <= OpAtomicXor) ||
+         isFPAtomicOpCode(OpCode);
+}
+
 inline bool isBinaryOpCode(Op OpCode) {
   return ((unsigned)OpCode >= OpIAdd && (unsigned)OpCode <= OpFMod) ||
          OpCode == OpDot || OpCode == OpIAddCarry || OpCode == OpISubBorrow ||

diff --git a/test/AtomicBuiltinsFloat.ll b/test/AtomicBuiltinsFloat.ll
@@ -4,6 +4,10 @@
 ; RUN: llvm-spirv %t.bc -o %t.spv
 ; RUN: spirv-val %t.spv
 
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+; RUN: spirv-val %t.spv
+
 ; CHECK-LABEL: Label
 ; CHECK: Store
 ; CHECK-COUNT-3: AtomicStore

diff --git a/test/AtomicCompareExchange.ll b/test/AtomicCompareExchange.ll
@@ -3,6 +3,10 @@
 ; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 ; RUN: spirv-val %t.spv
 
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: spirv-val %t.spv
+
 ; CHECK-SPIRV: TypeInt [[Int:[0-9]+]] 32 0
 ; CHECK-SPIRV: Constant [[Int]] [[MemScope_CrossDevice:[0-9]+]] 0
 ; CHECK-SPIRV: Constant [[Int]] [[MemSemEqual_SeqCst:[0-9]+]] 16

diff --git a/test/AtomicCompareExchange_cl20.ll b/test/AtomicCompareExchange_cl20.ll
@@ -1,8 +1,12 @@
 ; RUN: llvm-as %s -o %t.bc
-; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s
+; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefixes=CHECK,CHECK-TYPED-PTR
 ; RUN: llvm-spirv %t.bc -o %t.spv
 ; RUN: spirv-val %t.spv
 
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_KHR_untyped_pointers -o - | FileCheck %s --check-prefixes=CHECK,CHECK-UNTYPED-PTR
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+; RUN: spirv-val %t.spv
+
 target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir-unknown-unknown"
 
@@ -14,7 +18,8 @@ target triple = "spir-unknown-unknown"
 ; CHECK: 4 TypeInt [[int:[0-9]+]] 32 0
 ; CHECK: Constant [[int]] [[DeviceScope:[0-9]+]] 1
 ; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 16
-; CHECK: 4 TypePointer [[int_ptr:[0-9]+]] 8 [[int]]
+; CHECK-TYPED-PTR: 4 TypePointer [[int_ptr:[0-9]+]] 8 [[int]]
+; CHECK-UNTYPED-PTR: 3 TypeUntypedPointerKHR [[int_ptr:[0-9]+]] 8
 ; CHECK: 2 TypeBool [[bool:[0-9]+]]
 
 ; Function Attrs: nounwind

diff --git a/test/atomic-load-store.ll b/test/atomic-load-store.ll
@@ -3,6 +3,10 @@
 ; RUN: spirv-val %t.spv
 ; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
 
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
 ; CHECK-DAG: Constant [[#]] [[#CrossDeviceScope:]] 0
 ; CHECK-DAG: Constant [[#]] [[#Release:]] 4
 ; CHECK-DAG: Constant [[#]] [[#SequentiallyConsistent:]] 16
@@ -14,7 +18,7 @@ target triple = "spir64"
 ; Function Attrs: nounwind
 define dso_local spir_func void @test() {
 entry:
-; CHECK: Variable [[#]] [[#PTR:]]
+; CHECK: {{(Variable|UntypedVariableKHR)}} [[#]] [[#PTR:]]
   %0 = alloca i32
 
 ; CHECK: AtomicStore [[#PTR]] [[#CrossDeviceScope]] {{.+}} [[#]]

diff --git a/test/atomicrmw.ll b/test/atomicrmw.ll
@@ -3,6 +3,10 @@
 ; RUN: spirv-val %t.spv
 ; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
 
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_untyped_pointers -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
 ; CHECK: TypeInt [[Int:[0-9]+]] 32 0
 ; CHECK-DAG: Constant [[Int]] [[MemSem_Relaxed:[0-9]+]] 0
 ; CHECK-DAG: Constant [[Int]] [[MemSem_Acquire:[0-9]+]] 2
@@ -11,8 +15,8 @@
 ; CHECK-DAG: Constant [[Int]] [[MemSem_SequentiallyConsistent:[0-9]+]] 16
 ; CHECK-DAG: Constant [[Int]] [[Value:[0-9]+]] 42
 ; CHECK: TypeFloat [[Float:[0-9]+]] 32
-; CHECK: Variable {{[0-9]+}} [[Pointer:[0-9]+]]
-; CHECK: Variable {{[0-9]+}} [[FPPointer:[0-9]+]]
+; CHECK: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[Pointer:[0-9]+]]
+; CHECK: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[FPPointer:[0-9]+]]
 ; CHECK: Constant [[Float]] [[FPValue:[0-9]+]] 1109917696
 
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"

diff --git a/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl b/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
@@ -6,6 +6,13 @@
 // RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
 // RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
 
+// RUN: llvm-spirv %t.bc -spirv-text -o %t.txt --spirv-ext=+SPV_KHR_untyped_pointers
+// RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+// RUN: spirv-val %t.spv
+// RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
+// RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
+
 #define DEFINE_KERNEL(TYPE)                                                            \
 __kernel void testAtomicCompareExchangeExplicit_cl20_##TYPE(                           \
     volatile global atomic_##TYPE* object,                                             \

diff --git a/test/transcoding/AtomicCompareExchange_cl20.ll b/test/transcoding/AtomicCompareExchange_cl20.ll
@@ -4,6 +4,12 @@ target triple = "spir-unknown-unknown"
 
 ; RUN: llvm-as %s -o %t.bc
 ; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
+; RUN: llvm-dis < %t.bc | FileCheck %s
+
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+; RUN: spirv-val %t.spv
 ; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
 ; RUN: llvm-dis < %t.bc | FileCheck %s
 

diff --git a/test/transcoding/OpenCL/atomic_cmpxchg.cl b/test/transcoding/OpenCL/atomic_cmpxchg.cl
@@ -1,7 +1,12 @@
 // RUN: %clang_cc1 %s -triple spir -cl-std=CL1.2 -emit-llvm-bc -fdeclare-opencl-builtins -o %t.bc
 // RUN: llvm-spirv %t.bc -o %t.spv
 // RUN: spirv-val %t.spv
-// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-TYPED-PTRS
+// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
+
+// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+// RUN: spirv-val %t.spv
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-UNTYPED-PTRS
 // RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
 
 // This test checks that the translator is capable to correctly translate
@@ -19,7 +24,8 @@ __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
 
 // CHECK-SPIRV: Name [[TEST:[0-9]+]] "test_atomic_cmpxchg"
 // CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
-// CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
+// CHECK-SPIRV-TYPED-PTRS-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
+// CHECK-SPIRV-UNTYPED-PTRS-DAG: TypeUntypedPointerKHR [[UINT_PTR:[0-9]+]] 5
 //
 // In SPIR-V, atomic_cmpxchg is represented as OpAtomicCompareExchange [2],
 // which also includes memory scope and two memory semantic arguments. The

diff --git a/test/transcoding/OpenCL/atomic_legacy.cl b/test/transcoding/OpenCL/atomic_legacy.cl
@@ -1,7 +1,12 @@
 // RUN: %clang_cc1 %s -triple spir -cl-std=CL1.2 -emit-llvm-bc -fdeclare-opencl-builtins -o %t.bc
 // RUN: llvm-spirv %t.bc -o %t.spv
 // RUN: spirv-val %t.spv
-// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-TYPED-PTRS
+// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
+
+// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+// RUN: spirv-val %t.spv
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-UNTYPED-PTRS
 // RUN: llvm-spirv %t.spv -r --spirv-target-env=CL1.2 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
 
 // This test checks that the translator is capable to correctly translate
@@ -15,7 +20,8 @@ __kernel void test_legacy_atomics(__global int *p, int val) {
 
 // CHECK-SPIRV: Name [[TEST:[0-9]+]] "test_legacy_atomics"
 // CHECK-SPIRV-DAG: TypeInt [[UINT:[0-9]+]] 32 0
-// CHECK-SPIRV-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
+// CHECK-SPIRV-TYPED-PTRS-DAG: TypePointer [[UINT_PTR:[0-9]+]] 5 [[UINT]]
+// CHECK-SPIRV-UNTYPED-PTRS-DAG: TypeUntypedPointerKHR [[UINT_PTR:[0-9]+]] 5
 //
 // In SPIR-V, atomic_add is represented as OpAtomicIAdd [2], which also includes
 // memory scope and memory semantic arguments. The translator applies a default

diff --git a/test/transcoding/OpenCL/atomic_syncscope_test.ll b/test/transcoding/OpenCL/atomic_syncscope_test.ll
@@ -6,6 +6,13 @@
 ; RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o %t.rev.ll
 ; RUN: FileCheck < %t.rev.ll %s -check-prefix=CHECK-LLVM
 
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add,+SPV_KHR_untyped_pointers -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv %t.spv -to-text -o %t.spt
+; RUN: FileCheck < %t.spt %s -check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o %t.rev.ll
+; RUN: FileCheck < %t.rev.ll %s -check-prefix=CHECK-LLVM
+
 target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir64"
 

diff --git a/test/transcoding/atomic_explicit_arguments.cl b/test/transcoding/atomic_explicit_arguments.cl
@@ -5,6 +5,10 @@
 // RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 // RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
 
+// RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv %t.spv -r --spirv-target-env=CL2.0 -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
+
 int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
   return atomic_load_explicit(obj, order, scope);
 }
@@ -33,7 +37,7 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
 
 // CHECK-SPIRV: Function [[int]] [[TRANS_MEM_SCOPE]]
 // CHECK-SPIRV: FunctionParameter [[int]] [[KEY:[0-9]+]]
-// CHECK-SPIRV: Variable {{[0-9]+}} [[RES:[0-9]+]]
+// CHECK-SPIRV: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[RES:[0-9]+]]
 // CHECK-SPIRV: Switch [[KEY]] [[CASE_2:[0-9]+]] 0 [[CASE_0:[0-9]+]] 1 [[CASE_1:[0-9]+]] 2 [[CASE_2]] 3 [[CASE_3:[0-9]+]] 4 [[CASE_4:[0-9]+]]
 // CHECK-SPIRV: Label [[CASE_0]]
 // CHECK-SPIRV: Store [[RES]] [[FOUR]]
@@ -57,7 +61,7 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
 
 // CHECK-SPIRV: Function [[int]] [[TRANS_MEM_ORDER]]
 // CHECK-SPIRV: FunctionParameter [[int]] [[KEY:[0-9]+]]
-// CHECK-SPIRV: Variable {{[0-9]+}} [[RES:[0-9]+]]
+// CHECK-SPIRV: {{(Variable|UntypedVariableKHR)}} {{[0-9]+}} [[RES:[0-9]+]]
 // CHECK-SPIRV: Switch [[KEY]] [[CASE_5:[0-9]+]] 0 [[CASE_0:[0-9]+]] 2 [[CASE_2:[0-9]+]] 3 [[CASE_3:[0-9]+]] 4 [[CASE_4:[0-9]+]] 5 [[CASE_5]]
 // CHECK-SPIRV: Label [[CASE_0]]
 // CHECK-SPIRV: Store [[RES]] [[ZERO]]

diff --git a/test/transcoding/atomic_flag.cl b/test/transcoding/atomic_flag.cl
@@ -4,7 +4,6 @@
 // RUN: spirv-val %t.spv
 // RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.rev.bc
 // RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
-
 kernel void testAtomicFlag(global int *res) {
   atomic_flag f;
 

diff --git a/test/transcoding/atomic_load_store.ll b/test/transcoding/atomic_load_store.ll
@@ -10,6 +10,13 @@ target triple = "spir-unknown-unknown"
 ; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
 ; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
 
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_untyped_pointers -spirv-text -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_untyped_pointers -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -r --spirv-target-env=CL2.0 %t.spv -o %t.bc
+; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
+
 ; Check 'LLVM ==> SPIR-V ==> LLVM' conversion of atomic_load and atomic_store.
 
 

diff --git a/test/transcoding/atomics_1.2.ll b/test/transcoding/atomics_1.2.ll
@@ -7,6 +7,12 @@ target triple = "spir64-unknown-unknown"
 ; RUN: llvm-spirv -r %t.spv -o %t.bc
 ; RUN: llvm-dis < %t.bc | FileCheck %s
 
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_untyped_pointers
+; TODO: investigate why function parameters are decorated with Volatile multiple times.
+; R/UN: spirv-val %t.spv
+; RUN: llvm-spirv -r %t.spv -o %t.bc
+; RUN: llvm-dis < %t.bc | FileCheck %s
+
 ; Most of atomics lost information about the sign of the integer operand
 ; but since this concerns only built-ins  with two-complement's arithmetics
 ; it shouldn't cause any problems.