Skip to content

Commit

Permalink
Remove comma between fields in LoweringConfig assemblyFormat (iree-or…
Browse files Browse the repository at this point in the history
…g#9116)

The comma was not scaling to other fields see iree-org#9114
  • Loading branch information
nirvedhmeshram authored May 13, 2022
1 parent b46ff9e commit c8a68f1
Show file tree
Hide file tree
Showing 19 changed files with 54 additions and 54 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -47,7 +47,7 @@ hal.executable private @matmul_tensors {
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUTileFuseAndVectorize, workload_per_wg = [64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUTileFuseAndVectorize workload_per_wg = [64, 64]>
// CHECK: hal.executable.entry_point public @matmul_tensors
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[DEVICE:.+]]: !hal.device,
Expand Down Expand Up @@ -112,7 +112,7 @@ hal.executable private @add {
}
}
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [64, 64]>
// CHECK: hal.executable private @add
// CHECK: hal.executable.entry_point public @add
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -178,7 +178,7 @@ hal.executable private @add4D {
}
}
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [64, 64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable.entry_point public @add4D
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[DEVICE:.+]]: !hal.device,
Expand Down Expand Up @@ -238,7 +238,7 @@ hal.executable private @batch_matmul_tensors {
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUTileFuseAndVectorize, workload_per_wg = [64, 64, 1]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUTileFuseAndVectorize workload_per_wg = [64, 64, 1]>
// CHECK: hal.executable.entry_point public @batch_matmul_tensors
// CHECK-NEXT: (%[[DEVICE:.+]]: !hal.device,
// CHECK-SAME: %[[WORKLOAD_X:[a-zA-Z0-9_]+]]: index
Expand Down Expand Up @@ -290,7 +290,7 @@ hal.executable private @preset_config_matmul_tensors {
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 16)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [16, 32]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [16, 32]>
// CHECK: hal.executable.entry_point public @preset_config
// CHECK-NEXT: (%[[DEVICE:.+]]: !hal.device,
// CHECK-SAME: %[[WORKLOAD_X:[a-zA-Z0-9_]+]]: index
Expand Down Expand Up @@ -347,7 +347,7 @@ hal.executable public @copy_op {
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUBufferOpsTileAndVectorize, workload_per_wg = [64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUBufferOpsTileAndVectorize workload_per_wg = [64, 64]>
// CHECK: hal.executable.entry_point public @copy_op
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[DEVICE:.+]]: !hal.device,
Expand Down Expand Up @@ -398,7 +398,7 @@ hal.executable private @static_1d_fft_stage2 {
}
}
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault workload_per_wg = [64]>
// CHECK: hal.executable private @static_1d_fft_stage2
// CHECK: hal.executable.entry_point public @static_1d_fft_stage2
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -442,7 +442,7 @@ hal.executable private @static_3d_fft_stage3 {
}
}
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [64, 64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable private @static_3d_fft_stage3
// CHECK: hal.executable.entry_point public @static_3d_fft_stage3
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -512,7 +512,7 @@ hal.executable private @outs_fusion {
}
}
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [64, 64]>
// CHECK: hal.executable private @outs_fusion
// CHECK: hal.executable.entry_point public @outs_fusion_fn
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -577,7 +577,7 @@ hal.executable private @conv {
}
}
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [64, 64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable private @conv
// CHECK: hal.executable.entry_point public @conv
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -636,7 +636,7 @@ hal.executable private @conv_static {
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 48)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 40)>
// CHECK-DAG: #[[MAP2:.+]] = affine_map<()[s0] -> (s0 ceildiv 20)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [48, 40, 20]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault workload_per_wg = [48, 40, 20]>
// CHECK: hal.executable private @conv_static
// CHECK: hal.executable.entry_point public @conv_static
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -693,7 +693,7 @@ hal.executable private @generic_static {
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 16)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [32, 16]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [32, 16]>
// CHECK: hal.executable private @generic_static
// CHECK: hal.executable.entry_point public @generic_static
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -751,7 +751,7 @@ hal.executable private @matmul_static {
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 28)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUTileFuseAndVectorize, workload_per_wg = [8, 28]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUTileFuseAndVectorize workload_per_wg = [8, 28]>
// CHECK: hal.executable private @matmul_static
// CHECK: hal.executable.entry_point public @matmul_static
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -808,7 +808,7 @@ hal.executable private @restrict_num_workgroups {
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 7)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [64, 7, 1]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDefault workload_per_wg = [64, 7, 1]>
// CHECK: hal.executable private @restrict_num_workgroups
// CHECK: hal.executable.entry_point public @restrict_num_workgroups
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -872,7 +872,7 @@ hal.executable private @reduction {
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [4]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [4]>
// CHECK: hal.executable private @reduction
// CHECK: hal.executable.entry_point public @reduction
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -930,7 +930,7 @@ hal.executable private @gemm_unit_N {
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [64]>
// CHECK: hal.executable private @gemm_unit_N
// CHECK: hal.executable.entry_point public @gemm_unit_N
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -1045,7 +1045,7 @@ hal.executable private @generic_unit_dims {
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [64, 64, 64]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable private @generic_unit_dims
// CHECK: hal.executable.entry_point public @generic_unit_dims
// CHECK-SAME: translation_info = #[[TRANSLATION]]
Expand Down Expand Up @@ -1219,7 +1219,7 @@ hal.executable private @matmul_interchange {
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert, workload_per_wg = [64, 32]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUDoubleTilingExpert workload_per_wg = [64, 32]>
// CHECK: hal.executable.entry_point public @matmul_interchange
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[DEVICE:.+]]: !hal.device,
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -54,7 +54,7 @@ hal.executable private @dispatch_0 {
]>

// CHECK-LABEL: func.func @workgroup_tile_loop()
#translation = #iree_codegen.translation_info<LLVMGPUDistribute, workload_per_wg = [32]>
#translation = #iree_codegen.translation_info<LLVMGPUDistribute workload_per_wg = [32]>
hal.executable private @workgroup_tile_loop {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @workgroup_tile_loop layout(#executable_layout) {
Expand Down Expand Up @@ -88,7 +88,7 @@ hal.executable private @workgroup_tile_loop {
]>

// CHECK-LABEL: func.func @workgroup_tile_loop_negative()
#translation = #iree_codegen.translation_info<LLVMGPUDistribute, workload_per_wg = [16]>
#translation = #iree_codegen.translation_info<LLVMGPUDistribute workload_per_wg = [16]>
hal.executable private @workgroup_tile_loop_negative {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @workgroup_tile_loop_negative layout(#executable_layout) {
Expand Down Expand Up @@ -124,7 +124,7 @@ hal.executable private @workgroup_tile_loop_negative {
// CHECK-LABEL: func.func @both_workgroup_and_workitem()
// CHECK-NOT: scf.for
// CHECK: gpu.barrier
#translation = #iree_codegen.translation_info<LLVMGPUDistribute, workload_per_wg = [32, 8, 1]>
#translation = #iree_codegen.translation_info<LLVMGPUDistribute workload_per_wg = [32, 8, 1]>
hal.executable private @both_workgroup_and_workitem {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @both_workgroup_and_workitem layout(#executable_layout) {
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -328,7 +328,7 @@ hal.executable private @preset_config_matmul_tensors {
]>
]>
#executable_target_system_elf_x86_64_ = #hal.executable.target<"llvm", "system-elf-x86_64">
#translation = #iree_codegen.translation_info<CPUBufferOpsTileAndVectorize, workload_per_wg = [64, 64]>
#translation = #iree_codegen.translation_info<CPUBufferOpsTileAndVectorize workload_per_wg = [64, 64]>
hal.executable public @copy_op {
hal.executable.variant public @system_elf_x86_64, target = #executable_target_system_elf_x86_64_ {
hal.executable.entry_point public @copy_op layout(#executable_layout) {translation_info = #translation}
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -104,7 +104,7 @@ def IREECodegen_TranslationInfoAttr :
}];

let assemblyFormat = [{
`<` `` $passPipeline (`,` `workload_per_wg` `=` $workloadPerWorkgroup^)? `>`
`<` `` $passPipeline (`workload_per_wg` `=` $workloadPerWorkgroup^)? `>`
}];

let parameters = (ins
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -2,11 +2,11 @@

module {
func.func @test() attributes {
lowring_config = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [32, 42]>} {
lowring_config = #iree_codegen.translation_info<CPUDefault workload_per_wg = [32, 42]>} {
return
}
}
// CHECK: #translation = #iree_codegen.translation_info<CPUDefault, workload_per_wg = [32, 42]>
// CHECK: #translation = #iree_codegen.translation_info<CPUDefault workload_per_wg = [32, 42]>

// -----

Expand Down
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
// RUN: iree-opt --split-input-file --pass-pipeline='hal.executable(hal.executable.variant(builtin.module(func.func(iree-llvmgpu-tile-and-distribute))))' %s | FileCheck %s

#config = #iree_codegen.lowering_config<tile_sizes = [[2, 256, 4]]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt, workload_per_wg = [256, 2]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt workload_per_wg = [256, 2]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
Expand Down Expand Up @@ -85,7 +85,7 @@ hal.executable private @dot_dispatch_0 {

// -----

#translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt, workload_per_wg = [32, 8, 1]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt workload_per_wg = [32, 8, 1]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
Expand Down Expand Up @@ -160,7 +160,7 @@ builtin.module {
// -----

#config = #iree_codegen.lowering_config<tile_sizes = [[2, 32, 4]]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt, workload_per_wg = [32, 2]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulSimt workload_per_wg = [32, 2]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
Expand Down Expand Up @@ -294,7 +294,7 @@ hal.executable @reduction_dispatch {

// -----

#translation = #iree_codegen.translation_info<LLVMGPUVectorize, workload_per_wg = [256, 1, 1]>
#translation = #iree_codegen.translation_info<LLVMGPUVectorize workload_per_wg = [256, 1, 1]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -283,7 +283,7 @@ hal.executable private @static_3d_fft_stage3 {

#compilation = #iree_codegen.compilation_info<
lowering_config = <tile_sizes = [[32, 128, 64]]>,
translation_info = <LLVMGPUMatmulSimt, workload_per_wg = [256, 32]>,
translation_info = <LLVMGPUMatmulSimt workload_per_wg = [256, 32]>,
workgroup_size = [16, 8, 1]>
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
Expand Down Expand Up @@ -320,7 +320,7 @@ hal.executable.variant public @cuda_nvptx_fb, target = <"cuda", "cuda-nvptx-fb">
}

// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering_config<tile_sizes = {{\[}}[32, 128, 64]{{\]}}
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<LLVMGPUMatmulSimt, workload_per_wg = [256, 32]>
// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<LLVMGPUMatmulSimt workload_per_wg = [256, 32]>
// CHECK: hal.executable.entry_point public @_lowering_config_test_dispatch_1
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [16 : index, 8 : index, 1 : index]
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -289,7 +289,7 @@ hal.executable private @matmul_tensors {
// -----

#config = #iree_codegen.lowering_config<tile_sizes = [[2, 32, 32, 16]]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulTensorCore, workload_per_wg = [32, 8, 1]>
#translation = #iree_codegen.translation_info<LLVMGPUMatmulTensorCore workload_per_wg = [32, 8, 1]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@
#map6 = affine_map<(d0, d1, d2) -> (d0, d1)>

#config = #iree_codegen.lowering_config<tile_sizes = [[8, 16], [1, 1], [0, 0, 1]]>
#translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [8, 16]>
#translation = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [8, 16]>
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
Expand Down Expand Up @@ -79,7 +79,7 @@ hal.executable private @matmul {
// -----

#config = #iree_codegen.lowering_config<tile_sizes = [[1, 4, 32], [1, 1, 1]]>
#translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [32, 4, 1]>
#translation = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [32, 4, 1]>
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
Expand Down Expand Up @@ -159,7 +159,7 @@ hal.executable private @conv_1d {
#map7 = affine_map<(d0)[s0] -> (32, -d0 + s0)>

#config = #iree_codegen.lowering_config<tile_sizes = [[0, 1, 4, 32], [0, 1, 1, 1], [0, 0, 0, 0, 1, 1, 4]]>
#translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [32, 4, 1]>
#translation = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [32, 4, 1]>
#executable_layout = #hal.executable.layout<push_constants = 9, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
Expand Down Expand Up @@ -274,7 +274,7 @@ hal.executable private @conv_2d {
// -----

#config = #iree_codegen.lowering_config<tile_sizes = [[0, 0, 1, 4, 32], [0, 0, 1, 1, 1]]>
#translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [32, 4, 1]>
#translation = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [32, 4, 1]>
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
Expand Down Expand Up @@ -344,7 +344,7 @@ hal.executable private @conv_3d {
#map7 = affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 1092 + s0 + d1 * 78 + d2 * 6 + d3)>

#config = #iree_codegen.lowering_config<tile_sizes = [[1, 4, 32], [1, 1, 1]]>
#translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [32, 4, 1]>
#translation = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [32, 4, 1]>
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
Expand Down Expand Up @@ -411,7 +411,7 @@ module {
// -----

#config = #iree_codegen.lowering_config<tile_sizes = [[32], [1]]>
#translation = #iree_codegen.translation_info<SPIRVDistribute, workload_per_wg = [32]>
#translation = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [32]>
#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
Expand Down
Loading

0 comments on commit c8a68f1

Please sign in to comment.