Revert " Add builtin/intrinsic global_(load|store)_b128"

Author: ronlieb

clang/docs/LanguageExtensions.rst

@@ -4901,52 +4901,6 @@ for a concrete target, and shall reflect the latter's identity and features.
 Thus, it is possible to author high-level code, in e.g. HIP, that is target
 adaptive in a dynamic fashion, contrary to macro based mechanisms.
 
-__builtin_amdgcn_global_load_b128 and __builtin_amdgcn_global_store_b128
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-
-Signature:
-
-.. code-block:: c
-
-    typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u;
-    typedef v4u __attribute__((address_space(1))) *global_ptr_to_v4u;
-
-    v4u __builtin_amdgcn_global_load_b128(
-       v4u __attribute__((address_space(1))) *src,
-       const char                            *scope);
-
-    void __builtin_amdgcn_global_store_b128(
-       v4u __attribute__((address_space(1))) *dst,
-       v4u                                    data,
-       const char                            *scope);
-
-Load or store a vector of 4 unsigned integers from or to global memory with
-cache behavior specified by `scope` which must be a string literal.
-
-Valid values for `scope` are:
-
-===================== ==========================================================
-scope                 architecture name
-===================== ==========================================================
-``"wavefront"``       wave
-
-``"workgroup"``       group
-
-``"agent"``           device
-
-``""`` (empty string) system
-===================== ==========================================================
-
-These builtins are only supported on gfx942 and gfx950 devices.
-
-For semantics on gfx942, see Tables 47 and 48 in section 9.1.10 "Memory Scope
-and Temporal Controls" of the "AMD Instinct MI300" Instruction Set Architecture
-Reference.
-
-For semantics on gfx950, see Tables 49 and 50 in section 9.1.10 "Memory Scope
-and Temporal Controls" of the CDNA4 Instruction Set Architecture Reference.
-
-
 ARM/AArch64 Language Extensions
 -------------------------------
 

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -165,9 +165,6 @@ BUILTIN(__builtin_amdgcn_raw_buffer_load_b128, "V4UiQbiiIi", "n")
 
 TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi", "t", "vmem-to-lds-load-insts")
 
-TARGET_BUILTIN(__builtin_amdgcn_global_load_b128, "V4UiV4Ui*1cC*", "n", "gfx940-insts")
-TARGET_BUILTIN(__builtin_amdgcn_global_store_b128, "vV4Ui*1V4UicC*", "n", "gfx940-insts")
-
 //===----------------------------------------------------------------------===//
 // Ballot builtins.
 //===----------------------------------------------------------------------===//

clang/include/clang/Sema/SemaAMDGPU.h

@@ -30,8 +30,6 @@ class SemaAMDGPU : public SemaBase {
 
   bool CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall);
 
-  bool checkScopedMemAccessFunctionCall(CallExpr *TheCall);
-
   bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
                                unsigned NumDataArgs);
 

clang/lib/CodeGen/CGBuiltin.cpp

@@ -20116,26 +20116,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy});
     return Builder.CreateCall(F, {Addr});
   }
-  case AMDGPU::BI__builtin_amdgcn_global_load_b128:
-  case AMDGPU::BI__builtin_amdgcn_global_store_b128: {
-    const bool IsStore =
-        BuiltinID == AMDGPU::BI__builtin_amdgcn_global_store_b128;
-    LLVMContext &Ctx = CGM.getLLVMContext();
-    SmallVector<Value *, 5> Args = {EmitScalarExpr(E->getArg(0))}; // addr
-    if (IsStore)
-      Args.push_back(EmitScalarExpr(E->getArg(1))); // data
-    const unsigned ScopeIdx = E->getNumArgs() - 1;
-    StringRef ScopeLit =
-        cast<StringLiteral>(E->getArg(ScopeIdx)->IgnoreParenCasts())
-            ->getString();
-    llvm::MDNode *MD =
-        llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, ScopeLit)});
-    Args.push_back(llvm::MetadataAsValue::get(Ctx, MD)); // scope
-    llvm::Function *F =
-        CGM.getIntrinsic(IsStore ? Intrinsic::amdgcn_global_store_b128
-                                 : Intrinsic::amdgcn_global_load_b128);
-    return Builder.CreateCall(F, Args);
-  }
   case AMDGPU::BI__builtin_amdgcn_get_fpenv: {
     Function *F = CGM.getIntrinsic(Intrinsic::get_fpenv,
                                    {llvm::Type::getInt64Ty(getLLVMContext())});

clang/lib/Sema/SemaAMDGPU.cpp

@@ -84,9 +84,6 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_update_dpp: {
     return checkMovDPPFunctionCall(TheCall, 6, 2);
   }
-  case AMDGPU::BI__builtin_amdgcn_global_load_b128:
-  case AMDGPU::BI__builtin_amdgcn_global_store_b128:
-    return checkScopedMemAccessFunctionCall(TheCall);
   default:
     return false;
   }
@@ -132,19 +129,6 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
   return false;
 }
 
-bool SemaAMDGPU::checkScopedMemAccessFunctionCall(CallExpr *TheCall) {
-  bool Fail = false;
-  // Last argument is a string literal
-  Expr *Arg = TheCall->getArg(TheCall->getNumArgs() - 1);
-  auto Scope = dyn_cast<StringLiteral>(Arg->IgnoreParenCasts());
-  if (!Scope) {
-    Fail = true;
-    Diag(TheCall->getBeginLoc(), diag::err_expr_not_string_literal)
-        << Arg->getSourceRange();
-  }
-  return Fail;
-}
-
 bool SemaAMDGPU::checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
                                          unsigned NumDataArgs) {
   assert(NumDataArgs <= 2);

clang/lib/Sema/SemaExpr.cpp

@@ -6247,8 +6247,7 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema *Sema, ASTContext &Context,
       return nullptr;
     Expr *Arg = ArgRes.get();
     QualType ArgType = Arg->getType();
-    if (!ParamType->isPointerType() ||
-        ParamType->getPointeeType().hasAddressSpace() ||
+    if (!ParamType->isPointerType() || ParamType.hasAddressSpace() ||
         !ArgType->isPointerType() ||
         !ArgType->getPointeeType().hasAddressSpace() ||
         isPtrSizeAddressSpace(ArgType->getPointeeType().getAddressSpace())) {
@@ -6257,6 +6256,9 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema *Sema, ASTContext &Context,
     }
 
     QualType PointeeType = ParamType->getPointeeType();
+    if (PointeeType.hasAddressSpace())
+      continue;
+
     NeedsNewDecl = true;
     LangAS AS = ArgType->getPointeeType().getAddressSpace();