[llvm] AMDGPU: Make llvm.amdgcn.endpgm convergent (PR #74555)

Matt Arsenault via llvm-commits llvm-commits at lists.llvm.org
Tue Dec 5 20:12:38 PST 2023


https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/74555

I don't believe this makes any practical difference.

Fixes #64013

>From 16fe53385bc5864711f9b40601e1edaf597d160d Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Wed, 15 Nov 2023 16:32:45 +0900
Subject: [PATCH] AMDGPU: Make llvm.amdgcn.endpgm convergent

I don't believe this makes any practical difference.

Fixes #64013
---
 llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 06f9c0445bcea..bc9f99783d98f 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2105,7 +2105,8 @@ def int_amdgcn_wqm_vote : Intrinsic<[llvm_i1_ty],
 def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], [IntrNoCallback, IntrNoFree]>;
 
 def int_amdgcn_endpgm : ClangBuiltin<"__builtin_amdgcn_endpgm">,
-  Intrinsic<[], [], [IntrNoReturn, IntrCold, IntrNoMem, IntrHasSideEffects, IntrNoCallback, IntrNoFree]
+  Intrinsic<[], [], [IntrNoReturn, IntrCold, IntrNoMem, IntrHasSideEffects, IntrConvergent,
+                     IntrNoCallback, IntrNoFree]
 >;
 
 // If false, mark all active lanes as helper lanes until the end of program.



More information about the llvm-commits mailing list