[CodeGenC] Handle GlobalVar callee as internal function call

[Lunderberg]

Analogous to #14901, treat GlobalVar callees as internal function calls in CodeGenC. This specific PR doesn't provide new end-to-end functionality, as the target="c" backend isn't compiled. It does lead into allowing subroutines in any target whose codegen derives from CodeGenC, which will depend on the single-module lowering flow in #14985.

[tvm-bot]

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

• No users to tag found in teams: codegenc _{See #10317 for details}

_{Generated by tvm-bot}

[csullivan]

Approved with only a few questions and nitpicks.

[MasterJH5574]

Hello @Lunderberg, unfortunately this PR/commit completely breaks our Metal codegen.

You are able to reproduce this issue by following the steps on Mac below:

• open tests/python/unittest/test_allreduce_cuda.py (I just picked an arbitrary test for CUDA),
• replace all occurrences of cuda with metal,
• run this test_allreduce_cuda.py.

Then you will see the following error:

~/W/tvm/t/p/unittest @4d7e93cd *1 !2 ❯ python test_allreduce_cuda.py                                 python310 ruihang-macstudio@GS22385 08:12:05 PM
Traceback (most recent call last):
  File "/Users/ruihang-macstudio/Workspace/tvm/tests/python/unittest/test_allreduce_cuda.py", line 104, in <module>
    test_allreduce_metal()
  File "/Users/ruihang-macstudio/Workspace/tvm/tests/python/unittest/test_allreduce_cuda.py", line 99, in test_allreduce_metal
    check_sum(d1, d2, d3)
  File "/Users/ruihang-macstudio/Workspace/tvm/tests/python/unittest/test_allreduce_cuda.py", line 69, in check_sum
    f(a, b)
  File "/Users/ruihang-macstudio/Workspace/tvm/python/tvm/runtime/module.py", line 201, in __call__
    return self.entry_func(*args)
  File "/Users/ruihang-macstudio/Workspace/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line 239, in __call__
    raise_last_ffi_error()
  File "/Users/ruihang-macstudio/Workspace/tvm/python/tvm/_ffi/base.py", line 476, in raise_last_ffi_error
    raise py_err
  File "/Users/ruihang-macstudio/Workspace/tvm/src/runtime/library_module.cc", line 76, in tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::$_0::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const
    ICHECK_EQ(ret, 0) << TVMGetLastError();
tvm.error.InternalError: Traceback (most recent call last):
  0: tvm::runtime::WrapPackedFunc(int (*)(TVMValue*, int*, int, TVMValue*, int*, void*), tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::$_0::operator()(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) const
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/library_module.cc:76
  6: tvm::runtime::ModuleNode::GetFuncFromEnv(tvm::runtime::String const&)
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/module.cc:114
  5: tvm::runtime::ModuleNode::GetFunction(tvm::runtime::String const&, bool)
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/module.cc:66
  4: tvm::runtime::MetalModuleNode::GetFunction(tvm::runtime::String const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/metal/metal_module.mm:246
  3: void tvm::runtime::metal::AutoReleasePoolWrapper::operator<<<tvm::runtime::MetalModuleNode::GetFunction(tvm::runtime::String const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::$_0>(tvm::runtime::MetalModuleNode::GetFunction(tvm::runtime::String const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::$_0 const&)
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/metal/metal_common.h:89
  2: tvm::runtime::MetalModuleNode::GetFunction(tvm::runtime::String const&, tvm::runtime::ObjectPtr<tvm::runtime::Object> const&)::$_0::operator()() const
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/metal/metal_module.mm:257
  1: tvm::runtime::MetalWrappedFunc::Init(tvm::runtime::MetalModuleNode*, tvm::runtime::ObjectPtr<tvm::runtime::Object>, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char> > const&, unsigned long, unsigned long, std::__1::vector<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char> >, std::__1::allocator<std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char> > > > const&)
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/metal/metal_module.mm:187
  0: tvm::runtime::MetalModuleNode::GetPipelineState(unsigned long, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char> > const&)
        at /Users/ruihang-macstudio/Workspace/tvm/src/runtime/metal/metal_module.mm:109
  File "/Users/ruihang-macstudio/Workspace/tvm/src/runtime/metal/metal_module.mm", line 109
  File "/Users/ruihang-macstudio/Workspace/tvm/src/runtime/library_module.cc", line 76
InternalError: Check failed: ret == 0 (-1 vs. 0) : TVMError: Fail to compile metal source:program_source:10:1: error: expected parameter declarator
;
^
program_source:10:1: error: expected ')'
program_source:8:36: note: to match this '('
kernel void default_function_kernel(  device float* A [[ buffer(0) ]],
                                   ^
program_source:8:58: warning: 'buffer' attribute ignored on function declaration
kernel void default_function_kernel(  device float* A [[ buffer(0) ]],
                                                         ^
program_source:9:22: warning: 'buffer' attribute ignored on function declaration
  device float* B [[ buffer(1) ]],
                     ^
program_source:11:20: error: 'threadgroup_position_in_grid' attribute only applies to parameters
  uint3 blockIdx [[threadgroup_position_in_grid]],
                   ^
program_source:11:9: error: program scope variable must reside in constant address space
  uint3 blockIdx [[threadgroup_position_in_grid]],
        ^
program_source:11:50: error: expected ';' at end of declaration
  uint3 blockIdx [[threadgroup_position_in_grid]],
                                                 ^
                                                 ;
program_source:12:21: error: 'thread_position_in_threadgroup' attribute only applies to parameters
  uint3 threadIdx [[thread_position_in_threadgroup]]
                    ^
program_source:12:9: error: program scope variable must reside in constant address space
  uint3 threadIdx [[thread_position_in_threadgroup]]
        ^
program_source:12:53: error: expected ';' after top level declarator
  uint3 threadIdx [[thread_position_in_threadgroup]]
                                                    ^
                                                    ;

It says that the generated Metal code cannot be parsed.

By reverting this commit, I am able to get this test run (though it will fail on numerical comparison due to precision issue, which is a nonissue for here).

Could you make a fix for this issue? I appreciate it if the problem can get resolved shortly. Thank you!

[junrushao]

We might need some basic testings for Metal to prevent this issue from happening...For now, to quickly ensure that MLC LLM is not impacted, @MasterJH5574 shall we submit a PR to temporarily revert this commit? Thanks a lot!