Fix wrong fence scope generated for group barrier

YuriPlyakhin · igcbot · commit adb0b1790f0f · 2024-08-03T00:36:32.000+02:00
- Modify GenISA_memoryfence intrinsic to accept memory scope
  parameter and take it into account, when generating LSC fence
  (previously GPU scope was always used for global fence and
  GROUP scope was always used for local fence).
- Modify other passes and built-ins which are using this
  intrinsic to pass through scope parameter, where it is
  available and makes sense.
diff --git a/IGC/AdaptorCommon/DivergentBarrierPass.cpp b/IGC/AdaptorCommon/DivergentBarrierPass.cpp
@@ -56,6 +56,14 @@ void DivergentBarrierPass::updateFenceArgs(
         else
             return true;
     };
+    auto updateScope = [](const Value* V)
+    {
+        // If we don't know at compile-time, conservatively assume LSC_SCOPE_GPU
+        if (auto* C = dyn_cast<ConstantInt>(V))
+            return static_cast<uint32_t>(C->getValue().getZExtValue());
+        else
+            return static_cast<uint32_t>(LSC_SCOPE_GPU);
+    };
 
     // Always CommitEnable
 
@@ -66,6 +74,7 @@ void DivergentBarrierPass::updateFenceArgs(
     Args.Global                 |= update(I->getOperand(5));
     Args.L1_Invalidate          |= update(I->getOperand(6));
     Args.L1_Evict               |= update(I->getOperand(7));
+    Args.Scope                   = updateScope(I->getOperand(8));
 }
 
 CallInst* DivergentBarrierPass::insertFence(
@@ -83,6 +92,7 @@ CallInst* DivergentBarrierPass::insertFence(
         IRB.getInt1(FA.Global),
         IRB.getInt1(FA.L1_Invalidate),
         IRB.getInt1(FA.L1_Evict),
+        IRB.getInt32(FA.Scope)
     };
     return IRB.CreateCall(FenceFn, Args);
 }
diff --git a/IGC/AdaptorCommon/DivergentBarrierPass.h b/IGC/AdaptorCommon/DivergentBarrierPass.h
@@ -49,6 +49,9 @@ namespace IGC
             bool Global                 = false;
             bool L1_Invalidate          = false;
             bool L1_Evict               = false;
+            // init Scope with GROUP which should be used for SLM barriers
+            // to be consistent with Global argument initialized to false
+            uint Scope                  = LSC_SCOPE_GROUP;
         };
 
         typedef llvm::DenseMap<uint64_t, WIAnalysis::WIDependancy> SlotDepMap;
diff --git a/IGC/BiFModule/Implementation/IGCBiF_Intrinsics.cl b/IGC/BiFModule/Implementation/IGCBiF_Intrinsics.cl
@@ -113,8 +113,7 @@ void     __builtin_IB_write_2d_f(int, int2, float4, int);
 local uchar* __builtin_IB_AllocLocalMemPool(bool allocAllWorkgroups, uint numAdditionalElements, uint elementSize);
 
 // Memory fences
-// See GenISAIntrinsics.td for documentation
-void     __builtin_IB_memfence(bool commitEnable, bool flushRW, bool flushConstant, bool flushTexture, bool flushIcache, bool isGlobal, bool invalidateL1, bool evictL1);
+void     __builtin_IB_memfence(bool commitEnable, bool flushRW, bool flushConstant, bool flushTexture, bool flushIcache, bool isGlobal, bool invalidateL1, bool evictL1, uint scope);
 void     __builtin_IB_flush_sampler_cache(void);
 void     __builtin_IB_typedmemfence(bool invalidateCache);
 
diff --git a/IGC/BiFModule/Implementation/IGCBiF_Intrinsics_Lsc.cl b/IGC/BiFModule/Implementation/IGCBiF_Intrinsics_Lsc.cl
@@ -263,29 +263,29 @@ bool __builtin_IB_lsc_load_status_global_ulong8(const __global ulong8 *base, int
 ///////////////////////////////////////////////////////////////////////
 
 // FS - Fence Scope
-enum LSC_FS {
+typedef enum {
     LSC_FS_THREAD_GROUP,
     LSC_FS_LOCAL,
     LSC_FS_TILE,
     LSC_FS_GPU,
     LSC_FS_GPUs,
     LSC_FS_SYSTEM_RELEASE,
     LSC_FS_SYSTEM_ACQUIRE
-};
+} LSC_FS;
 
 // FT - Fence Type
-enum LSC_FT {
+typedef enum {
     LSC_FT_DEFAULT,
     LSC_FT_EVICT,
     LSC_FT_INVALIDATE,
     LSC_FT_DISCARD,
     LSC_FT_CLEAN,
     LSC_FT_L3
-};
+} LSC_FT;
 
-void  __builtin_IB_lsc_fence_global_untyped(enum LSC_FS scope, enum LSC_FT flushType);   // Mem Port - UGM
-void  __builtin_IB_lsc_fence_global_untyped_cross_tile(enum LSC_FS scope, enum LSC_FT flushType);  // Mem Port - UGML
-void  __builtin_IB_lsc_fence_global_typed(enum LSC_FS scope, enum LSC_FT flushType);     // Mem Port - TGM
+void  __builtin_IB_lsc_fence_global_untyped(LSC_FS scope, LSC_FT flushType);             // Mem Port - UGM
+void  __builtin_IB_lsc_fence_global_untyped_cross_tile(LSC_FS scope, LSC_FT flushType);  // Mem Port - UGML
+void  __builtin_IB_lsc_fence_global_typed(LSC_FS scope, LSC_FT flushType);               // Mem Port - TGM
 void  __builtin_IB_lsc_fence_local();                                                    // Mem Port - SLM
 void  __builtin_IB_lsc_fence_evict_to_memory();                                          // Mem Port - UGM
 
diff --git a/IGC/BiFModule/Implementation/atomics.cl b/IGC/BiFModule/Implementation/atomics.cl
@@ -20,7 +20,7 @@ SPDX-License-Identifier: MIT
   __local int* __builtin_IB_get_local_lock();
   __global int* __builtin_IB_get_global_lock();
   void __builtin_IB_eu_thread_pause(uint value);
-  void __intel_memfence_handler(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1);
+  void __intel_memfence_handler(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1, Scope_t scope);
 
 #define LOCAL_SPINLOCK_START() \
   { \
@@ -51,14 +51,14 @@ SPDX-License-Identifier: MIT
   if( ( (Semantics) & ( SEMANTICS_PRE_OP_NEED_FENCE ) ) > 0 )                         \
   {                                                                                   \
       bool flushL3 = (isGlobal) && ((Scope) == Device || (Scope) == CrossDevice);     \
-      __intel_memfence_handler(flushL3, isGlobal, false, isGlobal);                   \
+      __intel_memfence_handler(flushL3, isGlobal, false, isGlobal, Scope);            \
   }
 
 #define FENCE_POST_OP(Scope, Semantics, isGlobal)                                     \
   if( ( (Semantics) & ( SEMANTICS_POST_OP_NEEDS_FENCE ) ) > 0 )                       \
   {                                                                                   \
       bool flushL3 = (isGlobal) && ((Scope) == Device || (Scope) == CrossDevice);     \
-      __intel_memfence_handler(flushL3, isGlobal, isGlobal, false);                   \
+      __intel_memfence_handler(flushL3, isGlobal, isGlobal, false, Scope);            \
   }
 
 // This fencing scheme allows us to obey the memory model when coherency is
diff --git a/IGC/BiFModule/Implementation/barrier.cl b/IGC/BiFModule/Implementation/barrier.cl
@@ -16,39 +16,62 @@ SPDX-License-Identifier: MIT
 
 // MEMFENCE IMPLEMENTATION
 
-void OPTNONE __intel_memfence_optnone(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1)
+// Use macro to make conversion of scope parameter a compile-time constant for O0 path
+// Below choice of some HW scopes is explained.
+// 1. CrossDevice -> LSC_FS_SYSTEM_RELEASE:
+//     sycl::memory_scope::system is lowered to SPIR-V CrossDevice. SYCL SPEC 2020 for system scope:
+//     "applies to any work-item or host thread in the system that is currently permitted to access
+//     the memory allocation containing the referenced object".
+//     Hence, conservatively mapping to “system”.
+// 2. Subgroup and Invocation --> LSC_FS_THREAD_GROUP:
+//     our HW spec doesn’t have corresponding scope for Invocation or Subgroup, hence mapping to
+//     lowest possible scope
+#define CONVERT_SCOPE_SPIRV_TO_VISA(scope)            \
+    ((scope) == CrossDevice ? LSC_FS_SYSTEM_RELEASE : \
+     (scope) == Device      ? LSC_FS_GPU :            \
+     LSC_FS_THREAD_GROUP)
+
+void OPTNONE __intel_memfence_optnone(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1, Scope_t scope)
 {
-#define MEMFENCE_IF(V1, V5, V6, V7)                                           \
-if (flushRW == V1 && isGlobal == V5 && invalidateL1 == V6 && evictL1 == V7)   \
-{                                                                             \
-    __builtin_IB_memfence(true, V1, false, false, false, V5, V6, V7);         \
+#define MEMFENCE_IF(V1, V5, V6, V7, V8)                                                      \
+if (flushRW == V1 && isGlobal == V5 && invalidateL1 == V6 && evictL1 == V7 && scope == V8)   \
+{                                                                                            \
+    LSC_FS lsc_scope = CONVERT_SCOPE_SPIRV_TO_VISA(V8);                                      \
+    __builtin_IB_memfence(true, V1, false, false, false, V5, V6, V7, lsc_scope);             \
 } else
 
 // Generate combinations for all MEMFENCE_IF cases, e.g.:
 // true, true, true
 // true, true, false etc.
+#define MF_L4(...) MF_L3(__VA_ARGS__,false) MF_L3(__VA_ARGS__,true)
 #define MF_L3(...) MF_L2(__VA_ARGS__,false) MF_L2(__VA_ARGS__,true)
 #define MF_L2(...) MF_L1(__VA_ARGS__,false) MF_L1(__VA_ARGS__,true)
-#define MF_L1(...) MEMFENCE_IF(__VA_ARGS__,false) MEMFENCE_IF(__VA_ARGS__,true)
-MF_L3(false)
-MF_L3(true) {}
+#define MF_L1(...) MEMFENCE_IF(__VA_ARGS__,CrossDevice) \
+                   MEMFENCE_IF(__VA_ARGS__,Device)      \
+                   MEMFENCE_IF(__VA_ARGS__,Workgroup)   \
+                   MEMFENCE_IF(__VA_ARGS__,Subgroup)    \
+                   MEMFENCE_IF(__VA_ARGS__,Invocation)
+MF_L4(false)
+MF_L4(true) {}
 
 #undef MEMFENCE_IF
+#undef MF_L4
 #undef MF_L3
 #undef MF_L2
 #undef MF_L1
 }
-void __intel_memfence(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1)
+void __intel_memfence(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1, Scope_t scope)
 {
-    __builtin_IB_memfence(true, flushRW, false, false, false, isGlobal, invalidateL1, evictL1);
+    LSC_FS lsc_scope = CONVERT_SCOPE_SPIRV_TO_VISA(scope);
+    __builtin_IB_memfence(true, flushRW, false, false, false, isGlobal, invalidateL1, evictL1, lsc_scope);
 }
 
-void __intel_memfence_handler(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1)
+void __intel_memfence_handler(bool flushRW, bool isGlobal, bool invalidateL1, bool evictL1, Scope_t scope)
 {
     if (BIF_FLAG_CTRL_GET(OptDisable))
-        __intel_memfence_optnone(flushRW, isGlobal, invalidateL1, evictL1);
+        __intel_memfence_optnone(flushRW, isGlobal, invalidateL1, evictL1, scope);
     else
-        __intel_memfence(flushRW, isGlobal, invalidateL1, evictL1);
+        __intel_memfence(flushRW, isGlobal, invalidateL1, evictL1, scope);
 }
 
 // TYPEDMEMFENCE IMPLEMENTATION
@@ -98,19 +121,19 @@ static void __intel_atomic_work_item_fence( Scope_t Memory, uint Semantics )
         // although on some platforms they may be elided; platform-specific checks are performed in codegen
         if (Semantics & WorkgroupMemory)
         {
-           __intel_memfence_handler(false, false, false, false);
+           __intel_memfence_handler(false, false, false, false, Memory);
         }
         if (Semantics & CrossWorkgroupMemory)
         {
            if (Memory == Device || Memory == CrossDevice)
            {
-               __intel_memfence_handler(true, true, invalidateL1, evictL1);
+               __intel_memfence_handler(true, true, invalidateL1, evictL1, Memory);
            }
            else
            {
                // Single workgroup executes on one DSS and shares the same L1 cache.
                // If scope doesn't reach outside of workgroup, L1 flush can be skipped.
-               __intel_memfence_handler(false, true, false, false);
+               __intel_memfence_handler(false, true, false, false, Memory);
            }
         }
     }
diff --git a/IGC/Compiler/CISACodeGen/EmitVISAPass.cpp b/IGC/Compiler/CISACodeGen/EmitVISAPass.cpp
@@ -15865,8 +15865,9 @@ LSC_FENCE_OP EmitPass::getLSCMemoryFenceOp(bool IsGlobalMemFence, bool Invalidat
 
 void EmitPass::emitMemoryFence(llvm::Instruction* inst)
 {
-    static constexpr int ExpectedNumberOfArguments = 8;
-    IGC_ASSERT(IGCLLVM::getNumArgOperands(cast<CallInst>(inst)) == ExpectedNumberOfArguments);
+    static constexpr int ExpectedNumberOfBoolArguments = 8;
+    static constexpr int ScopeArgumentNumber = 8;
+    IGC_ASSERT(IGCLLVM::getNumArgOperands(cast<CallInst>(inst)) == (ExpectedNumberOfBoolArguments + 1));
     CodeGenContext* ctx = getAnalysis<CodeGenContextWrapper>().getCodeGenContext();
 
     // If passed a non-constant value for any of the parameters,
@@ -15882,8 +15883,9 @@ void EmitPass::emitMemoryFence(llvm::Instruction* inst)
     // Note: this flag is respected only for LSC case currently.
     // TODO: add support for non-LSC and typed fence.
     bool L1_Evict = true;
+    LSC_SCOPE Scope = LSC_SCOPE_GPU;
 
-    std::array<reference_wrapper<bool>, ExpectedNumberOfArguments> MemFenceArguments{
+    std::array<reference_wrapper<bool>, ExpectedNumberOfBoolArguments> MemFenceArguments{
       CommitEnable,
       L3_Flush_RW_Data,
       L3_Flush_Constant_Data,
@@ -15901,6 +15903,11 @@ void EmitPass::emitMemoryFence(llvm::Instruction* inst)
         }
     }
 
+    if (ConstantInt* CI = llvm::dyn_cast<llvm::ConstantInt>(inst->getOperand(ScopeArgumentNumber)))
+    {
+        Scope = (LSC_SCOPE) CI->getValue().getZExtValue();
+    }
+
     bool EmitFence = true;
 
     // Check whether we know this is a local fence. If we do, don't emit fence for a BDW+SKL/BXT only.
@@ -15929,8 +15936,11 @@ void EmitPass::emitMemoryFence(llvm::Instruction* inst)
     {
         // tgm should use GenISA_typedmemoryfence
         LSC_SFID sfid = Global_Mem_Fence ? LSC_UGM : LSC_SLM;
-        // ToDo: replace with fence instrinsics that take scope/op
-        LSC_SCOPE scope = Global_Mem_Fence ? LSC_SCOPE_GPU : LSC_SCOPE_GROUP;
+        // SLM only gets .group and .none
+        if(!Global_Mem_Fence && Scope >= LSC_SCOPE_LOCAL)
+        {
+            Scope = LSC_SCOPE_GROUP;
+        }
         // Do L1 evict only when default L1 cache policy is write-back.
         if (L1_Evict && m_pCtx->type == ShaderType::OPENCL_SHADER)
         {
@@ -15939,19 +15949,32 @@ void EmitPass::emitMemoryFence(llvm::Instruction* inst)
                 L1_Evict = static_cast<LSC_L1_L3_CC>(CLCtx->m_InternalOptions.StoreCacheDefault) == LSC_L1IAR_WB_L3C_WB;
         }
         // Change the scope from `GPU` to `Tile` on single-tile platforms to avoid L3 flush on DG2 and MTL and ARL.
-        if (scope == LSC_SCOPE_GPU &&
+        if (Scope == LSC_SCOPE_GPU &&
             !m_currShader->m_Platform->hasMultiTile() &&
             m_currShader->m_Platform->hasL3FlushOnGPUScopeInvalidate() &&
             IGC_IS_FLAG_DISABLED(EnableGPUFenceScopeOnSingleTileGPUs))
         {
-            scope = LSC_SCOPE_TILE;
+            Scope = LSC_SCOPE_TILE;
         }
+
+        // On some products cache is not flushed for scope < tile.
+        if (Global_Mem_Fence && (L1_Invalidate || L1_Evict) &&
+            (Scope < LSC_SCOPE_TILE) && (
+               m_currShader->m_Platform->getPlatformInfo().eProductFamily == IGFX_DG2
+            || m_currShader->m_Platform->getPlatformInfo().eProductFamily == IGFX_PVC
+            || m_currShader->m_Platform->getPlatformInfo().eProductFamily == IGFX_METEORLAKE
+            || m_currShader->m_Platform->getPlatformInfo().eProductFamily == IGFX_ARROWLAKE
+            ))
+        {
+            Scope = LSC_SCOPE_TILE;
+        }
+
         LSC_FENCE_OP op = getLSCMemoryFenceOp(Global_Mem_Fence, L1_Invalidate, L1_Evict);
         if (inst->getMetadata("forceFlushNone") || sfid == LSC_SLM)
         {
             op = LSC_FENCE_OP_NONE;
         }
-        m_encoder->LSC_Fence(sfid, scope, op);
+        m_encoder->LSC_Fence(sfid, Scope, op);
         m_encoder->Push();
         return;
     }
diff --git a/IGC/Compiler/Optimizer/OpenCLPasses/Atomics/ResolveOCLAtomics.cpp b/IGC/Compiler/Optimizer/OpenCLPasses/Atomics/ResolveOCLAtomics.cpp
@@ -323,7 +323,7 @@ void ResolveOCLAtomics::findLockUsers(Value* V)
 // br label %init_spinlock_var.end
 //
 // init_spinlock_var.end:                            ; preds = %init_spinlock_var.start, %entry
-//     call void @llvm.genx.GenISA.memoryfence(i1 true, i1 false, i1 false, i1 false, i1 false, i1 false, i1 true)
+//     call void @llvm.genx.GenISA.memoryfence(i1 true, i1 false, i1 false, i1 false, i1 false, i1 false, i1 true, i1 false, i32 0)
 //     call void @llvm.genx.GenISA.threadgroupbarrier()
 void ResolveOCLAtomics::generateLockInitilization(Function* F)
 {
@@ -361,11 +361,12 @@ void ResolveOCLAtomics::generateLockInitilization(Function* F)
     m_builder->CreateStore(m_localLock->getInitializer(), m_localLock);
     m_builder->CreateBr(initSpinLockEndBB);
 
-    // insert call void @llvm.genx.GenISA.memoryfence(i1 true, i1 false, i1 false, i1 false, i1 false, i1 false, i1 true, i1 false)
+    // insert call void @llvm.genx.GenISA.memoryfence(i1 true, i1 false, i1 false, i1 false, i1 false, i1 false, i1 true, i1 false, i32 0)
     //        call void @llvm.genx.GenISA.threadgroupbarrier()
     // to guarantee synchronization in accessing spin lock variable
     Value* trueValue = m_builder->getTrue();
     Value* falseValue = m_builder->getFalse();
+    Value* groupScopeValue = m_builder->getInt32(LSC_SCOPE_GROUP);
     Value* localMemFenceArgs[] =
     {
         trueValue,
@@ -376,6 +377,7 @@ void ResolveOCLAtomics::generateLockInitilization(Function* F)
         falseValue,
         trueValue,
         falseValue,
+        groupScopeValue
     };
     m_builder->SetInsertPoint(initSpinLockEndBB, initSpinLockEndBB->getFirstInsertionPt());
     Function* localMemFence = GenISAIntrinsic::getDeclaration(m_pModule, GenISAIntrinsic::GenISA_memoryfence);
diff --git a/IGC/Compiler/Optimizer/OpenCLPasses/NamedBarriers/NamedBarriersResolution.cpp b/IGC/Compiler/Optimizer/OpenCLPasses/NamedBarriers/NamedBarriersResolution.cpp
@@ -241,11 +241,16 @@ void NamedBarriersResolution::HandleNamedBarrierSyncHW(CallInst& NBarrierSyncCal
 
     Value* trueValue = IRB.getInt1(true);
     Value* falseValue = IRB.getInt1(false);
+    Value* gpuScopeValue = IRB.getInt32(LSC_SCOPE_GPU);
+    Value* groupScopeValue = IRB.getInt32(LSC_SCOPE_GROUP);
 
     ConstantInt* memFenceType = cast<ConstantInt>(NBarrierSyncCall.getArgOperand(1));
     // LOCAL = 1
     // GLOBAL = 2
-    Value* isGlobal = ( (int)memFenceType->getValue().getSExtValue() & 2 ) == 2 ? trueValue : falseValue;
+    bool isGlobal = ( (int)memFenceType->getValue().getSExtValue() & 2 ) == 2;
+    Value* isGlobalValue = isGlobal ? trueValue : falseValue;
+    // Conservatively for local barrier set GROUP scope, for global barrier set GPU scope
+    Value* scopeValue = isGlobal ? gpuScopeValue : groupScopeValue;
 
     GenIntrinsicInst::Create(
         GenISAIntrinsic::getDeclaration(module, GenISAIntrinsic::GenISA_memoryfence),
@@ -255,9 +260,10 @@ void NamedBarriersResolution::HandleNamedBarrierSyncHW(CallInst& NBarrierSyncCal
             falseValue, // bool flushConstant
             falseValue, // bool flushTexture
             falseValue, // bool flushIcache
-            isGlobal,   // bool isGlobal
+            isGlobalValue,   // bool isGlobal
             falseValue, // bool invalidateL1
             falseValue, // bool evictL1
+            scopeValue // int memory scope
         },
         "",
         &(NBarrierSyncCall));
diff --git a/IGC/Compiler/Optimizer/SynchronizationObjectCoalescing.cpp b/IGC/Compiler/Optimizer/SynchronizationObjectCoalescing.cpp
@@ -630,6 +630,8 @@ void SynchronizationObjectCoalescing::EraseRedundantGlobalScope(llvm::Instructio
     {
         constexpr uint32_t globalMemFenceArg = 5;
         pGenIntrinsicInst->setOperand(globalMemFenceArg, llvm::ConstantInt::getFalse(pGenIntrinsicInst->getOperand(globalMemFenceArg)->getType()));
+        constexpr uint32_t scopeMemFenceArg = 8;
+        pGenIntrinsicInst->setOperand(scopeMemFenceArg, llvm::ConstantInt::get(pGenIntrinsicInst->getOperand(scopeMemFenceArg)->getType(), static_cast<uint32_t>(LSC_SCOPE::LSC_SCOPE_GROUP)));
         break;
     }
     case llvm::GenISAIntrinsic::GenISA_LSCFence:
diff --git a/IGC/Compiler/tests/ResolveOCLAtomics/local_spin_lock.ll b/IGC/Compiler/tests/ResolveOCLAtomics/local_spin_lock.ll
diff --git a/IGC/GenISAIntrinsics/generator/input/Intrinsic_definitions.yml b/IGC/GenISAIntrinsics/generator/input/Intrinsic_definitions.yml
diff --git a/IGC/LLVM3DBuilder/BuiltinsFrontendDefinitions.hpp b/IGC/LLVM3DBuilder/BuiltinsFrontendDefinitions.hpp
diff --git a/visa/include/visa_igc_common_header.h b/visa/include/visa_igc_common_header.h

Original file line number	Diff line number	Diff line change
`@@ -630,6 +630,8 @@ void SynchronizationObjectCoalescing::EraseRedundantGlobalScope(llvm::Instructio`
`630`	`630`	`{`
`631`	`631`	`constexpr uint32_t globalMemFenceArg = 5;`
`632`	`632`	`pGenIntrinsicInst->setOperand(globalMemFenceArg, llvm::ConstantInt::getFalse(pGenIntrinsicInst->getOperand(globalMemFenceArg)->getType()));`
	`633`	`+ constexpr uint32_t scopeMemFenceArg = 8;`
	`634`	`+ pGenIntrinsicInst->setOperand(scopeMemFenceArg, llvm::ConstantInt::get(pGenIntrinsicInst->getOperand(scopeMemFenceArg)->getType(), static_cast<uint32_t>(LSC_SCOPE::LSC_SCOPE_GROUP)));`
`633`	`635`	`break;`
`634`	`636`	`}`
`635`	`637`	`case llvm::GenISAIntrinsic::GenISA_LSCFence:`