[TOPI] Use f-strings for string formatting, NFC (apache#14839)

python/tvm/topi/adreno/conv2d_alter_op.py

@@ -130,8 +130,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
                 dtype=kernel_tensor.dtype,
             )
             new_workload = autotvm.task.args_to_workload(
-                [new_data, new_weight, strides, padding, dilation, out_dtype],
-                wkl_name,
+                [new_data, new_weight, strides, padding, dilation, out_dtype], wkl_name
             )
             dispatch_ctx.update(target, new_workload, cfg)
             return relay.nn.contrib_conv2d_winograd_without_weight_transform(
@@ -165,18 +164,17 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
 
         if in_channel_block != 4 or num_filter_block != 4:
             new_workload = autotvm.task.args_to_workload(
-                [new_data, new_weight, strides, padding, dilation, out_dtype],
-                wkl_name,
+                [new_data, new_weight, strides, padding, dilation, out_dtype], wkl_name
             )
             dispatch_ctx.update(target, new_workload, cfg)
             return relay.nn.contrib_conv2d_winograd_without_weight_transform(
                 inputs[0], weight, **new_attrs
             )
 
-        new_attrs["data_layout"] = "NCHW%dc" % in_channel_block
+        new_attrs["data_layout"] = f"NCHW{in_channel_block}c"
         # (oc, ic, h, w) -> (h, w, ic, oc // 4, oc % 4)
-        new_attrs["kernel_layout"] = "HWIO%do" % num_filter_block
-        new_attrs["out_layout"] = "NCHW%dc" % num_filter_block
+        new_attrs["kernel_layout"] = f"HWIO{num_filter_block}o"
+        new_attrs["out_layout"] = f"NCHW{num_filter_block}c"
         # Store altered operator's config
         new_data = te.placeholder(
             (N, CI // in_channel_block, H, W, in_channel_block), dtype=data_dtype
@@ -186,15 +184,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
             dtype=kernel_tensor.dtype,
         )
         new_workload = autotvm.task.args_to_workload(
-            [
-                new_data,
-                new_weight,
-                strides,
-                padding,
-                dilation,
-                out_dtype,
-            ],
-            wkl_name,
+            [new_data, new_weight, strides, padding, dilation, out_dtype], wkl_name
         )
         dispatch_ctx.update(target, new_workload, cfg)
         return relay.nn.contrib_conv2d_winograd_without_weight_transform(
@@ -226,8 +216,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
                 dtype=kernel_tensor.dtype,
             )
             new_workload = autotvm.task.args_to_workload(
-                [new_data, new_weight, strides, padding, dilation, out_dtype],
-                wkl_name,
+                [new_data, new_weight, strides, padding, dilation, out_dtype], wkl_name
             )
             dispatch_ctx.update(target, new_workload, cfg)
             return relay.nn.contrib_conv2d_winograd_without_weight_transform(
@@ -259,18 +248,17 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
 
         if in_channel_block != 4 or num_filter_block != 4:
             new_workload = autotvm.task.args_to_workload(
-                [new_data, new_weight, strides, padding, dilation, out_dtype],
-                wkl_name,
+                [new_data, new_weight, strides, padding, dilation, out_dtype], wkl_name
             )
             dispatch_ctx.update(target, new_workload, cfg)
             return relay.nn.contrib_conv2d_winograd_without_weight_transform(
                 inputs[0], weight, **new_attrs
             )
 
-        new_attrs["data_layout"] = "NHWC%dc" % in_channel_block
+        new_attrs["data_layout"] = f"NHWC{in_channel_block}c"
         # (oc, ic, h, w) -> (h, w, ic, oc // 4, oc % 4)
-        new_attrs["kernel_layout"] = "HWIO%do" % num_filter_block
-        new_attrs["out_layout"] = "NHWC%dc" % num_filter_block
+        new_attrs["kernel_layout"] = f"HWIO{num_filter_block}o"
+        new_attrs["out_layout"] = f"NHWC{num_filter_block}c"
         # Store altered operator's config
         new_data = te.placeholder(
             (N, H, W, CI // in_channel_block, in_channel_block), dtype=data_dtype
@@ -280,15 +268,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
             dtype=kernel_tensor.dtype,
         )
         new_workload = autotvm.task.args_to_workload(
-            [
-                new_data,
-                new_weight,
-                strides,
-                padding,
-                dilation,
-                out_dtype,
-            ],
-            wkl_name,
+            [new_data, new_weight, strides, padding, dilation, out_dtype], wkl_name
         )
         dispatch_ctx.update(target, new_workload, cfg)
         return relay.nn.contrib_conv2d_winograd_without_weight_transform(
@@ -316,12 +296,12 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
             # update new attrs
             new_attrs["channels"] = out_channel
             if in_channel_block == 4:
-                new_attrs["data_layout"] = "NCHW%dc" % in_channel_block
+                new_attrs["data_layout"] = f"NCHW{in_channel_block}c"
             else:
                 new_attrs["data_layout"] = "NCHW"
             # (oc, ic, h, w) -> (OC, ic, h, w, oc)
-            new_attrs["kernel_layout"] = "OIHW%do" % num_filter_block
-            new_attrs["out_layout"] = "NCHW%dc" % num_filter_block
+            new_attrs["kernel_layout"] = f"OIHW{num_filter_block}o"
+            new_attrs["out_layout"] = f"NCHW{num_filter_block}c"
 
             # Store altered operator's config for applying of tuned AutoTVM statistics
             if in_channel_block == 4:
@@ -336,14 +316,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
                 dtype=kernel_tensor.dtype,
             )
             new_workload = autotvm.task.args_to_workload(
-                [
-                    new_data,
-                    new_kernel,
-                    strides,
-                    padding,
-                    dilation,
-                    out_dtype,
-                ],
+                [new_data, new_kernel, strides, padding, dilation, out_dtype],
                 topi_tmpl,  # "conv2d_nchwc.image2d",
             )
             dispatch_ctx.update(target, new_workload, cfg)
@@ -376,15 +349,15 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
             # update new attrs
             new_attrs["channels"] = out_channles
             if in_channel_block == 4:
-                new_attrs["data_layout"] = "NHWC%dc" % in_channel_block
+                new_attrs["data_layout"] = f"NHWC{in_channel_block}c"
             else:
                 new_attrs["data_layout"] = "NHWC"
             # (h, w, ic, oc) -> (h, w, ic, OC, oc)
             if kernel_layout == "HWIO":
-                new_attrs["kernel_layout"] = "HWIO%do" % num_filter_block
+                new_attrs["kernel_layout"] = f"HWIO{num_filter_block}o"
             else:
-                new_attrs["kernel_layout"] = "HWOI%do" % num_filter_block
-            new_attrs["out_layout"] = "NHWC%dc" % num_filter_block
+                new_attrs["kernel_layout"] = f"HWOI{num_filter_block}o"
+            new_attrs["out_layout"] = f"NHWC{num_filter_block}c"
 
             # Store altered operator's config for applying of tuned AutoTVM statistics
             if in_channel_block == 4:
@@ -423,15 +396,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
                     dtype=kernel_tensor.dtype,
                 )
             new_workload = autotvm.task.args_to_workload(
-                [
-                    new_data,
-                    new_kernel,
-                    strides,
-                    padding,
-                    dilation,
-                    out_dtype,
-                ],
-                topi_tmpl,
+                [new_data, new_kernel, strides, padding, dilation, out_dtype], topi_tmpl
             )
             dispatch_ctx.update(target, new_workload, cfg)
         else:

python/tvm/topi/adreno/pooling.py

@@ -120,7 +120,7 @@ def traverse(OP):
             Pool = OP.output(0)
             _schedule_global(Pool, layout)
         else:
-            raise RuntimeError("Unsupported operator: %s" % OP.tag)
+            raise RuntimeError(f"Unsupported operator: {OP.tag}")
 
         scheduled_ops.append(OP)
 
@@ -188,7 +188,7 @@ def traverse(OP):
             Pool = OP.output(0)
             _schedule(PaddedInput, Pool)
         else:
-            raise RuntimeError("Unsupported operator: %s" % OP.tag)
+            raise RuntimeError(f"Unsupported operator: {OP.tag}")
 
         scheduled_ops.append(OP)
 

python/tvm/topi/arm_cpu/bitserial_dense.py

@@ -205,7 +205,7 @@ def traverse(op):
             unipolar = output.op.tag == "bitserial_dense_unipolar"
             _schedule(cfg, s, data_vec, weight_vec, output, unipolar)
         else:
-            raise RuntimeError("Unsupported operator: %s" % op.tag)
+            raise RuntimeError(f"Unsupported operator: {op.tag}")
 
     traverse(outs[0].op)
     return s

python/tvm/topi/arm_cpu/tensor_intrin.py

@@ -466,42 +466,36 @@ def dot_int8_int8_int32_neon_82(int32_lanes, dtype="uint"):
     """
     num_int8_elements = 4  # 4 int8 elements in int32
 
-    data = te.placeholder((num_int8_elements,), dtype="%s8" % dtype, name="data")
-    kernel = te.placeholder((int32_lanes, num_int8_elements), dtype="%s8" % dtype, name="kernel")
+    data = te.placeholder((num_int8_elements,), dtype=f"{dtype}8", name="data")
+    kernel = te.placeholder((int32_lanes, num_int8_elements), dtype=f"{dtype}8", name="kernel")
 
     k = te.reduce_axis((0, num_int8_elements), name="k")
     C = te.compute(
         (int32_lanes,),
-        lambda i: te.sum(
-            data[k].astype("%s32" % dtype) * kernel[i, k].astype("%s32" % dtype), axis=k
-        ),
+        lambda i: te.sum(data[k].astype(f"{dtype}32") * kernel[i, k].astype(f"{dtype}32"), axis=k),
         name="C",
     )
 
     a_buffer = tvm.tir.decl_buffer(
-        data.shape, dtype="%s8" % dtype, name="a_buffer", offset_factor=1, strides=[1]
+        data.shape, dtype=f"{dtype}8", name="a_buffer", offset_factor=1, strides=[1]
     )
     b_buffer = tvm.tir.decl_buffer(
-        kernel.shape,
-        dtype="%s8" % dtype,
-        name="b_buffer",
-        offset_factor=1,
-        strides=[te.var("s"), 1],
+        kernel.shape, dtype=f"{dtype}8", name="b_buffer", offset_factor=1, strides=[te.var("s"), 1]
     )
 
     def _intrin_func(ins, outs):
         def _instr(index):
             ib = tvm.tir.ir_builder.create()
             if index == 1:
-                ib.emit(outs[0].vstore(0, tvm.tir.const(0, "%s32x%d" % (dtype, int32_lanes))))
+                ib.emit(outs[0].vstore(0, tvm.tir.const(0, f"{dtype}32x{int32_lanes}")))
                 return ib.get()
 
-            dtype_a = "%s8x%d" % (dtype, num_int8_elements)
-            dtype_b = "%s8x%d" % (dtype, int32_lanes * num_int8_elements)
-            dtype_c = "%s32x%d" % (dtype, int32_lanes)
+            dtype_a = f"{dtype}8x{num_int8_elements}"
+            dtype_b = f"{dtype}8x{int32_lanes * num_int8_elements}"
+            dtype_c = f"{dtype}32x{int32_lanes}"
 
             a_int8 = ins[0].vload([0], dtype_a)
-            re_int32 = tvm.tir.call_intrin("%s32" % dtype, "tir.reinterpret", a_int8)
+            re_int32 = tvm.tir.call_intrin(f"{dtype}32", "tir.reinterpret", a_int8)
             # broadcast a
             vec_ai32 = re_int32.astype(dtype_c)
 
@@ -805,12 +799,7 @@ def _instr(index):
                 #                           a*2+b*6+c*10+d*14,
                 #                           a*3+b*7+c*11+d*15]
                 vdot = tvm.tir.call_llvm_intrin(
-                    "int32x4",
-                    llvm_intrin,
-                    tvm.tir.const(3, "uint32"),
-                    vec_c,
-                    vec_b,
-                    vec_aa[i],
+                    "int32x4", llvm_intrin, tvm.tir.const(3, "uint32"), vec_c, vec_b, vec_aa[i]
                 )
 
                 # Store the result
@@ -885,11 +874,7 @@ def gemm_acc_nx16_int8_int8_int32(dtype, rows):
         A.shape, dtype, name="aa_buffer", offset_factor=1, strides=[te.var("sa"), 1]
     )
     bb_buffer = tvm.tir.decl_buffer(
-        B.shape,
-        dtype,
-        name="bb_buffer",
-        offset_factor=1,
-        strides=[te.var("sb0"), te.var("sb1"), 1],
+        B.shape, dtype, name="bb_buffer", offset_factor=1, strides=[te.var("sb0"), te.var("sb1"), 1]
     )
     cc_buffer = tvm.tir.decl_buffer(
         C.shape, dtype="int32", name="cc_buffer", offset_factor=1, strides=[te.var("sc"), 1]
@@ -936,12 +921,7 @@ def _instr(index):
                         #                           a*2+b*18+c*34+d*50,
                         #                           a*3+b*19+c*35+d*51]
                         vdot = tvm.tir.call_llvm_intrin(
-                            "int32x4",
-                            llvm_intrin,
-                            tvm.tir.const(3, "uint32"),
-                            vec_c,
-                            vec_b,
-                            vec_aa,
+                            "int32x4", llvm_intrin, tvm.tir.const(3, "uint32"), vec_c, vec_b, vec_aa
                         )
                         ib.emit(outs[0].vstore([k, 4 * j], vdot))
             return ib.get()
@@ -977,27 +957,17 @@ def smlal_int16_int32():
     A = te.placeholder((int16_lanes,), dtype="int16", name="A")
     B = te.placeholder((int16_lanes, 1), dtype="int16", name="B")
     C = te.compute(
-        (int16_lanes,),
-        lambda i: A[i].astype("int32") * B[i, 0].astype("int32"),
-        name="C",
+        (int16_lanes,), lambda i: A[i].astype("int32") * B[i, 0].astype("int32"), name="C"
     )
 
     a_buffer = tvm.tir.decl_buffer(
         A.shape, dtype="int16", name="a_buffer", offset_factor=1, strides=[1]
     )
     b_buffer = tvm.tir.decl_buffer(
-        B.shape,
-        dtype="int16",
-        name="b_buffer",
-        offset_factor=1,
-        strides=[te.var("sb"), 1],
+        B.shape, dtype="int16", name="b_buffer", offset_factor=1, strides=[te.var("sb"), 1]
     )
     c_buffer = tvm.tir.decl_buffer(
-        C.shape,
-        dtype="int32",
-        name="c_buffer",
-        offset_factor=1,
-        strides=[1],
+        C.shape, dtype="int32", name="c_buffer", offset_factor=1, strides=[1]
     )
 
     def _intrin_func(ins, outs):
@@ -1122,12 +1092,7 @@ def _instr(index):
             #          i*1 + j*3 + k*5 + l*7 +m*9 + n*11 + o*13 + p*15]
             vec_c = outs[0].vload([0, 0], "int32x4")
             vmmla = tvm.tir.call_llvm_intrin(
-                "int32x4",
-                llvm_intrin,
-                tvm.tir.const(3, "uint32"),
-                vec_c,
-                vec_a,
-                vec_b,
+                "int32x4", llvm_intrin, tvm.tir.const(3, "uint32"), vec_c, vec_a, vec_b
             )
             # Store the result
             ib.emit(outs[0].vstore([0, 0], vmmla))

python/tvm/topi/bifrost/conv2d.py

@@ -509,7 +509,7 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
         CO, _, KH, KW = get_const_tuple(kernel.shape)
         VC = cfg["tile_co"].size[-1]
 
-        new_attrs["kernel_layout"] = "OIHW%do" % VC
+        new_attrs["kernel_layout"] = f"OIHW{VC}o"
 
         new_data = data
         new_kernel = te.placeholder((idxd(CO, VC), CI, KH, KW, VC), dtype=kernel.dtype)

python/tvm/topi/cuda/batch_matmul_tensorcore.py

@@ -118,7 +118,7 @@ def _schedule(cfg, s, C):
             wmma_m = wmma_n = 8
             wmma_k = 32
         else:
-            raise ValueError("data dtype %s is not yet supported" % data_dtype)
+            raise ValueError(f"data dtype {data_dtype} is not yet supported")
 
         warp_size = 32
         block_row_warps = cfg["block_row_warps"].val

python/tvm/topi/cuda/conv2d.py

@@ -59,7 +59,7 @@ def conv2d_cudnn(
         tensor_format = 1  # CUDNN_TENSOR_NHWC
         N, H, W, _ = get_const_tuple(data.shape)
     else:
-        raise ValueError("Unsupported layout %s in cudnn" % layout)
+        raise ValueError(f"Unsupported layout {layout} in cudnn")
     CO, CI, KH, KW = get_const_tuple(kernel.shape)
 
     # handle dilation

python/tvm/topi/cuda/conv2d_hwcn.py

@@ -155,7 +155,7 @@ def traverse(operator):
             B = operator.output(0)
             schedule(Apad, W, B)
         else:
-            raise RuntimeError("Unsupported operator: %s" % operator.tag)
+            raise RuntimeError(f"Unsupported operator: {operator.tag}")
 
         scheduled_ops.append(operator)
 

python/tvm/topi/cuda/conv3d.py

@@ -197,7 +197,7 @@ def conv3d_cudnn(
         tensor_format = 1  # CUDNN_TENSOR_NHWC
         N, D, H, W, _ = get_const_tuple(data.shape)
     else:
-        raise ValueError("Unsupported layout %s in cudnn" % layout)
+        raise ValueError(f"Unsupported layout {layout} in cudnn")
     CO, CI, KD, KH, KW = get_const_tuple(kernel.shape)
 
     assert groups == 1, "conv3d_cudnn does not support groups"

python/tvm/topi/cuda/dense_tensorcore.py

@@ -152,7 +152,7 @@ def _schedule_dense_tensorcore(cfg, s, C):
         wmma_m = wmma_n = 8
         wmma_k = 32
     else:
-        raise ValueError("data dtype %s is not yet supported" % data_dtype)
+        raise ValueError(f"data dtype {data_dtype} is not yet supported")
 
     warp_size = 32
     block_row_warps = cfg["block_row_warps"].val