diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp
index 2340d98fb..4fc8b132d 100644
--- a/lib/SPIRV/OCLToSPIRV.cpp
+++ b/lib/SPIRV/OCLToSPIRV.cpp
@@ -42,15 +42,21 @@
 #include "SPIRVInternal.h"
 #include "libSPIRV/SPIRVDebug.h"
 
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/Analysis/ValueTracking.h"
+#include "llvm/IR/Constants.h"
 #include "llvm/IR/IRBuilder.h"
 #include "llvm/IR/Instruction.h"
 #include "llvm/IR/Instructions.h"
+#include "llvm/IR/Operator.h"
 #include "llvm/IR/PatternMatch.h"
+#include "llvm/IR/TypedPointerType.h"
 #include "llvm/Support/Debug.h"
 
 #include <algorithm>
+#include <optional>
 #include <regex>
 #include <set>
 
@@ -62,6 +68,96 @@ using namespace SPIRV;
 using namespace OCLUtil;
 
 namespace SPIRV {
+
+static std::optional<unsigned> getAddressSpaceFromType(Type *Ty) {
+  assert(Ty && "Can't deduce pointer AS");
+  if (auto *TypedPtr = dyn_cast<TypedPointerType>(Ty))
+    return TypedPtr->getAddressSpace();
+  if (auto *Ptr = dyn_cast<PointerType>(Ty))
+    return Ptr->getAddressSpace();
+  return std::nullopt;
+}
+
+// Performs an address space inference analysis.
+static std::optional<unsigned> getAddressSpaceFromValue(Value *Ptr) {
+  assert(Ptr && "Can't deduce pointer AS");
+
+  SmallPtrSet<Value *, 8> Visited;
+  SmallVector<Value *, 8> Worklist;
+  Worklist.push_back(Ptr);
+  std::optional<unsigned> GenericAS;
+
+  while (!Worklist.empty()) {
+    Value *Current = Worklist.pop_back_val();
+    if (!Visited.insert(Current).second)
+      continue;
+
+    if (auto MaybeAS = getAddressSpaceFromType(Current->getType())) {
+      if (*MaybeAS != SPIRAS_Generic)
+        return MaybeAS;
+      GenericAS = MaybeAS;
+    }
+
+    if (isa<AllocaInst>(Current))
+      // It's safe to assume, that generic alloca is actually a function
+      // storage allocation.
+      return SPIRAS_Private;
+
+    // Find origins of the pointer and add to the worklist.
+    if (auto *Op = dyn_cast<Operator>(Current)) {
+      switch (Op->getOpcode()) {
+      case Instruction::AddrSpaceCast:
+      case Instruction::BitCast:
+      case Instruction::GetElementPtr:
+        Worklist.push_back(Op->getOperand(0));
+        break;
+      case Instruction::Select:
+        Worklist.push_back(Op->getOperand(1));
+        Worklist.push_back(Op->getOperand(2));
+        break;
+      case Instruction::PHI: {
+        auto *Phi = cast<PHINode>(Op);
+        for (Value *Incoming : Phi->incoming_values())
+          Worklist.push_back(Incoming);
+        break;
+      }
+      default:
+        break;
+      }
+    }
+  }
+
+  return GenericAS;
+}
+
+// Sets memory semantic mask of an atomic depending on a pointer argument
+// address space.
+static unsigned getAtomicPointerMemorySemanticsMemoryMask(Value *Ptr,
+                                                          Type *RecordedType) {
+  std::optional<unsigned> AddrSpace = getAddressSpaceFromType(RecordedType);
+  if ((!AddrSpace || *AddrSpace == SPIRAS_Generic) && Ptr)
+    if (auto MaybeAS = getAddressSpaceFromValue(Ptr))
+      AddrSpace = MaybeAS;
+
+  if (!AddrSpace)
+    return MemorySemanticsMaskNone;
+
+  switch (*AddrSpace) {
+  case SPIRAS_Global:
+  case SPIRAS_GlobalDevice:
+  case SPIRAS_GlobalHost:
+    return MemorySemanticsCrossWorkgroupMemoryMask;
+  case SPIRAS_Local:
+    return MemorySemanticsWorkgroupMemoryMask;
+  case SPIRAS_Generic:
+    return MemorySemanticsCrossWorkgroupMemoryMask |
+           MemorySemanticsWorkgroupMemoryMask |
+           MemorySemanticsSubgroupMemoryMask;
+  default:
+    return MemorySemanticsMaskNone;
+  }
+}
+
 static size_t getOCLCpp11AtomicMaxNumOps(StringRef Name) {
   return StringSwitch<size_t>(Name)
       .Cases({"load", "flag_test_and_set", "flag_clear"}, 3)
@@ -700,6 +796,11 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
   const size_t ScopeIdx = ArgsCount - 1;
   const size_t OrderIdx = ScopeIdx - NumOrder;
 
+  unsigned PtrMemSemantics = MemorySemanticsMaskNone;
+  if (Mutator.arg_size() > 0)
+    PtrMemSemantics = getAtomicPointerMemorySemanticsMemoryMask(
+        Mutator.getArg(0), Mutator.getType(0));
+
   if (NeedsNegate) {
     Mutator.mapArg(1, [=](Value *V) {
       IRBuilder<> IRB(CI);
@@ -710,9 +811,20 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
     return transOCLMemScopeIntoSPIRVScope(V, OCLMS_device, CI);
   });
   for (size_t I = 0; I < NumOrder; ++I) {
-    Mutator.mapArg(OrderIdx + I, [=](Value *V) {
-      return transOCLMemOrderIntoSPIRVMemorySemantics(V, OCLMO_seq_cst, CI);
-    });
+    Mutator.mapArg(
+        OrderIdx + I, [=](IRBuilder<> &Builder, Value *V) -> Value * {
+          Value *MemSem =
+              transOCLMemOrderIntoSPIRVMemorySemantics(V, OCLMO_seq_cst, CI);
+          if (PtrMemSemantics == MemorySemanticsMaskNone)
+            return MemSem;
+
+          auto *MemSemTy = cast<IntegerType>(MemSem->getType());
+          auto *Mask = ConstantInt::get(MemSemTy, PtrMemSemantics);
+          if (auto *Const = dyn_cast<ConstantInt>(MemSem))
+            return static_cast<Value *>(ConstantInt::get(
+                MemSemTy, Const->getZExtValue() | PtrMemSemantics));
+          return Builder.CreateOr(MemSem, Mask);
+        });
   }
 
   // Order of args in SPIR-V:
diff --git a/test/AtomicCompareExchangeExplicit.ll b/test/AtomicCompareExchangeExplicit.ll
index 2cfac5a58..dc728a09b 100644
--- a/test/AtomicCompareExchangeExplicit.ll
+++ b/test/AtomicCompareExchangeExplicit.ll
@@ -8,7 +8,9 @@ target triple = "spir64"
 
 ; CHECK-DAG: 4 TypeInt [[#int:]] 32 0
 ; CHECK-DAG: Constant [[#int]] [[#DeviceScope:]] 4
-; CHECK-DAG: Constant [[#int]] [[#SequentiallyConsistent_MS:]] 0
+; Memory semantics: 256 = WorkgroupMemory (256) | SequentiallyConsistent (0)
+; Local address space (3) maps to WorkgroupMemory storage class
+; CHECK-DAG: Constant [[#int]] [[#SequentiallyConsistent_MS:]] 256
 ; CHECK-DAG: 4 TypePointer [[#int_ptr:]] 4 [[#int]]
 ; CHECK-DAG: 2 TypeBool [[#bool:]]
 
diff --git a/test/AtomicCompareExchange_cl20.ll b/test/AtomicCompareExchange_cl20.ll
index be6023c3d..c1c539069 100644
--- a/test/AtomicCompareExchange_cl20.ll
+++ b/test/AtomicCompareExchange_cl20.ll
@@ -17,7 +17,9 @@ 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
+; For generic AS with SequentiallyConsistent: 912 = 896 (storage class) + 16 (SeqCst)
+; Where 896 = CrossWorkgroupMemory (512) | WorkgroupMemory (256) | SubgroupMemory (128)
+; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 912
 ; 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]+]]
diff --git a/test/extensions/EXT/SPV_EXT_shader_atomic_float_/AtomicFSubEXTForOCL.ll b/test/extensions/EXT/SPV_EXT_shader_atomic_float_/AtomicFSubEXTForOCL.ll
index 87c05b8a8..e4e0f8de7 100644
--- a/test/extensions/EXT/SPV_EXT_shader_atomic_float_/AtomicFSubEXTForOCL.ll
+++ b/test/extensions/EXT/SPV_EXT_shader_atomic_float_/AtomicFSubEXTForOCL.ll
@@ -29,7 +29,8 @@ define dso_local spir_func void @test_atomic_float(ptr addrspace(1) %a) local_un
 entry:
   ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] 13 7 10 11 [[NEGATIVE_229]]
   ; CHECK-LLVM-CL20: call spir_func float @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(ptr addrspace(4) %a.as, float -2.290000e+02, i32 0, i32 1) #0
-  ; CHECK-LLVM-SPV: call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fiif(ptr addrspace(1) %a, i32 2, i32 0, float -2.290000e+02) #0
+  ; Memory semantics: 512 = CrossWorkgroupMemory (512) | Relaxed (0) for global AS
+  ; CHECK-LLVM-SPV: call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fiif(ptr addrspace(1) %a, i32 2, i32 512, float -2.290000e+02) #0
   %call2 = tail call spir_func float @_Z25atomic_fetch_sub_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(ptr addrspace(1) noundef %a, float noundef 2.290000e+02, i32 noundef 0, i32 noundef 1) #2
   ret void
 }
@@ -43,7 +44,8 @@ define dso_local spir_func void @test_atomic_double(ptr addrspace(1) %a) local_u
 entry:
   ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] 21 18 10 11 [[NEGATIVE_334]]
   ; CHECK-LLVM-CL20: call spir_func double @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicdd12memory_order12memory_scope(ptr addrspace(4) %a.as, double -3.340000e+02, i32 0, i32 1) #0
-  ; CHECK-LLVM-SPV: call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1diid(ptr addrspace(1) %a, i32 2, i32 0, double -3.340000e+02) #0
+  ; Memory semantics: 512 = CrossWorkgroupMemory (512) | Relaxed (0) for global AS
+  ; CHECK-LLVM-SPV: call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1diid(ptr addrspace(1) %a, i32 2, i32 512, double -3.340000e+02) #0
   %call = tail call spir_func double @_Z25atomic_fetch_sub_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(ptr addrspace(1) noundef %a, double noundef 3.340000e+02, i32 noundef 0, i32 noundef 1) #2
   ret void
 }
diff --git a/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl b/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
index 45111b8cd..7f2119a64 100644
--- a/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
+++ b/test/transcoding/AtomicCompareExchangeExplicit_cl20.cl
@@ -49,9 +49,13 @@ DEFINE_KERNEL(double)
 //; Constants below correspond to the SPIR-V spec
 //CHECK-SPIRV-DAG: Constant [[int32]] [[DeviceScope:[0-9]+]] 1
 //CHECK-SPIRV-DAG: Constant [[int32]] [[WorkgroupScope:[0-9]+]] 2
-//CHECK-SPIRV-DAG: Constant [[int32]] [[ReleaseMemSem:[0-9]+]] 4
-//CHECK-SPIRV-DAG: Constant [[int32]] [[RelaxedMemSem:[0-9]+]] 0
-//CHECK-SPIRV-DAG: Constant [[int32]] [[AcqRelMemSem:[0-9]+]] 8
+//; Memory semantics include both memory order and storage class bits
+//; 516 = CrossWorkgroupMemory (512) | Release (4)
+//; 512 = CrossWorkgroupMemory (512) | Relaxed (0)
+//; 520 = CrossWorkgroupMemory (512) | AcqRel (8)
+//CHECK-SPIRV-DAG: Constant [[int32]] [[ReleaseMemSem:[0-9]+]] 516
+//CHECK-SPIRV-DAG: Constant [[int32]] [[RelaxedMemSem:[0-9]+]] 512
+//CHECK-SPIRV-DAG: Constant [[int32]] [[AcqRelMemSem:[0-9]+]] 520
 
 //CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[DeviceScope]] [[ReleaseMemSem]] [[RelaxedMemSem]]
 //CHECK-SPIRV: AtomicCompareExchange [[int32]] {{[0-9]+}} {{[0-9]+}} [[WorkgroupScope]] [[AcqRelMemSem]] [[RelaxedMemSem]]
diff --git a/test/transcoding/OpenCL/atomic_cmpxchg.cl b/test/transcoding/OpenCL/atomic_cmpxchg.cl
index aeffc8367..37dd05600 100644
--- a/test/transcoding/OpenCL/atomic_cmpxchg.cl
+++ b/test/transcoding/OpenCL/atomic_cmpxchg.cl
@@ -35,9 +35,8 @@ __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
 // 0x2 Workgroup
 // CHECK-SPIRV-DAG: Constant [[UINT]] [[WORKGROUP_SCOPE:[0-9]+]] 2
 //
-// 0x0 Relaxed
-// TODO: do we need CrossWorkgroupMemory here as well?
-// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 0
+// 0x0 Relaxed | 0x200 CrossWorkgroupMemory
+// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 512
 //
 // CHECK-SPIRV: Function {{[0-9]+}} [[TEST]]
 // CHECK-SPIRV: FunctionParameter [[UINT_PTR]] [[PTR:[0-9]+]]
diff --git a/test/transcoding/OpenCL/atomic_legacy.cl b/test/transcoding/OpenCL/atomic_legacy.cl
index 31a2cc3ef..47febef61 100644
--- a/test/transcoding/OpenCL/atomic_legacy.cl
+++ b/test/transcoding/OpenCL/atomic_legacy.cl
@@ -31,8 +31,9 @@ __kernel void test_legacy_atomics(__global int *p, int val) {
 // 0x2 Workgroup
 // CHECK-SPIRV-DAG: Constant [[UINT]] [[WORKGROUP_SCOPE:[0-9]+]] 2
 //
-// 0x0 Relaxed
-// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 0
+// 0x200 CrossWorkgroupMemory | 0x0 Relaxed = 512
+// Global address space (AS 1) maps to CrossWorkgroupMemory storage class
+// CHECK-SPIRV-DAG: Constant [[UINT]] [[RELAXED:[0-9]+]] 512
 //
 // CHECK-SPIRV: Function {{[0-9]+}} [[TEST]]
 // CHECK-SPIRV: FunctionParameter [[UINT_PTR]] [[PTR:[0-9]+]]
diff --git a/test/transcoding/OpenCL/atomic_syncscope_test.ll b/test/transcoding/OpenCL/atomic_syncscope_test.ll
index 46ebd9660..8248a9b05 100644
--- a/test/transcoding/OpenCL/atomic_syncscope_test.ll
+++ b/test/transcoding/OpenCL/atomic_syncscope_test.ll
@@ -33,19 +33,23 @@ target triple = "spir64"
 ; 4 - sub_group
 
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt0:]] 0
-; CHECK-SPIRV-DAG: Constant [[#]] [[#SequentiallyConsistent:]] 16
+; CHECK-SPIRV-DAG: Constant [[#]] [[#SCPrivate:]] 16
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt1:]] 1
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt2:]] 2
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt3:]] 3
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt4:]] 4
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#Const2Power30:]] 1073741824
 ; CHECK-SPIRV-DAG: Constant [[#]] [[#ConstInt42:]] 42
+; Note: Storage class bits (SCGlobal, SCLocal, etc.) are not added for plain LLVM IR atomics
+; Only OpenCL builtin atomics get the storage class memory semantics bits from the patch
 
 ; AtomicLoad ResTypeId ResId PtrId MemScopeId MemSemanticsId
-; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]]
-; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt1]] [[#SequentiallyConsistent]]
-; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SequentiallyConsistent]]
-; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt3]] [[#SequentiallyConsistent]]
+; Note: Plain LLVM atomic loads don't get storage class bits added (only OpenCL builtins do)
+; These use SCPrivate (16) which is SequentiallyConsistent without storage class bits
+; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]]
+; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt1]] [[#SCPrivate]]
+; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SCPrivate]]
+; CHECK-SPIRV: AtomicLoad [[#]] [[#]] [[#]] [[#ConstInt3]] [[#SCPrivate]]
 
 ; CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 1)
 ; CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 2)
@@ -62,8 +66,9 @@ entry:
 }
 
 ; AtomicStore PtrId MemScopeId MemSemanticsId ValueId
-; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt3]] [[#SequentiallyConsistent]] [[#ConstInt1]]
-; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#ConstInt1]]
+; Plain LLVM IR store atomic instructions don't get storage class bits
+; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt3]] [[#SCPrivate]] [[#ConstInt1]]
+; CHECK-SPIRV: AtomicStore [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#ConstInt1]]
 ; CHECK-LLVM: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 4)
 ; CHECK-LLVM: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 1)
 
@@ -75,11 +80,11 @@ entry:
 }
 
 ; Atomic* ResTypeId ResId PtrId MemScopeId MemSemanticsId ValueId
-; CHECK-SPIRV: AtomicAnd [[#]] [[#]] [[#]] [[#ConstInt4]] [[#SequentiallyConsistent]] [[#ConstInt1]]
-; CHECK-SPIRV: AtomicSMin [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SequentiallyConsistent]] [[#ConstInt1]]
-; CHECK-SPIRV: AtomicSMax [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SequentiallyConsistent]] [[#ConstInt1]]
-; CHECK-SPIRV: AtomicUMin [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#ConstInt1]]
-; CHECK-SPIRV: AtomicUMax [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#ConstInt1]]
+; CHECK-SPIRV: AtomicAnd [[#]] [[#]] [[#]] [[#ConstInt4]] [[#SCPrivate]] [[#ConstInt1]]
+; CHECK-SPIRV: AtomicSMin [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SCPrivate]] [[#ConstInt1]]
+; CHECK-SPIRV: AtomicSMax [[#]] [[#]] [[#]] [[#ConstInt0]] [[#SCPrivate]] [[#ConstInt1]]
+; CHECK-SPIRV: AtomicUMin [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#ConstInt1]]
+; CHECK-SPIRV: AtomicUMax [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#ConstInt1]]
 
 ; CHECK-LLVM: call spir_func i32 @_Z25atomic_fetch_and_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1, i32 5, i32 0)
 ; CHECK-LLVM: call spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1, i32 5, i32 3)
@@ -109,7 +114,7 @@ entry:
 }
 
 ; AtomicExchange ResTypeId ResId PtrId MemScopeId MemSemanticsId ValueId
-; CHECK-SPIRV: AtomicExchange [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#Const2Power30]]
+; CHECK-SPIRV: AtomicExchange [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SCPrivate]] [[#Const2Power30]]
 ; CHECK-LLVM: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1073741824, i32 5, i32 1)
 
 define dso_local float @ff3(ptr captures(none) noundef %d) local_unnamed_addr #0 {
@@ -120,6 +125,7 @@ entry:
 }
 
 ; AtomicFAddEXT ResTypeId ResId PtrId MemScopeId MemSemanticsId ValueId
+; Plain LLVM atomicrmw fadd doesn't get storage class bits
 ; CHECK-SPIRV: AtomicFAddEXT [[#]] [[#]] [[#]] [[#ConstInt2]] [[#ConstInt0]] [[#]]
 ; CHECK-LLVM: call spir_func float @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(ptr{{.*}}, i32 0, i32 1)
 
diff --git a/test/transcoding/atomic_explicit_arguments.cl b/test/transcoding/atomic_explicit_arguments.cl
index 21b16c16b..56b90d28a 100644
--- a/test/transcoding/atomic_explicit_arguments.cl
+++ b/test/transcoding/atomic_explicit_arguments.cl
@@ -18,6 +18,8 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
 // CHECK-SPIRV: Name [[TRANS_MEM_ORDER:[0-9]+]] "__translate_ocl_memory_order"
 
 // CHECK-SPIRV: TypeInt [[int:[0-9]+]] 32 0
+// Memory semantics for generic AS: 896 = CrossWorkgroupMemory | WorkgroupMemory | SubgroupMemory = 512+256+128
+// CHECK-SPIRV-DAG: Constant [[int]] [[GENERIC_STORAGE_MASK:[0-9]+]] 896
 // CHECK-SPIRV-DAG: Constant [[int]] [[ZERO:[0-9]+]] 0
 // CHECK-SPIRV-DAG: Constant [[int]] [[ONE:[0-9]+]] 1
 // CHECK-SPIRV-DAG: Constant [[int]] [[TWO:[0-9]+]] 2
@@ -31,8 +33,11 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
 // CHECK-SPIRV: FunctionParameter {{[0-9]+}} [[OCL_ORDER:[0-9]+]]
 // CHECK-SPIRV: FunctionParameter {{[0-9]+}} [[OCL_SCOPE:[0-9]+]]
 
+//
 // CHECK-SPIRV: FunctionCall [[int]] [[SPIRV_SCOPE:[0-9]+]] [[TRANS_MEM_SCOPE]] [[OCL_SCOPE]]
-// CHECK-SPIRV: FunctionCall [[int]] [[SPIRV_ORDER:[0-9]+]] [[TRANS_MEM_ORDER]] [[OCL_ORDER]]
+// CHECK-SPIRV: FunctionCall [[int]] [[SPIRV_ORDER_BASE:[0-9]+]] [[TRANS_MEM_ORDER]] [[OCL_ORDER]]
+// The translated memory order is combined with storage class semantics for generic AS
+// CHECK-SPIRV: BitwiseOr [[int]] [[SPIRV_ORDER:[0-9]+]] [[SPIRV_ORDER_BASE]] [[GENERIC_STORAGE_MASK]]
 // CHECK-SPIRV: AtomicLoad [[int]] {{[0-9]+}} [[OBJECT]] [[SPIRV_SCOPE]] [[SPIRV_ORDER]]
 
 // CHECK-SPIRV: Function [[int]] [[TRANS_MEM_SCOPE]]
@@ -86,5 +91,7 @@ int load (volatile atomic_int* obj, memory_order order, memory_scope scope) {
 
 // CHECK-LLVM: define spir_func i32 @load(ptr addrspace(4) %[[obj:[0-9a-zA-Z._]+]], i32 %[[order:[0-9a-zA-Z._]+]], i32 %[[scope:[0-9a-zA-Z._]+]]) #0 {
 // CHECK-LLVM: entry:
-// CHECK-LLVM:  call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr addrspace(4) %[[obj]], i32 %[[order]], i32 %[[scope]])
+// With the patch, memory order is ORed with storage class bits (896) and translated back
+// CHECK-LLVM:  %[[#]] = or i32 %{{[0-9a-zA-Z._]+}}, 896
+// CHECK-LLVM:  call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr addrspace(4) %[[obj]], i32 %{{[0-9a-zA-Z._]+}}, i32 %[[scope]])
 // CHECK-LLVM: }