[AMDGPU] - Mark inverse.ballot as not convergent

[OutOfCache]

Summary

inverse.ballot checks if a cc bit is set for the current
lane. Therefore, it is not convergent.

[sameerds]

The change description and other places should say "inverse.ballot is not convergent". The keyword "divergent" is the opposite of "uniform". The opposite of convergent is simply "not convergent".

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1954

This is updating the wrong intrinsic. readlane is definitely convergent.

[foad]

Mark inverse.ballot as divergent

I would say "not convergent". "divergent" is the opposite of "uniform". I don't think we use it to mean the opposite of "convergent", do we?

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1954

You've changed the wrong intrinsic.

[foad]

Snap!

[OutOfCache]

Modifying the correct intrinsic

[sameerds]

Surprising that this didn't affect any lit test. If there are no tests, then they need to be added.

[arsenm]

In D156088#4528503, @sameerds wrote:

Surprising that this didn't affect any lit test. If there are no tests, then they need to be added.

I believe we don't have a dedicated test-this-is-convergent test for all the intrinsics anywhere

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
              ImmArg<ArgIndex<2>>, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 def int_amdgcn_ballot :
   Intrinsic<[llvm_anyint_ty], [llvm_i1_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 def int_amdgcn_inverse_ballot :
   Intrinsic<[llvm_i1_ty], [llvm_anyint_ty],
-            [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
+            [IntrNoMem, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 class AMDGPUWaveReduce<LLVMType data_ty = llvm_anyint_ty> : Intrinsic<
     [data_ty], 
     [
       LLVMMatchType<0>,   // llvm value to reduce (SGPR/VGPR)
       llvm_i32_ty         // Reduction Strategy Switch for lowering ( 0: Default,
                           //                                          1: Iterative strategy, and
                           //                                          2. DPP)
     ],
     [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, ImmArg<ArgIndex<1>>]>;
 
 def int_amdgcn_wave_reduce_umin : AMDGPUWaveReduce;
 def int_amdgcn_wave_reduce_umax : AMDGPUWaveReduce;
 
 def int_amdgcn_readfirstlane :
   ClangBuiltin<"__builtin_amdgcn_readfirstlane">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 // The lane argument must be uniform across the currently active threads of the
 // current wave. Otherwise, the result is undefined.
 def int_amdgcn_readlane :
   ClangBuiltin<"__builtin_amdgcn_readlane">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
 // The value to write and lane select arguments must be uniform across the
 // currently active threads of the current wave. Otherwise, the result is
 // undefined.
 def int_amdgcn_writelane :
   ClangBuiltin<"__builtin_amdgcn_writelane">,
   Intrinsic<[llvm_i32_ty], [
     llvm_i32_ty,    // uniform value to write: returned by the selected lane