[Mlir-commits] [clang] [llvm] [mlir] [AMDGPU] Fix missing cache bypass scope on seq_cst buffer fat pointer atomics (PR #213585)

Krzysztof Drewniak llvmlistbot at llvm.org
Thu Aug 6 08:55:16 PDT 2026


================
@@ -2180,19 +2194,19 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
             EmitScalarExpr(E->getArg(4))});
   }
   case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
-    return emitBuiltinWithOneOverloadedType<5>(
+    return emitAMDGPURawPtrBufferBuiltinWithScope(
         *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_add);
   case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
   case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16:
-    return emitBuiltinWithOneOverloadedType<5>(
+    return emitAMDGPURawPtrBufferBuiltinWithScope(
         *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fadd);
   case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
   case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64:
-    return emitBuiltinWithOneOverloadedType<5>(
+    return emitAMDGPURawPtrBufferBuiltinWithScope(
----------------
krzysz00 wrote:

Is there some way in C++ to specify the scope? Should we add one? Maybe give this thing an optional string argument?

https://github.com/llvm/llvm-project/pull/213585


More information about the Mlir-commits mailing list