René's URL Explorer Experiment

{"@context":"https://schema.org","@type":"DiscussionForumPosting","headline":"InternalError: Check failed: func-\u003ebuffer_map.size() == 0 (3 vs. 0) : This pass must be called after MakePackedAPI","articleBody":"When I used bitblas to tune a matmul operator on the A100, I encountered the following error.\n![Image](https://github.com/user-attachments/assets/c7548aed-9076-43ed-b33a-b70a369c70a6)\n\nEnvironment Configuration:\nUbuntu 22.04.2 LTS\nbitblas: 2bd1deeb7977ea27019c8acd3d926dec9ada9eed\ncuda12.1\n\nThe specific test script is as follows:\n```\nfrom bitblas.gpu.matmul_analysis import get_tensorized_func_and_tags\nfrom bitblas.base.roller.policy import TensorCorePolicy, DefaultPolicy\nfrom bitblas.base.arch import CUDA\nfrom bitblas.base.utils import apply_and_build\nimport tvm\nfrom tvm.script import tir as T\nimport bitblas\n\n@tvm.script.ir_module\nclass FusedSingleOp:\n    @T.prim_func(private=True)\n    def dense1(lv11: T.Buffer((T.int64(1024), T.int64(1024)), \"float16\"), B: T.Buffer((T.int64(1024), T.int64(1024)), \"float16\"), T_matmul_NT: T.Buffer((T.int64(1024), T.int64(1024)), \"float16\")):\n        T.func_attr({\"layout_free_buffers\": [1], \"op_attrs\": {\"op_name\": \"nn.dense\", \"out_dtype\": \"float16\", \"units\": None}, \"op_pattern\": 4, \"tir.noalias\": T.bool(True)})\n        # with T.block(\"root\"):\n        for i0, i1, k in T.grid(T.int64(1024), T.int64(1024), T.int64(1024)):\n            with T.block(\"T_matmul_NT\"):\n                v_i0, v_i1, v_k = T.axis.remap(\"SSR\", [i0, i1, k])\n                T.reads(lv11[v_i0, v_k], B[v_i1, v_k])\n                T.writes(T_matmul_NT[v_i0, v_i1])\n                with T.init():\n                    T_matmul_NT[v_i0, v_i1] = T.float16(0)\n                T_matmul_NT[v_i0, v_i1] = T_matmul_NT[v_i0, v_i1] + lv11[v_i0, v_k] * B[v_i1, v_k]\n\nir_module = FusedSingleOp\nfunc = ir_module[\"dense1\"]\ntarget = tvm.target.Target(\"cuda\")\narch = CUDA(target)\npolicy = DefaultPolicy(func=func, arch=arch)\ntry:\n    tensorized_func, tags = get_tensorized_func_and_tags(func, arch.target)\nexcept Exception:\n    tags = None\n# Tune with Tensor Core if possible\nif tags:\n    policy = TensorCorePolicy(func=tensorized_func, arch=arch, tags=tags)\n\nconfigs = policy.emit_config(topk=20)\n\ncpresults, best = apply_and_build(func, configs, arch, parallel_build=False)\n\nprint(best.code)\n```\n\n\nWhen I tested the same script on the V100 with same Environment Configuration, the error`InternalError: Check failed: func-\u003ebuffer_map.size() == 0 (3 vs. 0) : This pass must be called after MakePackedAPI`did not occur. Instead, it resulted in:\n![Image](https://github.com/user-attachments/assets/f0342e51-97a3-46e9-9481-bcffb6a5bf10)\n\nThe specific test script is as follows:\n```\nfrom bitblas.gpu.matmul_analysis import get_tensorized_func_and_tags\nfrom bitblas.base.roller.policy import TensorCorePolicy, DefaultPolicy\nfrom bitblas.base.arch import CUDA\nfrom bitblas.base.utils import apply_and_build\nimport tvm\nfrom tvm.script import tir as T\nimport bitblas\n\n@tvm.script.ir_module\nclass FusedSingleOp:\n    @T.prim_func(private=True)\n    def dense1(lv11: T.Buffer((T.int64(1024), T.int64(1024)), \"float16\"), B: T.Buffer((T.int64(1024), T.int64(1024)), \"float16\"), T_matmul_NT: T.Buffer((T.int64(1024), T.int64(1024)), \"float16\")):\n        T.func_attr({\"layout_free_buffers\": [1], \"op_attrs\": {\"op_name\": \"nn.dense\", \"out_dtype\": \"float16\", \"units\": None}, \"op_pattern\": 4, \"tir.noalias\": T.bool(True)})\n        # with T.block(\"root\"):\n        for i0, i1, k in T.grid(T.int64(1024), T.int64(1024), T.int64(1024)):\n            with T.block(\"T_matmul_NT\"):\n                v_i0, v_i1, v_k = T.axis.remap(\"SSR\", [i0, i1, k])\n                T.reads(lv11[v_i0, v_k], B[v_i1, v_k])\n                T.writes(T_matmul_NT[v_i0, v_i1])\n                with T.init():\n                    T_matmul_NT[v_i0, v_i1] = T.float16(0)\n                T_matmul_NT[v_i0, v_i1] = T_matmul_NT[v_i0, v_i1] + lv11[v_i0, v_k] * B[v_i1, v_k]\n\nir_module = FusedSingleOp\nfunc = ir_module[\"dense1\"]\ntarget = tvm.target.Target(\"cuda\")\narch = CUDA(target)\npolicy = DefaultPolicy(func=func, arch=arch)\ntry:\n    tensorized_func, tags = get_tensorized_func_and_tags(func, arch.target)\nexcept Exception:\n    tags = None\n# Tune with Tensor Core if possible\nif tags:\n    policy = TensorCorePolicy(func=tensorized_func, arch=arch, tags=tags)\n\nconfigs = policy.emit_config(topk=20)\n\ncpresults, best = apply_and_build(func, configs, arch, parallel_build=False)\n\nprint(best.code)\n```","author":{"url":"https://github.com/Cunxiao2002","@type":"Person","name":"Cunxiao2002"},"datePublished":"2024-12-08T13:36:14.000Z","interactionStatistic":{"@type":"InteractionCounter","interactionType":"https://schema.org/CommentAction","userInteractionCount":4},"url":"https://github.com/259/BitBLAS/issues/259"}