Revert "[Dynamic][Enhancement] Convert div and mod including symbolva…

…rs to fast int div/mod" (#463) Reverts CentML/hidet#405
hidet-org · Dec 19, 2024 · e251184 · e251184
1 parent 4d3c48f
commit e251184
Show file tree

Hide file tree

Showing 12 changed files with 4 additions and 427 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -25,7 +25,6 @@ add_library(hidet_runtime SHARED
         src/hidet/runtime/callbacks.cpp
         src/hidet/runtime/logging.cpp
         src/hidet/runtime/symbols.cpp
-        src/hidet/runtime/int_fastdiv.cpp
         src/hidet/runtime/llm/tokenizer/decoders.cpp
         src/hidet/runtime/llm/tokenizer/models.cpp
         src/hidet/runtime/llm/tokenizer/normalizers.cpp

diff --git a/include/hidet/runtime/int_fastdiv.h b/include/hidet/runtime/int_fastdiv.h
diff --git a/python/hidet/backend/codegen.py b/python/hidet/backend/codegen.py
@@ -712,7 +712,7 @@ def require_headers(self) -> Doc:
         doc += Text('#include <hidet/runtime/cuda/complex.h>') + NewLine()
         doc += Text('#include <hidet/runtime/cuda/context.h>') + NewLine()
         doc += Text("#include <hidet/runtime/logging.h>") + NewLine()
-        doc += Text('#include <hidet/runtime/int_fastdiv.h>') + NewLine()
+
         for header in self.ir_module.include_headers:
             doc += Text('#include <{}>').format(header) + NewLine()
 
@@ -801,7 +801,6 @@ def require_headers(self) -> Doc:
         doc += Text('#include <hidet/runtime/cpu/context.h>') + NewLine()
         doc += Text('#include <hidet/runtime/cpu/float32.h>') + NewLine()
         doc += Text("#include <hidet/runtime/logging.h>") + NewLine()
-        doc += Text('#include <hidet/runtime/int_fastdiv.h>') + NewLine()
 
         if self.require_complex:
             doc += Text('#include <hidet/runtime/cpu/complex.h>') + NewLine()

diff --git a/python/hidet/ir/primitives/cuda/fastintdiv.py b/python/hidet/ir/primitives/cuda/fastintdiv.py
diff --git a/python/hidet/ir/primitives/runtime.py b/python/hidet/ir/primitives/runtime.py
@@ -56,11 +56,6 @@ def register_functions():
     register_primitive_function(
         name='get_nccl_comm', func_or_type=FuncType([int32], void_p), codegen_name='get_nccl_comm'
     )
-    register_primitive_function(
-        name='calculate_magic_numbers',
-        func_or_type=FuncType([int32, int32, int32, int32], void_p),
-        codegen_name='calculate_magic_numbers',
-    )
 
 
 def get_cuda_stream() -> void_p:
@@ -101,7 +96,3 @@ def memory_planner_used(idx: Union[int, Expr]):
 
 def get_nccl_comm(idx: int) -> void_p:
     return call_primitive_func('get_nccl_comm', [idx])
-
-
-def calculate_magic_numbers(divisor: int, m: int, s: int, ads: int):
-    return call_primitive_func('calculate_magic_numbers', [divisor, m, s, ads])
diff --git a/python/hidet/lang/cuda.py b/python/hidet/lang/cuda.py
@@ -26,7 +26,6 @@
 from hidet.ir.primitives.cuda.time import nano_sleep
 from hidet.ir.primitives.cuda.memcpy import memcpy_async, memcpy
 from hidet.ir.primitives.cuda.atomic import atomic_add, atomic_sub, atomic_min, atomic_max, atomic_exchange, atomic_cas
-from hidet.ir.primitives.cuda.fastintdiv import fast_intdiv, fast_intmod
 from hidet.ir.primitives.cuda.shfl import shfl_sync, shfl_up_sync, shfl_xor_sync, shfl_down_sync
 from hidet.ir.primitives.cuda.mutex import acquire_lock, release_lock, acquire_seq_semaphore, release_seq_semaphore
 from hidet.lang.constructs.declare import register_tensor, shared_tensor

diff --git a/python/hidet/testing/torch_utils.py b/python/hidet/testing/torch_utils.py
@@ -84,7 +84,6 @@ def init_hidet(self):
         # hidet.option.debug_cache_tuning(True)
         # hidet.option.save_lower_ir(True)
         # hidet.option.debug_show_verbose_flow_graph(True)
-        # hidet.torch.dynamo_config.dump_graph_ir("./graph_ir")
 
         # Initialise compiler server
         if os.environ.get('CI_CS_HOSTNAME'):

diff --git a/python/hidet/transforms/__init__.py b/python/hidet/transforms/__init__.py
@@ -16,7 +16,6 @@
 from .instruments import PassInstrument, SaveIRInstrument, ProfileInstrument
 
 from .unify_global_objects import unify_global_objects_pass
-from .convert_div_to_fastintdiv import convert_div_to_fastintdiv_pass
 from .flatten_tensor_slice import flatten_tensor_slice_pass
 from .flatten_tensor_index import flatten_tensor_index_pass
 from .generate_launch_func import generate_launch_func_pass
@@ -93,8 +92,6 @@ def lower(ir_module: IRModule) -> IRModule:
         add_explicit_cast_pass(),
         declare_to_let_pass(),
         instantiate_symbols_pass(),
-        convert_div_to_fastintdiv_pass(),
-        import_primitive_functions_pass(),
         check_launch_configuration_pass(),
         # simplification
         expand_let_expr_pass(),

diff --git a/python/hidet/transforms/convert_div_to_fastintdiv.py b/python/hidet/transforms/convert_div_to_fastintdiv.py