Page MenuHomePhabricator

[AMDGPU] Use AssumptionCacheTracker in the divrem32 expansion
ClosedPublic

Authored by rampitec on Jul 24 2018, 3:00 PM.

Details

Summary

This allows to use a short div/rem expansion given a kernel with __builtin_assume() call:

kernel void divrem24_uint(global int *a, uint x) {
  uint id = get_local_id(0);
  __builtin_assume(x < 42);
  a[id / x] = 0;
}

Diff Detail

Repository
rL LLVM

Event Timeline

rampitec created this revision.Jul 24 2018, 3:00 PM
artem.tamazov accepted this revision.Jul 25 2018, 7:47 AM

LGTM

test/CodeGen/AMDGPU/divrem24-assume.ll
31 ↗(On Diff #157137)

The usage of the Answer to the Ultimate Question of Life, the Universe, and Everything to resolve the divrem problem is definitely overkill ))

This revision is now accepted and ready to land.Jul 25 2018, 7:47 AM
This revision was automatically updated to reflect the committed changes.