cxx122 opened a new issue, #15196:
URL: https://github.com/apache/tvm/issues/15196
`te.comm_reducer` does not support asymmetric operations such as division
and subtraction. If use these operation, it will cause inconsistent problem.
```
def te_test():
mysub = te.comm_reducer(lambda x, y: x - y,
lambda t: tvm.tir.const(0, dtype=t), name="mydiv")
A = te.placeholder((64, 64), name="A")
k = te.reduce_axis((0, 64), name="k")
B = te.compute((64,), lambda i: mysub(A[i, k], axis=k), name="B")
return [A, B]
```
### Expected behavior
Same results.
### Actual behavior
Different results.
Before Schedule
```
@main = primfn(A_1: handle, B_1: handle) -> ()
attr = {"from_legacy_te_schedule": True, "global_symbol": "main",
"tir.noalias": True}
buffers = {A: Buffer(A_2: Pointer(float32), float32, [4096], []),
B: Buffer(B_2: Pointer(float32), float32, [64], [])}
buffer_map = {A_1: A, B_1: B}
preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [64, 64], []),
B_1: B_3: Buffer(B_2, float32, [64], [])} {
for (i: int32, 0, 64) {
B[i] = 0f32
for (k: int32, 0, 64) {
B[i] = (B[i] - A[((i*64) + k)])
}
}
}
```
After Schedule
```
@main = primfn(A_1: handle, B_1: handle) -> ()
attr = {"from_legacy_te_schedule": True, "global_symbol": "main",
"tir.noalias": True}
buffers = {A: Buffer(A_2: Pointer(float32), float32, [4096], []),
B: Buffer(B_2: Pointer(float32), float32, [64], [])}
buffer_map = {A_1: A, B_1: B}
preflattened_buffer_map = {A_1: A_3: Buffer(A_2, float32, [64, 64], []),
B_1: B_3: Buffer(B_2, float32, [64], [])} {
allocate(B.rf: Pointer(global float32), float32, [2048]), storage_scope =
global {
for (i: int32, 0, 64) "parallel" {
for (k.inner.init: int32, 0, 32) {
B.rf_1: Buffer(B.rf, float32, [2048], [])[((i*32) + k.inner.init)] =
0f32
}
for (k.outer: int32, 0, 2) {
for (k.inner: int32, 0, 32) {
let cse_var_1: int32 = ((i*32) + k.inner)
B.rf_1[cse_var_1] = (B.rf_1[cse_var_1] - A[(((i*64) +
(k.outer*32)) + k.inner)])
}
}
}
for (ax0: int32, 0, 64) "parallel" {
B[ax0] = 0f32
for (k.inner.v: int32, 0, 32) {
B[ax0] = (B[ax0] - B.rf_1[((ax0*32) + k.inner.v)])
}
}
}
}
```
### Environment
Operating System: Ubuntu 18.04
TVM version: v0.10.dev0
### Steps to reproduce
```
import tvm
import random
import numpy as np
from tvm import te
from tvm import tir
from tvm import testing
from tvm import auto_scheduler
from tvm.auto_scheduler.workload_registry import register_workload_tensors
POLICY_PARAMS = {
"eps_greedy": 0.05,
"retry_search_one_round_on_empty": 1,
"sample_init_min_population": 3,
"sample_init_use_measured_ratio": 0.2,
"evolutionary_search_population": 5,
"evolutionary_search_num_iters": 4,
"evolutionary_search_mutation_prob": 0.85,
"cpu_multi_level_tiling_structure": "SSRSRS",
"gpu_multi_level_tiling_structure": "SSSRRSRS",
# Notice: the default thread bind policy of GPU assumes the tiling
structure to have at
# least 3 spatial tiling levels in outermost
"max_innermost_split_factor": 64,
"max_vectorize_size": 16,
"disable_change_compute_location": 0,
}
# Division.
# def te_test():
# mydiv = te.comm_reducer(lambda x, y: tir.div(x, y),
# lambda t: tvm.tir.const(0, dtype=t), name="mydiv")
# A = te.placeholder((64, 64), name="A")
# k = te.reduce_axis((0, 64), name="k")
# B = te.compute((64,), lambda i: mydiv(A[i, k], axis=k), name="B")
# return [A, B]
# Subtraction
def te_test():
mysub = te.comm_reducer(lambda x, y: x - y,
lambda t: tvm.tir.const(0, dtype=t), name="mydiv")
A = te.placeholder((64, 64), name="A")
k = te.reduce_axis((0, 64), name="k")
B = te.compute((64,), lambda i: mysub(A[i, k], axis=k), name="B")
return [A, B]
# Get dag and print it.
tensors = te_test()
dag = auto_scheduler.ComputeDAG(tensors)
# Get inputs.
inputs = []
for tensor in dag.tensors:
shape = [x.value if 'value' in dir(x) and isinstance(x.value, int) else
1 for x in tensor.shape]
inputs.append((2 * np.random.uniform(size=shape)+1).astype(tensor.dtype))
# Get program with no schedule.
results = []
mod_list = []
pre_schedule_list = dag.apply_steps_from_state(dag.get_init_state())
pre_mod = tvm.lower(pre_schedule_list[0], pre_schedule_list[1],
simple_mode=True)
mod_list.append(pre_mod)
with tvm.transform.PassContext(opt_level=0):
mod_exec = tvm.build(pre_mod)
new_inputs = [tvm.nd.array(x, tvm.cpu()) for x in inputs.copy()]
mod_exec(*new_inputs)
result = []
for x in new_inputs:
try:
result.append(x.numpy() if isinstance(
x, tvm.runtime.NDArray) else x)
except (ValueError, tvm.TVMError):
result.append(None)
results.append(result)
# Get program with schedule.
register_workload_tensors(dag.workload_key(), tensors)
task = auto_scheduler.SearchTask(workload_key=dag.workload_key(),
target=tvm.target.Target("llvm"))
policy = auto_scheduler.SketchPolicy(task, verbose=0, params=POLICY_PARAMS)
states = policy.sample_initial_population()
for state in states:
schedule_list = dag.apply_steps_from_state(state)
mod = tvm.lower(schedule_list[0], schedule_list[1], simple_mode=True)
mod_list.append(mod)
with tvm.transform.PassContext(opt_level=0):
mod_exec = tvm.build(mod)
new_inputs = [tvm.nd.array(x, tvm.cpu()) for x in inputs.copy()]
mod_exec(*new_inputs)
result = []
for x in new_inputs:
try:
result.append(x.numpy() if isinstance(
x, tvm.runtime.NDArray) else x)
except (ValueError, tvm.TVMError):
result.append(None)
results.append(result)
for i in range(1, len(results)):
result = results[i]
for compare_idex in [1]:
try:
tvm.testing.assert_allclose(results[0][compare_idex],
result[compare_idex], rtol=1e-5)
except AssertionError as e:
print(e)
print(mod_list[0])
print(mod_list[i])
break
```
### Triage
* tune:auto_scheduler
* tir:schedule
* tir:transform
--
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]
