symbol.py 3.6 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101
  1. # mypy: allow-untyped-defs
  2. """
  3. This file contains canonical definitions for our symbol naming conventions,
  4. across torch.fx.experimental.symbolic_shapes and torch._inductor. The
  5. intention is:
  6. 1. To make it easily greppable where all the sites we use a prefix are
  7. 2. Make it possible to easily tell if we can introduce a new prefix without
  8. introducing a conflict
  9. You can occasionally test if prefixes have been hardcoded by renaming prefixes
  10. in this file and seeing what breaks.
  11. """
  12. from collections.abc import Iterable
  13. from enum import auto, Enum
  14. from typing import Union
  15. import sympy
  16. class SymT(Enum):
  17. SIZE = auto()
  18. FLOAT = auto()
  19. UNBACKED_INT = auto()
  20. UNBACKED_FLOAT = auto()
  21. # Inductor: The intermediates in inner_fn tmp0, one generated per ops call.
  22. # If one of these shows up in an indexing expression, that means an
  23. # indirect load is happening.
  24. TMP = auto()
  25. # Inductor: Placeholder variable that is later replaced with TMP
  26. INDIRECT = auto()
  27. # Inductor: Some size expressions are replaced with a precomputed size ps0
  28. # which is computed host side, and then directly reused in the kernel, so
  29. # we don't repeatedly recompute it on device.
  30. PRECOMPUTED_SIZE = auto()
  31. # Inductor: An indexing variable i0 in loops IR which ranges over non-reduced
  32. # dim in the loop
  33. INDEX = auto()
  34. # Inductor: A reduction indexing (r0, r1) variables in loops IR which ranges over
  35. # reduced dim(s) in the loop
  36. R0_INDEX = auto()
  37. R1_INDEX = auto()
  38. # Inductor: In templated kernels torch._inductor.kernel, we have a hook to
  39. # store the final output and append epilogue fusions. To do this, we must
  40. # know what the indexes the outputs range over. NB: These will also
  41. # advertise as INDEX, this is... probably OK?
  42. TEMPLATE_INDEX = auto()
  43. # Inductor: iteration domain for blockIdx.x/blockIdx.y
  44. XBLOCK = auto()
  45. YBLOCK = auto()
  46. ZBLOCK = auto()
  47. # Inductor: this is used solely for dynamic_reshape_indexer
  48. VIEW = auto()
  49. # Alternate (non-modular) indexing used in halide kernels
  50. HALIDE = auto()
  51. # Invariant: there must not be a prefix which is a prefix of another string,
  52. # as this introduces ambiguity
  53. prefix_str = {
  54. SymT.SIZE: "s", # integer
  55. SymT.UNBACKED_INT: "u", # integer
  56. # Prefix z here is chosen to avoid false aliasing in symbol_is_type test
  57. # DO NOT add a "z" type. You also need to avoid conflicts on these
  58. # prefixes but this is somewhat easier to manage
  59. SymT.FLOAT: "zf",
  60. SymT.UNBACKED_FLOAT: "zuf",
  61. SymT.TMP: "tmp",
  62. SymT.PRECOMPUTED_SIZE: "ps",
  63. SymT.INDEX: "i",
  64. SymT.R0_INDEX: "r0_",
  65. SymT.R1_INDEX: "r1_",
  66. SymT.TEMPLATE_INDEX: "idx",
  67. SymT.XBLOCK: "x",
  68. SymT.YBLOCK: "y",
  69. SymT.ZBLOCK: "z",
  70. SymT.INDIRECT: "indirect", # false aliasing?
  71. SymT.VIEW: "view",
  72. SymT.HALIDE: "h",
  73. }
  74. def make_symbol(prefix: SymT, idx: int, **kwargs) -> sympy.Symbol:
  75. # TODO: maybe put the assumptions here directly
  76. return sympy.Symbol(f"{prefix_str[prefix]}{idx}", **kwargs)
  77. # This type is a little wider than it should be, because free_symbols says
  78. # that it contains Basic, rather than Symbol
  79. def symbol_is_type(sym: sympy.Basic, prefix: Union[SymT, Iterable[SymT]]) -> bool:
  80. assert isinstance(sym, sympy.Symbol)
  81. name_str = sym.name.lower() # Match capitalized names like XBLOCK, RBLOCK
  82. if isinstance(prefix, SymT):
  83. return name_str.startswith(prefix_str[prefix])
  84. else:
  85. return name_str.startswith(tuple(prefix_str[p] for p in prefix))
  86. def free_symbol_is_type(e: sympy.Expr, prefix: Union[SymT, Iterable[SymT]]) -> bool:
  87. return any(symbol_is_type(v, prefix) for v in e.free_symbols)