[PATCH] D137142: [WIP] DivergenceAnalysis: Infer uniformity from assume calls

[Matt Arsenault via Phabricator via llvm-commits]

arsenm created this revision.
arsenm added reviewers: AMDGPU, sameerds, nhaehnle, simoll, jdoerfert.
Herald added subscribers: kosarev, foad, kerbowa, hiraditya, jvesely.
Herald added a project: All.
arsenm requested review of this revision.
Herald added a subscriber: wdng.
Herald added a project: LLVM.

I believe this patch is OK as is, but is currently useless in practice and
I'm not sure how useful this really will be. Theoretically this should allow
something like:

      

kernel void foo(global int* global* arg_ptr) {

  global int* ptr = arg_ptr[get_global_id(0)];
  __builtin_assume(sub_group_all(ptr != NULL));
  if (ptr != NULL) {
      *ptr += 1;
  }

}

      

to use a scalar branch around the pointer dereference. There are a few obstacles
to this working today. First, using sub_group_all generates this warning
for some reason:

      

warning: the argument to '__builtin_assume' has side effects that will be discarded

Second, the device libraries are still using the legacy llvm.amdgcn.icmp intrinsics
instead of ballot.

      

Third, the device libraries are still using an inline assembly hack in lieu of
convergence tokens.

      

Fourth, even if those issues are avoided, the branch is still treated
as divergent when ultimately selected.


https://reviews.llvm.org/D137142

Files:
  llvm/include/llvm/Analysis/DivergenceAnalysis.h
  llvm/include/llvm/Analysis/TargetTransformInfo.h
  llvm/include/llvm/Analysis/TargetTransformInfoImpl.h
  llvm/include/llvm/CodeGen/BasicTTIImpl.h
  llvm/lib/Analysis/AssumptionCache.cpp
  llvm/lib/Analysis/DivergenceAnalysis.cpp
  llvm/lib/Analysis/LegacyDivergenceAnalysis.cpp
  llvm/lib/Analysis/TargetTransformInfo.cpp
  llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.cpp
  llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.h
  llvm/test/Analysis/DivergenceAnalysis/AMDGPU/assume.ll
  llvm/unittests/Analysis/DivergenceAnalysisTest.cpp

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D137142.472197.patch
Type: text/x-patch
Size: 26382 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20221101/82bde9b1/attachment.bin>