https://github.com/raymondytian updated
https://github.com/llvm/llvm-project/pull/95957
>From 8bd331321b2121321df37a0452e7bece7023053c 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..653337683f351
--- /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}}
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
