LeiWang1999 commented on PR #14901:
URL: https://github.com/apache/tvm/pull/14901#issuecomment-1598863567
looks like after this pr was merged, the ret val is
code to reproduce:
```python
import tvm
from tvm import te
import numpy as np
import tvm.testing
from tvm.script import tir as T
from tvm.tir import TensorIntrin
M = 64
N = 64
@tvm.script.ir_module
class MyModule:
@T.prim_func
def main(a: T.handle, b: T.handle):
T.func_attr({"global_symbol": "main"})
A = T.match_buffer(a, (M, N), dtype="float32")
B = T.match_buffer(b, (M, N), dtype="float32")
for i, j in T.grid(M, N):
with T.block("B"):
vi, vj = T.axis.remap("SS", [i, j])
B[vi, vj] = A[vi, vj] * 2.0
ir_module = MyModule
sch = tvm.tir.Schedule(ir_module, debug_mask="all")
block_b = sch.get_block("B")
i, j = sch.get_loops(block_b)
sch.bind(i, "blockIdx.x")
sch.bind(j, "threadIdx.x")
print(sch.mod["main"].script())
ctx = tvm.rocm(0)
with tvm.transform.PassContext():
rocm_mod = tvm.build(sch.mod, target="rocm")
'''
Traceback (most recent call last):
File "../../tvm_rocm/memory_copy.py", line 38, in <module>
rocm_mod = tvm.build(sch.mod, target="rocm -mcpu=gfx90a")
File "/home/aiscuser/v-leiwang3/tvm/python/tvm/driver/build_module.py",
line 281, in build
rt_mod_host = _driver_ffi.tir_to_runtime(annotated_mods, target_host)
File
"/home/aiscuser/v-leiwang3/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line
238, in __call__
raise get_last_ffi_error()
tvm._ffi.base.TVMError: Traceback (most recent call last):
7: TVMFuncCall
6:
tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::runtime::Module
(tvm::runtime::Map<tvm::Target, tvm::IRModule, void, void> const&,
tvm::Target)>::AssignTypedLambda<tvm::{lambda(tvm::runtime::Map<tvm::Target,
tvm::IRModule, void, void> const&,
tvm::Target)#6}>(tvm::{lambda(tvm::runtime::Map<tvm::Target, tvm::IRModule,
void, void> const&, tvm::Target)#6}, std::__cxx11::basic_string<char,
std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs
const&, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj
const*, std::__cxx11::basic_string<char, std::char_traits<char>,
std::allocator<char> >, tvm::runtime::TVMRetValue)
5: tvm::TIRToRuntime(tvm::runtime::Map<tvm::Target, tvm::IRModule, void,
void> const&, tvm::Target const&)
4: tvm::codegen::Build(tvm::IRModule, tvm::Target)
3:
tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::TypedPackedFunc<tvm::runtime::Module
(tvm::IRModule, tvm::Target)>::AssignTypedLambda<tvm::runtime::Module
(*)(tvm::IRModule, tvm::Target)>(tvm::runtime::Module (*)(tvm::IRModule,
tvm::Target), std::__cxx11::basic_string<char, std::char_traits<char>,
std::allocator<char> >)::{lambda(tvm::runtime::TVMArgs const&,
tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*,
tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)
2: tvm::codegen::BuildAMDGPU(tvm::IRModule, tvm::Target)
1: tvm::codegen::CodeGenLLVM::Finish()
0: tvm::codegen::CodeGenLLVM::Verify() const
File "/home/aiscuser/v-leiwang3/tvm/src/target/llvm/codegen_llvm.cc", line
361
TVMError: LLVM module verification failed with the following errors:
Calling convention requires void return type
ptr @main_kernel
Function return type does not match operand type of return inst!
ret void
i32
'''
```
this can be resolved through:
```c++
src/target/llvm/codegen_amdgpu.cc
class CodeGenAMDGPU : public CodeGenLLVM {
public:
CodeGenAMDGPU() = default;
virtual ~CodeGenAMDGPU() = default;
llvm::Function* DeclareFunction(const GlobalVar& gvar, const PrimFunc& f) {
// amd gpu kernel function should be void return
return this->DeclareFunctionInternal(gvar, f, true);
}
```
because amd gpu kernels should be void return but by default it wll call a
llvm codegen implementation which returns int by default, Previously there were
no issues, until this code snippet was added.
```c++
// src/target/llvm/codegen_llvm.cc
if (auto it = functions_.find(gvar.get()); it != functions_.end()) {
return it->second;
}
```
and if we wanna leverage shared memory, there's another issue:
```
TVMError: LLVM module verification failed with the following errors:
Global is external, but doesn't have external or weak linkage!
ptr addrspace(3) @shmem
Global is external, but doesn't have external or weak linkage!
ptr addrspace(3) @shmem.1
```
Please take a look @Lunderberg
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
