Author: Raymond Tian Date: 2024-07-01T14:54:04-04:00 New Revision: 8477ca6e8e2a2c5adcfd600780a5eecb2ca0b153
URL: https://github.com/llvm/llvm-project/commit/8477ca6e8e2a2c5adcfd600780a5eecb2ca0b153 DIFF: https://github.com/llvm/llvm-project/commit/8477ca6e8e2a2c5adcfd600780a5eecb2ca0b153.diff LOG: [HIP][Clang][Sema] Fix crash when calling builtins with pointer arguments (#95957) 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. Added: clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu Modified: clang/lib/Sema/SemaExpr.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 6ecf9e177ba42..cc811a40185ce 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6500,7 +6500,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..034bfaec25cfd --- /dev/null +++ b/clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu @@ -0,0 +1,28 @@ +// REQUIRES: amdgpu-registered-target +// REQUIRES: x86-registered-target + +// 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 at least 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}} +} + _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
