From 8477ca6e8e2a2c5adcfd600780a5eecb2ca0b153 Mon Sep 17 00:00:00 2001 From: Raymond Tian Date: Mon, 1 Jul 2024 11:54:04 -0700 Subject: [PATCH] [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. --- clang/lib/Sema/SemaExpr.cpp | 3 +- .../SemaCUDA/amdgpu-builtins-pointer-args.cu | 28 +++++++++++++++++++ 2 files changed, 30 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 6ecf9e177ba4..cc811a40185c 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 000000000000..034bfaec25cf --- /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}} +} +