Author: Joe Nash Date: 2025-06-06T09:31:32-04:00 New Revision: 7ae6c4319eb49a58a13c423686a110afa554e92d
URL: https://github.com/llvm/llvm-project/commit/7ae6c4319eb49a58a13c423686a110afa554e92d DIFF: https://github.com/llvm/llvm-project/commit/7ae6c4319eb49a58a13c423686a110afa554e92d.diff LOG: [Sema] Fix bug in builtin AS override (#138141) 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. Added: Modified: clang/include/clang/Basic/BuiltinsAMDGPU.def clang/lib/Sema/SemaExpr.cpp Removed: ################################################################################ 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 d67e29d8d606e..9c322deb55e00 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
