Skip to content

T
tic
Unverified Commit a5d7bdab by Haichen Shen Committed by GitHub
Options

Adjust strategy plevel to achieve expected performance by default (#5118)

parent d2bc94d9
Showing with 19 additions and 21 deletions
python/tvm/relay/op/strategy/arm_cpu.py
......@@ -67,13 +67,13 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
wrap_compute_conv2d(topi.arm_cpu.conv2d_nchw_winograd),
wrap_topi_schedule(topi.arm_cpu.schedule_conv2d_nchw_winograd),
name="conv2d_nchw_winograd.arm_cpu",
plevel=15)
plevel=5)
if "nnpack" in target.libs and pt == 1 and pb == 1 and pl == 1 and pr == 1:
strategy.add_implementation(
wrap_compute_conv2d(topi.arm_cpu.conv2d_nchw_winograd_nnpack),
wrap_topi_schedule(topi.arm_cpu.schedule_conv2d_nchw_winograd_nnpack),
name="conv2d_nchw_winograd_nnpack.arm_cpu",
plevel=13)
plevel=15)
elif re.match(r"OIHW\d*o", kernel_layout):
strategy.add_implementation(
wrap_compute_conv2d(topi.arm_cpu.conv2d_nchw_spatial_pack),
......@@ -177,7 +177,7 @@ def conv2d_winograd_without_weight_transfrom_strategy_arm_cpu(attrs, inputs, out
wrap_topi_schedule(
topi.arm_cpu.schedule_conv2d_nchw_winograd_nnpack_without_weight_transform),
name="conv2d_nchw_winograd_nnpack_withou_weight_transform.arm_cpu",
plevel=5)
plevel=15)
else:
raise RuntimeError("Unsupported kernel shape: {}".format(kernel.shape))
else:
......
python/tvm/relay/op/strategy/bifrost.py
......@@ -50,7 +50,7 @@ def conv2d_strategy_bifrost(attrs, inputs, out_type, target):
wrap_compute_conv2d(topi.bifrost.conv2d_nchw_winograd),
wrap_topi_schedule(topi.bifrost.schedule_conv2d_nchw_winograd),
name="conv2d_nchw_winograd.bifrost",
plevel=15)
plevel=5)
elif re.match(r"OIHW\d*o", kernel_layout):
strategy.add_implementation(
wrap_compute_conv2d(topi.bifrost.conv2d_nchw_spatial_pack),
......
python/tvm/relay/op/strategy/cuda.py
......@@ -135,7 +135,7 @@ def conv2d_strategy_cuda(attrs, inputs, out_type, target):
wrap_compute_conv2d(topi.cuda.conv2d_cudnn, True),
wrap_topi_schedule(topi.cuda.schedule_conv2d_cudnn),
name="conv2d_cudnn.cuda",
plevel=5)
plevel=15)
elif is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups):
if layout == "NCHW":
assert kernel_layout == "OIHW"
......@@ -295,13 +295,13 @@ def dense_strategy_cuda(attrs, inputs, out_type, target):
wrap_compute_dense(topi.cuda.dense_large_batch),
wrap_topi_schedule(topi.cuda.schedule_dense_large_batch),
name="dense_large_batch.cuda",
plevel=15)
plevel=5)
if target.target_name == "cuda" and "cublas" in target.libs:
strategy.add_implementation(
wrap_compute_dense(topi.cuda.dense_cublas),
wrap_topi_schedule(topi.cuda.schedule_dense_cublas),
name="dense_cublas.cuda",
plevel=20)
plevel=15)
return strategy
@batch_matmul_strategy.register(["cuda", "gpu"])
......
python/tvm/relay/op/strategy/mali.py
......@@ -49,7 +49,7 @@ def conv2d_strategy_mali(attrs, inputs, out_type, target):
wrap_compute_conv2d(topi.mali.conv2d_nchw_winograd),
wrap_topi_schedule(topi.mali.schedule_conv2d_nchw_winograd),
name="conv2d_nchw_winograd.mali",
plevel=15)
plevel=5)
elif re.match(r"OIHW\d*o", kernel_layout):
strategy.add_implementation(
wrap_compute_conv2d(topi.mali.conv2d_nchw_spatial_pack),
......
python/tvm/relay/op/strategy/rocm.py
......@@ -77,13 +77,12 @@ def conv2d_strategy_rocm(attrs, inputs, out_type, target):
else:
raise RuntimeError("Unsupported conv2d layout {} for CUDA".format(layout))
# add miopen implementation
if "miopen" in target.libs:
if layout == "NCHW":
strategy.add_implementation(
wrap_compute_conv2d(topi.rocm.conv2d_nchw_miopen, True),
wrap_topi_schedule(topi.rocm.schedule_conv2d_nchw_miopen),
name="conv2d_nchw_miopen.rocm",
plevel=15)
if "miopen" in target.libs and layout == "NCHW":
strategy.add_implementation(
wrap_compute_conv2d(topi.rocm.conv2d_nchw_miopen, True),
wrap_topi_schedule(topi.rocm.schedule_conv2d_nchw_miopen),
name="conv2d_nchw_miopen.rocm",
plevel=15)
elif is_depthwise_conv2d(data.shape, layout, kernel.shape, kernel_layout, groups):
if layout == "NCHW":
assert kernel_layout == "OIHW"
......@@ -120,9 +119,8 @@ def conv2d_strategy_rocm(attrs, inputs, out_type, target):
@dense_strategy.register("rocm")
def dense_strategy_rocm(attrs, inputs, out_type, target):
"""Dense strategy for ROCM"""
strategy = _op.OpStrategy()
assert len(inputs[0].shape) == 2 and len(inputs[1].shape) == 2, "Only support 2-dim dense"
strategy = _op.OpStrategy()
strategy.add_implementation(
wrap_compute_dense(topi.rocm.dense),
wrap_topi_schedule(topi.rocm.schedule_dense),
......@@ -133,5 +131,5 @@ def dense_strategy_rocm(attrs, inputs, out_type, target):
wrap_compute_dense(topi.rocm.dense_rocblas),
wrap_topi_schedule(topi.rocm.dense_rocblas),
name="dense_rocblas.rocm",
plevel=5)
plevel=15)
return strategy
python/tvm/relay/op/strategy/x86.py
......@@ -232,13 +232,13 @@ def dense_strategy_cpu(attrs, inputs, out_type, target):
strategy.add_implementation(wrap_compute_dense(topi.x86.dense_cblas),
wrap_topi_schedule(topi.x86.schedule_dense_cblas),
name="dense_cblas.x86",
plevel=5)
plevel=15)
with SpecializedCondition(m >= 16):
# this implementation may not be well-optimized, so use plevel=8 for now.
strategy.add_implementation(wrap_compute_dense(topi.x86.dense_pack),
wrap_topi_schedule(topi.x86.schedule_dense_pack),
name="dense_pack.x86",
plevel=8)
plevel=5)
return strategy
@batch_matmul_strategy.register("cpu")
......@@ -253,7 +253,7 @@ def batch_matmul_strategy_cpu(attrs, inputs, out_type, target):
strategy.add_implementation(wrap_compute_batch_matmul(topi.x86.batch_matmul_cblas),
wrap_topi_schedule(topi.x86.schedule_batch_matmul_cblas),
name="batch_matmul_cblas.x86",
plevel=5)
plevel=15)
return strategy
@schedule_sparse_dense.register("cpu")
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment