[clang] [HIP][Clang][Sema] Fix crash when calling builtins with pointer arguments (PR #95957)

[Raymond Tian via cfe-commits]

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 <rymdtian at 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}}
+}
+