Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[TIR] Add a new intrinsic count leading zeros for LLVM and SPIR-V #7825

Merged
merged 4 commits into from
Apr 16, 2021

Conversation

masahi
Copy link
Member

@masahi masahi commented Apr 12, 2021

This adds a new intrinsic clz to LLVM and SPIRV per discussion in #7669 (comment).

The motivation is to compute integer ceil(log2(x)) without using log2, since SPIR-V doesn't support log2 on fp64. Using this new intrinsic, we can compute integer ceil(log2(x)) as follows:

def ceil_log2(x):
    clz = tvm.tir.clz(x)
    bits = int(x.dtype[-2:])
    ceil_log2 = tvm.tir.if_then_else(x & (x - 1) == 0, bits - clz - 1, bits - clz)
    return ceil_log2

LLVM has ctlz intrinsic which we can use directly. SPIRV doesn't have clz intrin but it does have FindUMsb.
https://llvm.org/docs/LangRef.html#llvm-ctlz-intrinsic
https://www.khronos.org/registry/spir-v/specs/1.0/GLSL.std.450.html

Some discussion points:

  • The result is defined only for positive integers.
  • For now I didn't support clz on zero, because the LLVM doc above suggests clz(0) is not well supported by backends. See is_zero_undef argument in the llvm doc for more details.
  • Do we want to support clz on negative integers?
  • What about floating point?

@tqchen @junrushao1994 @vinx13 @zxybazh

src/target/intrin_rule.cc Outdated Show resolved Hide resolved
@masahi
Copy link
Member Author

masahi commented Apr 16, 2021

Good to go? Or any comments on the discussion points above @tqchen @junrushao1994 @vinx13 @zxybazh

Copy link
Member

@junrushao junrushao left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM :-)

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants