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

[via cfe-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-amdgpu

Author: Raymond Tian (raymondytian)

<details>
<summary>Changes</summary>

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.

---
Full diff: https://github.com/llvm/llvm-project/pull/95957.diff


2 Files Affected:

- (modified) clang/lib/Sema/SemaExpr.cpp (+2-1) 
- (added) clang/test/SemaCUDA/amdgpu-builtins-pointer-args.cu (+24) 


``````````diff
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

``````````

</details>


https://github.com/llvm/llvm-project/pull/95957