[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 5f6da1bf00c4b3bdb8aa75913629ebee262f3783 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..67d28028a4c62
--- /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 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}}
+}
+