https://github.com/rymdtian updated
https://github.com/llvm/llvm-project/pull/95957
>From f001ebc8b2670240930de82bb5b4692a0376248e 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 | 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 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..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
