https://github.com/raymondytian created
https://github.com/llvm/llvm-project/pull/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.
>From e0ab54fd8afa3ae0ee321504dd272bd787979e62 Mon Sep 17 00:00:00 2001
From: Raymond Tian <rymdt...@gmail.com>
Date: Tue, 18 Jun 2024 09:58:32 -0700
Subject: [PATCH] [HIP][Clang][Sema] Fix crash when calling builtins with
pointer arguments
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 | 24 +++++++++++++++++++
2 files changed, 26 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 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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
