[AMDGPU] Add sema check for global_atomic_fadd_v2f16 builtin #158145

tcgu-amd · 2025-09-11T20:15:16Z

The builtin expects a vector _Float16 of length 2, but clang does not emit errors when the wrong argument types is supplied (e.g. half2). This causes the compilation to pass but error during LTO.

This fix the issue by adding sema checks to the builtins.

github-actions · 2025-09-11T20:15:37Z

Thank you for submitting a Pull Request (PR) to the LLVM Project!

This PR will be automatically labeled and the relevant teams will be notified.

If you wish to, you can add reviewers by using the "Reviewers" section on this page.

If this is not working for you, it is probably because you do not have write permissions for the repository. In which case you can instead tag reviewers by name in a comment by using @ followed by their GitHub username.

If you have received no comments on your PR for a week, you can request a review by "ping"ing the PR by adding a comment “Ping”. The common courtesy "ping" rate is once a week. Please remember that you are asking for valuable time from other developers.

If you have further questions, they may be answered by the LLVM GitHub User Guide.

You can also ask questions in a comment on this PR, on the LLVM Discord or on the forums.

llvmbot · 2025-09-11T20:16:11Z

@llvm/pr-subscribers-backend-amdgpu

Author: Tim Gu (tcgu-amd)

Changes

Addresses ROCm/ROCm#5253.

The builtin expects a vector _Float16 of length 2, but clang does not emit errors when the wrong argument types is supplied (e.g. half2). This causes the compilation to pass but error during LTO.

This fix the issue by adding sema checks to the builtins.

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

2 Files Affected:

(modified) clang/include/clang/Sema/SemaAMDGPU.h (+2)
(modified) clang/lib/Sema/SemaAMDGPU.cpp (+32)

diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h
index bac812a9d4fcf..9ca4418349fff 100644
--- a/clang/include/clang/Sema/SemaAMDGPU.h
+++ b/clang/include/clang/Sema/SemaAMDGPU.h
@@ -31,6 +31,8 @@ class SemaAMDGPU : public SemaBase {
   bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
                                unsigned NumDataArgs);
 
+  bool checkAMDGCNAtomicFaddV2F16Type(CallExpr *TheCall);
+
   /// Create an AMDGPUWavesPerEUAttr attribute.
   AMDGPUFlatWorkGroupSizeAttr *
   CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min,
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index baba503239e9f..c0f17f185982e 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -109,6 +109,10 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
   case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
     return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
+  case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
+  case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
+    return checkAMDGCNAtomicFaddV2F16Type(TheCall);
   default:
     return false;
   }
@@ -436,4 +440,32 @@ void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
   addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
 }
 
+
+bool SemaAMDGPU::checkAMDGCNAtomicFaddV2F16Type(CallExpr *TheCall) {
+  // Check that the pointer argument is a pointer to v2f16
+
+  Expr *Arg = TheCall->getArg(1);
+  QualType ArgType = Arg->getType();
+
+  // Check if it's a vector type
+  if (!ArgType->isVectorType()) {
+    Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types)
+        << "expected _Float16 vector of length 2" << ArgType
+        << Arg->getSourceRange();
+    return true;
+  }
+
+  const VectorType *VT = ArgType->getAs<VectorType>();
+
+  // Check element type (should be _Float16) and vector length (should be 2)
+  QualType ElementType = VT->getElementType();
+  if (!ElementType->isFloat16Type() || VT->getNumElements() != 2) {
+    Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types)
+        << "expected _Float16 vector of length 2" << ArgType
+        << Arg->getSourceRange();
+    return true;
+  }
+
+  return false;
+}
 } // namespace clang

llvmbot · 2025-09-11T20:16:12Z

@llvm/pr-subscribers-clang

Author: Tim Gu (tcgu-amd)

Changes

Addresses ROCm/ROCm#5253.

The builtin expects a vector _Float16 of length 2, but clang does not emit errors when the wrong argument types is supplied (e.g. half2). This causes the compilation to pass but error during LTO.

This fix the issue by adding sema checks to the builtins.

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

2 Files Affected:

(modified) clang/include/clang/Sema/SemaAMDGPU.h (+2)
(modified) clang/lib/Sema/SemaAMDGPU.cpp (+32)

diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h
index bac812a9d4fcf..9ca4418349fff 100644
--- a/clang/include/clang/Sema/SemaAMDGPU.h
+++ b/clang/include/clang/Sema/SemaAMDGPU.h
@@ -31,6 +31,8 @@ class SemaAMDGPU : public SemaBase {
   bool checkMovDPPFunctionCall(CallExpr *TheCall, unsigned NumArgs,
                                unsigned NumDataArgs);
 
+  bool checkAMDGCNAtomicFaddV2F16Type(CallExpr *TheCall);
+
   /// Create an AMDGPUWavesPerEUAttr attribute.
   AMDGPUFlatWorkGroupSizeAttr *
   CreateAMDGPUFlatWorkGroupSizeAttr(const AttributeCommonInfo &CI, Expr *Min,
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index baba503239e9f..c0f17f185982e 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -109,6 +109,10 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
   case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
     return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
+  case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
+  case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
+    return checkAMDGCNAtomicFaddV2F16Type(TheCall);
   default:
     return false;
   }
@@ -436,4 +440,32 @@ void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
   addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
 }
 
+
+bool SemaAMDGPU::checkAMDGCNAtomicFaddV2F16Type(CallExpr *TheCall) {
+  // Check that the pointer argument is a pointer to v2f16
+
+  Expr *Arg = TheCall->getArg(1);
+  QualType ArgType = Arg->getType();
+
+  // Check if it's a vector type
+  if (!ArgType->isVectorType()) {
+    Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types)
+        << "expected _Float16 vector of length 2" << ArgType
+        << Arg->getSourceRange();
+    return true;
+  }
+
+  const VectorType *VT = ArgType->getAs<VectorType>();
+
+  // Check element type (should be _Float16) and vector length (should be 2)
+  QualType ElementType = VT->getElementType();
+  if (!ElementType->isFloat16Type() || VT->getNumElements() != 2) {
+    Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types)
+        << "expected _Float16 vector of length 2" << ArgType
+        << Arg->getSourceRange();
+    return true;
+  }
+
+  return false;
+}
 } // namespace clang

shiltian · 2025-09-11T20:45:18Z

I'm not sure what this PR is trying to fix, but we do have type enforcement in clang/include/clang/Basic/BuiltinsAMDGPU.def?

arsenm

Missing tests

arsenm · 2025-09-12T06:49:23Z

clang/lib/Sema/SemaAMDGPU.cpp

+
+  // Check element type (should be _Float16) and vector length (should be 2)
+  QualType ElementType = VT->getElementType();
+  if (!ElementType->isFloat16Type() || VT->getNumElements() != 2) {


The builtin already knows what type it is. You should not need to add custom type checking

Thanks for the comment! That's what I thought as well, but for some reason without explicit checking it here clang does not seem to catch the type mismatch. I am new to clang development so I don't know if that's expected.

arsenm · 2025-09-12T06:49:35Z

clang/lib/Sema/SemaAMDGPU.cpp

+  // Check if it's a vector type
+  if (!ArgType->isVectorType()) {
+    Diag(Arg->getBeginLoc(), diag::err_typecheck_call_different_arg_types)
+        << "expected _Float16 vector of length 2" << ArgType


This shouldn't require a custom message

There doesn't seem to be a "expect vector" message so I resorted to this. What do you suggest I use instead? Thanks!

Plus no diagnostic messages should be hardcoded strings like this.

You shouldn't need to do anything for the type checking. All of these builtins are marked with "t" for custom typechecking, but I don't see why. Can you just remove that from the builtin definition?

…oat16

shiltian · 2025-09-18T17:20:37Z

no test?

arsenm · 2025-09-19T11:48:22Z

clang/include/clang/Basic/BuiltinsAMDGPU.def

 TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst")
 TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")
-TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2xV2x*3V2x", "n", "atomic-ds-pk-add-16-insts")


Also make sure to test this with all of these similar intrinsics, the "t"s are all suspicious (I'm guessing if it was for anything, it was hacking around the address space of the pointer across the languages)

arsenm · 2025-09-25T06:59:59Z

clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c

+typedef _Float16 v2f16 __attribute__((ext_vector_type(2)));
+typedef float v2f32 __attribute__((ext_vector_type(2)));
+typedef _Float16 v4f16 __attribute__((ext_vector_type(4)));
+


Can you also test some of the cases with the type embedded in a structure like was crashing

I think the type was from hip though. How do I add a type from hip_fp16 to a clang test? Do I need to put the test in one of the hip/opencl folders? Thanks!

arsenm · 2025-09-25T07:00:51Z

clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c

+  __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16, val_v4f16); // expected-error{{passing 'v4f16'}}
+  __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16); // expected-error{{too few arguments to function call}}
+  __builtin_amdgcn_ds_atomic_fadd_v2f16(ptr_v2f16, val, val); // expected-error{{too many arguments to function call}}
+}


bf16 vector cases and f32 cases are also marked with t, and probably broken in the same way

arsenm · 2025-09-25T07:01:52Z

clang/test/Sema/builtin-amdgcn-atomic-fadd-v2f16-type-err.c

@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -verify %s


Can you verify we have test coverage using OpenCL plus one of the C/C++/HIP variants

llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Sep 11, 2025

lamb-j requested review from Pierre-vh, arsenm and shiltian September 11, 2025 23:05

arsenm requested changes Sep 12, 2025

View reviewed changes

Removed custom type checking & changed type string from __fp16 to _Fl…

2790a7a

…oat16

tcgu-amd force-pushed the users/tcgu/builtin-sema branch from 2192052 to 2790a7a Compare September 18, 2025 17:03

arsenm reviewed Sep 19, 2025

View reviewed changes

Adding tests.

f7a5b3b

arsenm reviewed Sep 25, 2025

View reviewed changes

tcgu-amd mentioned this pull request Sep 29, 2025

[Issue]: lld crashes when try to link the rdc sources with builtin half2 atomic function ROCm/ROCm#5253

Closed

		@@ -0,0 +1,46 @@
		// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -verify %s

[AMDGPU] Add sema check for global_atomic_fadd_v2f16 builtin #158145

Are you sure you want to change the base?

[AMDGPU] Add sema check for global_atomic_fadd_v2f16 builtin #158145

Conversation

tcgu-amd commented Sep 11, 2025

Uh oh!

github-actions bot commented Sep 11, 2025

Uh oh!

llvmbot commented Sep 11, 2025

Uh oh!

llvmbot commented Sep 11, 2025

Uh oh!

shiltian commented Sep 11, 2025

Uh oh!

arsenm left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

shiltian commented Sep 18, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

tcgu-amd Oct 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

tcgu-amd Oct 14, 2025 •

edited

Loading