https://github.com/Sisyph updated
https://github.com/llvm/llvm-project/pull/138141
>From f5cdefe8200d9c9f567d6b4a276a5587e44ac1fa Mon Sep 17 00:00:00 2001
From: Joe Nash <joseph.n...@amd.com>
Date: Thu, 1 May 2025 10:00:47 -0400
Subject: [PATCH] [Sema] Fix bug in builtin AS override
Fix the logic in rewriteBuiltinFunctionDecl to work when the builtin
has a pointer parameter with an address space and one without a fixed
address space. A builtin fitting these criteria was recently added.
Change the attribute string to perform type checking on it, so without
the sema change compilation would fail with a wrong number of arguments
error.
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 +-
clang/lib/Sema/SemaExpr.cpp | 6 ++----
2 files changed, 3 insertions(+), 5 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 730fd15913c11..802b4be42419d 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -257,7 +257,7 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16,
"V2sV2s*0V2s", "t", "at
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t",
"atomic-global-pk-add-bf16-inst")
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t",
"atomic-ds-pk-add-16-insts")
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t",
"atomic-ds-pk-add-16-insts")
-TARGET_BUILTIN(__builtin_amdgcn_load_to_lds, "vv*v*3IUiIiIUi", "t",
"vmem-to-lds-load-insts")
+TARGET_BUILTIN(__builtin_amdgcn_load_to_lds, "vv*v*3IUiIiIUi", "",
"vmem-to-lds-load-insts")
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t",
"vmem-to-lds-load-insts")
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 21bd7315e3dd4..85c7f995233a7 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6395,7 +6395,8 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema
*Sema, ASTContext &Context,
return nullptr;
Expr *Arg = ArgRes.get();
QualType ArgType = Arg->getType();
- if (!ParamType->isPointerType() || ParamType.hasAddressSpace() ||
+ if (!ParamType->isPointerType() ||
+ ParamType->getPointeeType().hasAddressSpace() ||
!ArgType->isPointerType() ||
!ArgType->getPointeeType().hasAddressSpace() ||
isPtrSizeAddressSpace(ArgType->getPointeeType().getAddressSpace())) {
@@ -6404,9 +6405,6 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema
*Sema, ASTContext &Context,
}
QualType PointeeType = ParamType->getPointeeType();
- if (PointeeType.hasAddressSpace())
- continue;
-
NeedsNewDecl = true;
LangAS AS = ArgType->getPointeeType().getAddressSpace();
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
