[clang] [llvm] [AMDGPU] Add `wave_id` and `wave_shuffle` Clang builtins. (PR #179492)

Matt Arsenault via cfe-commits cfe-commits at lists.llvm.org
Sat Feb 7 05:38:32 PST 2026


================
@@ -296,6 +296,31 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
     }
     return false;
   }
+  case AMDGPU::BI__builtin_amdgcn_wave_shuffle: {
+    Expr *Val = TheCall->getArg(0);
+    QualType ValTy = Val->getType();
+
+    if ((!ValTy->isIntegerType() && !ValTy->isFloatingType()) ||
+        SemaRef.getASTContext().getTypeSize(ValTy) > 32)
+      return Diag(Val->getExprLoc(), diag::err_builtin_invalid_arg_type)
+             << Val << /*scalar=*/1 << /*'int'=*/4 << /*floating point=*/2
+             << ValTy;
+
+    Expr *Idx = TheCall->getArg(1);
+    QualType IdxTy = Idx->getType();
+    if (!IdxTy->isIntegerType())
+      return Diag(Idx->getExprLoc(), diag::err_typecheck_expect_int) << IdxTy;
+    if (SemaRef.getASTContext().getTypeSize(IdxTy) > 32)
----------------
arsenm wrote:

Should only accept exactly i32?

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


More information about the cfe-commits mailing list