llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Raymond Tian (raymondytian) <details> <summary>Changes</summary> Crashed when the number of args passed was less than number of parameters in builtin definition, because we were indexing the list of args while iterating through the entire number of parameters. --- Full diff: https://github.com/llvm/llvm-project/pull/95957.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaExpr.cpp (+2-1) - (added) clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu (+24) ``````````diff diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index acaff304be193..83c3d094c496f 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6628,7 +6628,8 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc, // the parameter type. if (getLangOpts().HIP && getLangOpts().CUDAIsDevice && FD && FD->getBuiltinID()) { - for (unsigned Idx = 0; Idx < FD->param_size(); ++Idx) { + for (unsigned Idx = 0; Idx < ArgExprs.size() && Idx < FD->param_size(); + ++Idx) { ParmVarDecl *Param = FD->getParamDecl(Idx); if (!ArgExprs[Idx] || !Param || !Param->getType()->isPointerType() || !ArgExprs[Idx]->getType()->isPointerType()) diff --git a/clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu b/clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu new file mode 100644 index 0000000000000..d64fa02469a69 --- /dev/null +++ b/clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -fsyntax-only -verify %s + +void call_amdgpu_builtins() { + __builtin_amdgcn_fence(); // expected-error {{too few arguments to function call, expected 2, have 0}} + __builtin_amdgcn_atomic_inc32(); // expected-error {{too few arguments to function call, expected 4, have 0}} + __builtin_amdgcn_atomic_inc32(0); // expected-error {{too few arguments to function call, expected 4, have 1}} + __builtin_amdgcn_atomic_inc32(0, 0); // expected-error {{too few arguments to function call, expected 4, have 2}} + __builtin_amdgcn_atomic_inc32(0, 0, 0); // expected-error {{too few arguments to function call, expected 4, have 3}} + __builtin_amdgcn_atomic_inc64(); // expected-error {{too few arguments to function call, expected 4, have 0}} + __builtin_amdgcn_atomic_dec32(); // expected-error {{too few arguments to function call, expected 4, have 0}} + __builtin_amdgcn_atomic_dec64(); // expected-error {{too few arguments to function call, expected 4, have 0}} + __builtin_amdgcn_div_scale(); // expected-error {{too few arguments to function call, expected 4, have 0}} + __builtin_amdgcn_div_scale(0); // expected-error {{too few arguments to function call, expected 4, have 1}} + __builtin_amdgcn_div_scale(0, 0); // expected-error {{too few arguments to function call, expected 4, have 2}} + __builtin_amdgcn_div_scale(0, 0, 0); // expected-error {{too few arguments to function call, expected 4, have 3}} + __builtin_amdgcn_div_scalef(); // expected-error {{too few arguments to function call, expected 4, have 0}} + __builtin_amdgcn_ds_faddf(); // expected-error {{too few arguments to function call, expected 5, have 0}} + __builtin_amdgcn_ds_fminf(); // expected-error {{too few arguments to function call, expected 5, have 0}} + __builtin_amdgcn_ds_fmaxf(); // expected-error {{too few arguments to function call, expected 5, have 0}} + __builtin_amdgcn_ds_append(); // expected-error {{too few arguments to function call, expected 1, have 0}} + __builtin_amdgcn_ds_consume(); // expected-error {{too few arguments to function call, expected 1, have 0}} + __builtin_amdgcn_is_shared(); // expected-error {{too few arguments to function call, expected 1, have 0}} + __builtin_amdgcn_is_private(); // expected-error {{too few arguments to function call, expected 1, have 0}} +} \ No newline at end of file `````````` </details> https://github.com/llvm/llvm-project/pull/95957 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
