| /* |
| * Copyright 2016 Google Inc. |
| * |
| * Use of this source code is governed by a BSD-style license that can be |
| * found in the LICENSE file. |
| * |
| */ |
| |
| // |
| // |
| // |
| |
| #include <stdlib.h> |
| #include <stdbool.h> |
| #include <string.h> |
| #include <getopt.h> |
| #include <inttypes.h> |
| |
| // |
| // |
| // |
| |
| #include "networks.h" |
| #include "common/util.h" |
| #include "common/macros.h" |
| |
| // |
| // |
| // |
| |
| #undef HSG_OP_EXPAND_X |
| #define HSG_OP_EXPAND_X(t) #t , |
| |
| char const * const |
| hsg_op_type_string[] = |
| { |
| HSG_OP_EXPAND_ALL() |
| }; |
| |
| // |
| // |
| // |
| |
| #define EXIT() (struct hsg_op){ HSG_OP_TYPE_EXIT } |
| |
| #define END() (struct hsg_op){ HSG_OP_TYPE_END } |
| #define BEGIN() (struct hsg_op){ HSG_OP_TYPE_BEGIN } |
| #define ELSE() (struct hsg_op){ HSG_OP_TYPE_ELSE } |
| |
| #define TARGET_BEGIN() (struct hsg_op){ HSG_OP_TYPE_TARGET_BEGIN } |
| #define TARGET_END() (struct hsg_op){ HSG_OP_TYPE_TARGET_END } |
| |
| #define TRANSPOSE_KERNEL_PROTO() (struct hsg_op){ HSG_OP_TYPE_TRANSPOSE_KERNEL_PROTO } |
| #define TRANSPOSE_KERNEL_PREAMBLE() (struct hsg_op){ HSG_OP_TYPE_TRANSPOSE_KERNEL_PREAMBLE } |
| #define TRANSPOSE_KERNEL_BODY() (struct hsg_op){ HSG_OP_TYPE_TRANSPOSE_KERNEL_BODY } |
| |
| #define BS_KERNEL_PROTO(i) (struct hsg_op){ HSG_OP_TYPE_BS_KERNEL_PROTO, { i } } |
| #define BS_KERNEL_PREAMBLE(i) (struct hsg_op){ HSG_OP_TYPE_BS_KERNEL_PREAMBLE, { i } } |
| |
| #define BC_KERNEL_PROTO(i) (struct hsg_op){ HSG_OP_TYPE_BC_KERNEL_PROTO, { i } } |
| #define BC_KERNEL_PREAMBLE(i) (struct hsg_op){ HSG_OP_TYPE_BC_KERNEL_PREAMBLE, { i } } |
| |
| #define FM_KERNEL_PROTO(s,r) (struct hsg_op){ HSG_OP_TYPE_FM_KERNEL_PROTO, { s, r } } |
| #define FM_KERNEL_PREAMBLE(l,r) (struct hsg_op){ HSG_OP_TYPE_FM_KERNEL_PREAMBLE, { l, r } } |
| |
| #define HM_KERNEL_PROTO(s) (struct hsg_op){ HSG_OP_TYPE_HM_KERNEL_PROTO, { s } } |
| #define HM_KERNEL_PREAMBLE(l) (struct hsg_op){ HSG_OP_TYPE_HM_KERNEL_PREAMBLE, { l } } |
| |
| #define BX_REG_GLOBAL_LOAD(n,v) (struct hsg_op){ HSG_OP_TYPE_BX_REG_GLOBAL_LOAD, { n, v } } |
| #define BX_REG_GLOBAL_STORE(n) (struct hsg_op){ HSG_OP_TYPE_BX_REG_GLOBAL_STORE, { n } } |
| |
| #define FM_REG_GLOBAL_LOAD_LEFT(n,i) (struct hsg_op){ HSG_OP_TYPE_FM_REG_GLOBAL_LOAD_LEFT, { n, i } } |
| #define FM_REG_GLOBAL_STORE_LEFT(n,i) (struct hsg_op){ HSG_OP_TYPE_FM_REG_GLOBAL_STORE_LEFT, { n, i } } |
| #define FM_REG_GLOBAL_LOAD_RIGHT(n,i) (struct hsg_op){ HSG_OP_TYPE_FM_REG_GLOBAL_LOAD_RIGHT, { n, i } } |
| #define FM_REG_GLOBAL_STORE_RIGHT(n,i) (struct hsg_op){ HSG_OP_TYPE_FM_REG_GLOBAL_STORE_RIGHT, { n, i } } |
| #define FM_MERGE_RIGHT_PRED(n,s) (struct hsg_op){ HSG_OP_TYPE_FM_MERGE_RIGHT_PRED, { n, s } } |
| |
| #define HM_REG_GLOBAL_LOAD(n,i) (struct hsg_op){ HSG_OP_TYPE_HM_REG_GLOBAL_LOAD, { n, i } } |
| #define HM_REG_GLOBAL_STORE(n,i) (struct hsg_op){ HSG_OP_TYPE_HM_REG_GLOBAL_STORE, { n, i } } |
| |
| #define SLAB_FLIP(f) (struct hsg_op){ HSG_OP_TYPE_SLAB_FLIP, { f } } |
| #define SLAB_HALF(h) (struct hsg_op){ HSG_OP_TYPE_SLAB_HALF, { h } } |
| |
| #define CMP_FLIP(a,b,c) (struct hsg_op){ HSG_OP_TYPE_CMP_FLIP, { a, b, c } } |
| #define CMP_HALF(a,b) (struct hsg_op){ HSG_OP_TYPE_CMP_HALF, { a, b } } |
| |
| #define CMP_XCHG(a,b,p) (struct hsg_op){ HSG_OP_TYPE_CMP_XCHG, { a, b, p } } |
| |
| #define BS_REG_SHARED_STORE_V(m,i,r) (struct hsg_op){ HSG_OP_TYPE_BS_REG_SHARED_STORE_V, { m, i, r } } |
| #define BS_REG_SHARED_LOAD_V(m,i,r) (struct hsg_op){ HSG_OP_TYPE_BS_REG_SHARED_LOAD_V, { m, i, r } } |
| #define BC_REG_SHARED_LOAD_V(m,i,r) (struct hsg_op){ HSG_OP_TYPE_BC_REG_SHARED_LOAD_V, { m, i, r } } |
| |
| #define BX_REG_SHARED_STORE_LEFT(r,i,p) (struct hsg_op){ HSG_OP_TYPE_BX_REG_SHARED_STORE_LEFT, { r, i, p } } |
| #define BS_REG_SHARED_STORE_RIGHT(r,i,p) (struct hsg_op){ HSG_OP_TYPE_BS_REG_SHARED_STORE_RIGHT, { r, i, p } } |
| |
| #define BS_REG_SHARED_LOAD_LEFT(r,i,p) (struct hsg_op){ HSG_OP_TYPE_BS_REG_SHARED_LOAD_LEFT, { r, i, p } } |
| #define BS_REG_SHARED_LOAD_RIGHT(r,i,p) (struct hsg_op){ HSG_OP_TYPE_BS_REG_SHARED_LOAD_RIGHT, { r, i, p } } |
| |
| #define BC_REG_GLOBAL_LOAD_LEFT(r,i,p) (struct hsg_op){ HSG_OP_TYPE_BC_REG_GLOBAL_LOAD_LEFT, { r, i, p } } |
| |
| #define REG_F_PREAMBLE(s) (struct hsg_op){ HSG_OP_TYPE_REG_F_PREAMBLE, { s } } |
| #define REG_SHARED_STORE_F(r,i,s) (struct hsg_op){ HSG_OP_TYPE_REG_SHARED_STORE_F, { r, i, s } } |
| #define REG_SHARED_LOAD_F(r,i,s) (struct hsg_op){ HSG_OP_TYPE_REG_SHARED_LOAD_F, { r, i, s } } |
| #define REG_GLOBAL_STORE_F(r,i,s) (struct hsg_op){ HSG_OP_TYPE_REG_GLOBAL_STORE_F, { r, i, s } } |
| |
| #define BLOCK_SYNC() (struct hsg_op){ HSG_OP_TYPE_BLOCK_SYNC } |
| |
| #define BS_FRAC_PRED(m,w) (struct hsg_op){ HSG_OP_TYPE_BS_FRAC_PRED, { m, w } } |
| |
| #define BS_MERGE_H_PREAMBLE(i) (struct hsg_op){ HSG_OP_TYPE_BS_MERGE_H_PREAMBLE, { i } } |
| #define BC_MERGE_H_PREAMBLE(i) (struct hsg_op){ HSG_OP_TYPE_BC_MERGE_H_PREAMBLE, { i } } |
| |
| #define BX_MERGE_H_PRED(p) (struct hsg_op){ HSG_OP_TYPE_BX_MERGE_H_PRED, { p } } |
| |
| #define BS_ACTIVE_PRED(m,l) (struct hsg_op){ HSG_OP_TYPE_BS_ACTIVE_PRED, { m, l } } |
| |
| // |
| // DEFAULTS |
| // |
| |
| static |
| struct hsg_config hsg_config = |
| { |
| .merge = { |
| .flip = { |
| .warps = 1, |
| .lo = 1, |
| .hi = 1 |
| }, |
| .half = { |
| .warps = 1, |
| .lo = 1, |
| .hi = 1 |
| }, |
| }, |
| |
| .block = { |
| .warps_min = 1, // min warps for a block that uses smem barriers |
| .warps_max = UINT32_MAX, // max warps for the entire multiprocessor |
| .warps_mod = 2, // the number of warps necessary to load balance horizontal merging |
| |
| .smem_min = 0, |
| .smem_quantum = 1, |
| |
| .smem_bs = 49152, |
| .smem_bc = UINT32_MAX // implies field not set |
| }, |
| |
| .warp = { |
| .lanes = 32, |
| .lanes_log2 = 5, |
| }, |
| |
| .thread = { |
| .regs = 24, |
| .xtra = 0 |
| }, |
| |
| .type = { |
| .words = 2 |
| } |
| }; |
| |
| // |
| // ZERO HSG_MERGE STRUCT |
| // |
| |
| static |
| struct hsg_merge hsg_merge[MERGE_LEVELS_MAX_LOG2] = { 0 }; |
| |
| // |
| // STATS ON INSTRUCTIONS |
| // |
| |
| static hsg_op_type hsg_op_type_counts[HSG_OP_TYPE_COUNT] = { 0 }; |
| |
| // |
| // |
| // |
| |
| static |
| void |
| hsg_op_debug() |
| { |
| uint32_t total = 0; |
| |
| for (hsg_op_type t=HSG_OP_TYPE_EXIT; t<HSG_OP_TYPE_COUNT; t++) |
| { |
| uint32_t const count = hsg_op_type_counts[t]; |
| |
| total += count; |
| |
| fprintf(stderr,"%-37s : %u\n",hsg_op_type_string[t],count); |
| } |
| |
| fprintf(stderr,"%-37s : %u\n\n\n","TOTAL",total); |
| } |
| |
| // |
| // |
| // |
| |
| static |
| void |
| hsg_config_init_shared() |
| { |
| // |
| // The assumption here is that a proper smem_bs value was provided |
| // that represents the maximum fraction of the multiprocessor's |
| // available shared memory that can be accessed by the initial block |
| // sorting kernel. |
| // |
| // With CUDA devices this is 48KB out of 48KB, 64KB or 96KB. |
| // |
| // Intel subslices are a little trickier and the minimum allocation |
| // is 4KB and the maximum is 64KB on pre-Skylake IGPs. Sizes are |
| // allocated in 1KB increments. If a maximum of two block sorters |
| // can occupy a subslice then each should be assigned 32KB of shared |
| // memory. |
| // |
| // News Flash: apparently GEN9+ IGPs can allocate 1KB of SMEM per |
| // workgroup so all the previously written logic to support this |
| // issue is being removed. |
| // |
| uint32_t const bs_keys = hsg_config.block.smem_bs / (hsg_config.type.words * sizeof(uint32_t)); |
| |
| hsg_config.warp.skpw_bs = bs_keys / hsg_merge[0].warps; |
| } |
| |
| static |
| void |
| hsg_merge_levels_init_shared(struct hsg_merge * const merge) |
| { |
| { |
| // |
| // What is the max amount of shared in each possible bs block config? |
| // |
| // The provided smem_bs size will be allocated for each sorting block. |
| // |
| uint32_t const bs_threads = merge->warps << hsg_config.warp.lanes_log2; |
| uint32_t const bs_keys = hsg_config.block.smem_bs / (hsg_config.type.words * sizeof(uint32_t)); |
| uint32_t const bs_kpt = bs_keys / bs_threads; |
| uint32_t const bs_kpt_mod = (bs_kpt / hsg_config.block.warps_mod) * hsg_config.block.warps_mod; |
| uint32_t const bs_rows_even = bs_kpt_mod & ~1; // must be even because flip merge only works on row pairs |
| |
| // this is a showstopper |
| if (bs_rows_even < 2) |
| { |
| fprintf(stderr,"Error: need at least 2 rows of shared memory.\n"); |
| exit(-1); |
| } |
| |
| // clamp to number of registers |
| merge->rows_bs = MIN_MACRO(bs_rows_even, hsg_config.thread.regs); |
| } |
| |
| // |
| // smem key allocation rule for BC kernels is that a single block |
| // can't allocate more than smem_bs and must allocate at least |
| // smem_min in smem_quantum steps. |
| // |
| // Note that BC blocks will always be less than or equal to BS |
| // blocks. |
| // |
| { |
| // |
| // if merge->warps is not pow2 then we're going to skip creating a bc elsewhere |
| // |
| uint32_t const bc_warps_min = MAX_MACRO(merge->warps,hsg_config.block.warps_min); |
| uint32_t const bc_threads = bc_warps_min << hsg_config.warp.lanes_log2; |
| uint32_t const bc_block_rd = (((hsg_config.block.smem_bc * bc_warps_min) / hsg_config.block.warps_max) / |
| hsg_config.block.smem_quantum) * hsg_config.block.smem_quantum; |
| uint32_t const bc_block_max = MAX_MACRO(bc_block_rd,hsg_config.block.smem_min); |
| uint32_t const bc_block_smem = MIN_MACRO(bc_block_max,hsg_config.block.smem_bs); |
| |
| // what is the max amount of shared in each possible bc block config? |
| uint32_t const bc_keys = bc_block_smem / (hsg_config.type.words * sizeof(uint32_t)); |
| uint32_t const bc_kpt = bc_keys / bc_threads; |
| uint32_t const bc_kpt_mod = (bc_kpt / hsg_config.block.warps_mod) * hsg_config.block.warps_mod; |
| |
| merge->rows_bc = MIN_MACRO(bc_kpt_mod, hsg_config.thread.regs); |
| merge->skpw_bc = bc_keys / bc_warps_min; |
| } |
| } |
| |
| // |
| // |
| // |
| |
| static |
| void |
| hsg_merge_levels_init_1(struct hsg_merge * const merge, uint32_t const warps, uint32_t const level, uint32_t const offset) |
| { |
| uint32_t const even_odd = warps & 1; |
| |
| merge->levels[level].evenodds[even_odd]++; |
| merge->levels[level].networks[even_odd] = warps; |
| |
| if (warps == 1) |
| return; |
| |
| merge->levels[level].active.b64 |= BITS_TO_MASK_AT_64(warps,offset); |
| |
| uint32_t const count = merge->levels[level].count++; |
| uint32_t const index = (1 << level) + count; |
| uint32_t const bit = 1 << count; |
| |
| merge->levels[level].evenodd_masks[even_odd] |= bit; |
| |
| if (count > 0) |
| { |
| // offset from network to left of this network |
| uint32_t const diff = offset - merge->offsets[index-1]; |
| |
| uint32_t const diff_0 = merge->levels[level].diffs[0]; |
| uint32_t const diff_1 = merge->levels[level].diffs[1]; |
| |
| uint32_t diff_idx = UINT32_MAX; |
| |
| if ((diff_0 == 0) || (diff_0 == diff)) { |
| diff_idx = 0; |
| } else if ((diff_1 == 0) || (diff_1 == diff)) { |
| diff_idx = 1; |
| } else { |
| fprintf(stderr, "*** MORE THAN TWO DIFFS ***\n"); |
| exit(-1); |
| } |
| |
| merge->levels[level].diffs [diff_idx] = diff; |
| merge->levels[level].diff_masks[diff_idx] |= 1 << (count-1); |
| } |
| |
| merge->networks[index] = warps; |
| merge->offsets [index] = offset; |
| |
| uint32_t const l = (warps+1)/2; // lower/larger on left |
| uint32_t const r = (warps+0)/2; // higher/smaller on right |
| |
| hsg_merge_levels_init_1(merge,l,level+1,offset); |
| hsg_merge_levels_init_1(merge,r,level+1,offset+l); |
| } |
| |
| static |
| void |
| hsg_merge_levels_debug(struct hsg_merge * const merge) |
| { |
| for (uint32_t level=0; level<MERGE_LEVELS_MAX_LOG2; level++) |
| { |
| uint32_t count = merge->levels[level].count; |
| |
| if (count == 0) |
| break; |
| |
| fprintf(stderr, |
| "%-4u : %016" PRIX64 " \n", |
| count, |
| merge->levels[level].active.b64); |
| |
| fprintf(stderr, |
| "%-4u : %08X (%2u)\n" |
| "%-4u : %08X (%2u)\n", |
| merge->levels[level].diffs[0], |
| merge->levels[level].diff_masks[0], |
| POPCOUNT_MACRO(merge->levels[level].diff_masks[0]), |
| merge->levels[level].diffs[1], |
| merge->levels[level].diff_masks[1], |
| POPCOUNT_MACRO(merge->levels[level].diff_masks[1])); |
| |
| fprintf(stderr, |
| "EVEN : %08X (%2u)\n" |
| "ODD : %08X (%2u)\n", |
| merge->levels[level].evenodd_masks[0], |
| POPCOUNT_MACRO(merge->levels[level].evenodd_masks[0]), |
| merge->levels[level].evenodd_masks[1], |
| POPCOUNT_MACRO(merge->levels[level].evenodd_masks[1])); |
| |
| for (uint32_t ii=0; ii<2; ii++) |
| { |
| if (merge->levels[level].networks[ii] > 1) |
| { |
| fprintf(stderr, |
| "%-4s : ( %2u x %2u )\n", |
| (ii == 0) ? "EVEN" : "ODD", |
| merge->levels[level].evenodds[ii], |
| merge->levels[level].networks[ii]); |
| } |
| } |
| |
| uint32_t index = 1 << level; |
| |
| while (count-- > 0) |
| { |
| fprintf(stderr, |
| "[ %2u %2u ] ", |
| merge->offsets [index], |
| merge->networks[index]); |
| |
| index += 1; |
| } |
| |
| fprintf(stderr,"\n\n"); |
| } |
| } |
| |
| static |
| void |
| hsg_merge_levels_hint(struct hsg_merge * const merge, bool const autotune) |
| { |
| // clamp against merge levels |
| for (uint32_t level=0; level<MERGE_LEVELS_MAX_LOG2; level++) |
| { |
| // max network |
| uint32_t const n_max = MAX_MACRO(merge->levels[level].networks[0], |
| merge->levels[level].networks[1]); |
| |
| if (n_max <= (merge->rows_bs + hsg_config.thread.xtra)) |
| break; |
| |
| if (autotune) |
| { |
| hsg_config.thread.xtra = n_max - merge->rows_bs; |
| |
| uint32_t const r_total = hsg_config.thread.regs + hsg_config.thread.xtra; |
| uint32_t const r_limit = (hsg_config.type.words == 1) ? 120 : 58; |
| |
| if (r_total <= r_limit) |
| { |
| fprintf(stderr,"autotune: %u + %u\n", |
| hsg_config.thread.regs, |
| hsg_config.thread.xtra); |
| break; |
| } |
| else |
| { |
| fprintf(stderr,"skipping autotune: %u + %u > %u\n", |
| hsg_config.thread.regs, |
| hsg_config.thread.xtra, |
| r_limit); |
| exit(-1); |
| } |
| } |
| |
| fprintf(stderr,"*** HINT *** Try extra registers: %u\n", |
| n_max - merge->rows_bs); |
| |
| exit(-1); |
| } |
| } |
| |
| // |
| // |
| // |
| |
| static |
| struct hsg_op * |
| hsg_op(struct hsg_op * ops, struct hsg_op const opcode) |
| { |
| hsg_op_type_counts[opcode.type] += 1; |
| |
| *ops = opcode; |
| |
| return ops+1; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_exit(struct hsg_op * ops) |
| { |
| return hsg_op(ops,EXIT()); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_end(struct hsg_op * ops) |
| { |
| return hsg_op(ops,END()); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_begin(struct hsg_op * ops) |
| { |
| return hsg_op(ops,BEGIN()); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_else(struct hsg_op * ops) |
| { |
| return hsg_op(ops,ELSE()); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_network_copy(struct hsg_op * ops, |
| struct hsg_network const * const nets, |
| uint32_t const idx, |
| uint32_t const prefix) |
| { |
| uint32_t const len = nets[idx].length; |
| struct hsg_op const * const cxa = nets[idx].network; |
| |
| for (uint32_t ii=0; ii<len; ii++) |
| { |
| struct hsg_op const * const cx = cxa + ii; |
| |
| ops = hsg_op(ops,CMP_XCHG(cx->a,cx->b,prefix)); |
| } |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_sort(struct hsg_op * ops) |
| { |
| uint32_t const idx = hsg_config.thread.regs / 2 - 1; |
| |
| return hsg_network_copy(ops,hsg_networks_sorting,idx,UINT32_MAX); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_merge_prefix(struct hsg_op * ops, uint32_t const network, uint32_t const prefix) |
| { |
| if (network <= 1) |
| return ops; |
| |
| return hsg_network_copy(ops,hsg_networks_merging,network-2,prefix); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_merge(struct hsg_op * ops, uint32_t const network) |
| { |
| return hsg_thread_merge_prefix(ops,network,UINT32_MAX); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_merge_offset_prefix(struct hsg_op * ops, uint32_t const offset, uint32_t const network, uint32_t const prefix) |
| { |
| if (network <= 1) |
| return ops; |
| |
| uint32_t const idx = network - 2; |
| uint32_t const len = hsg_networks_merging[idx].length; |
| struct hsg_op const * const cxa = hsg_networks_merging[idx].network; |
| |
| for (uint32_t ii=0; ii<len; ii++) |
| { |
| struct hsg_op const * const cx = cxa + ii; |
| |
| ops = hsg_op(ops,CMP_XCHG(offset + cx->a,offset + cx->b,prefix)); |
| } |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_merge_offset(struct hsg_op * ops, uint32_t const offset, uint32_t const network) |
| { |
| return hsg_thread_merge_offset_prefix(ops,offset,network,UINT32_MAX); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_merge_left_right_prefix(struct hsg_op * ops, uint32_t const left, uint32_t const right, uint32_t const prefix) |
| { |
| for (uint32_t l=left,r=left+1; r<=left+right; l--,r++) |
| { |
| ops = hsg_op(ops,CMP_XCHG(l,r,prefix)); |
| } |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_thread_merge_left_right(struct hsg_op * ops, uint32_t const left, uint32_t const right) |
| { |
| return hsg_thread_merge_left_right_prefix(ops,left,right,UINT32_MAX); |
| } |
| |
| static |
| struct hsg_op * |
| hsg_warp_half_network(struct hsg_op * ops) |
| { |
| uint32_t const n = hsg_config.thread.regs; |
| |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,CMP_HALF(r-1,r)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_warp_half_downto(struct hsg_op * ops, uint32_t h) |
| { |
| // |
| // *** from h: downto[f/2,1) |
| // **** lane_half(h) |
| // |
| for (; h > 1; h/=2) |
| { |
| ops = hsg_begin(ops); |
| |
| ops = hsg_op(ops,SLAB_HALF(h)); |
| ops = hsg_warp_half_network(ops); |
| |
| ops = hsg_end(ops); |
| } |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_warp_flip_network(struct hsg_op * ops) |
| { |
| uint32_t const n = hsg_config.thread.regs; |
| |
| for (uint32_t r=1; r<=n/2; r++) |
| ops = hsg_op(ops,CMP_FLIP(r-1,r,n+1-r)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_warp_flip(struct hsg_op * ops, uint32_t f) |
| { |
| ops = hsg_begin(ops); |
| |
| ops = hsg_op(ops,SLAB_FLIP(f)); |
| ops = hsg_warp_flip_network(ops); |
| |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_bx_warp_load(struct hsg_op * ops, const int32_t vin_or_vout) |
| { |
| uint32_t const n = hsg_config.thread.regs; |
| |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,BX_REG_GLOBAL_LOAD(r,vin_or_vout)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_bx_warp_store(struct hsg_op * ops) |
| { |
| uint32_t const n = hsg_config.thread.regs; |
| |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,BX_REG_GLOBAL_STORE(r)); |
| |
| return ops; |
| } |
| |
| // |
| // |
| // |
| |
| static |
| struct hsg_op * |
| hsg_warp_transpose(struct hsg_op * ops) |
| { |
| // func proto |
| ops = hsg_op(ops,TRANSPOSE_KERNEL_PROTO()); |
| |
| // begin |
| ops = hsg_begin(ops); |
| |
| // preamble |
| ops = hsg_op(ops,TRANSPOSE_KERNEL_PREAMBLE()); |
| |
| // load |
| ops = hsg_bx_warp_load(ops,1); // 1 = load from vout[] |
| |
| // emit transpose blend and remap macros ... |
| ops = hsg_op(ops,TRANSPOSE_KERNEL_BODY()); |
| |
| // ... done! |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| // |
| // |
| // |
| |
| static |
| struct hsg_op * |
| hsg_warp_half(struct hsg_op * ops, uint32_t const h) |
| { |
| // |
| // *** from h: downto[f/2,1) |
| // **** lane_half(h) |
| // *** thread_merge |
| // |
| ops = hsg_warp_half_downto(ops,h); |
| ops = hsg_thread_merge(ops,hsg_config.thread.regs); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_warp_merge(struct hsg_op * ops) |
| { |
| // |
| // * from f: upto[2,warp.lanes] |
| // ** lane_flip(f) |
| // *** from h: downto[f/2,1) |
| // **** lane_half(h) |
| // *** thread_merge |
| // |
| uint32_t const level = hsg_config.warp.lanes; |
| |
| for (uint32_t f=2; f<=level; f*=2) |
| { |
| ops = hsg_warp_flip(ops,f); |
| ops = hsg_warp_half(ops,f/2); |
| } |
| |
| return ops; |
| } |
| |
| // |
| // |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bc_half_merge_level(struct hsg_op * ops, |
| struct hsg_merge const * const merge, |
| uint32_t const r_lo, |
| uint32_t const s_count) |
| { |
| // guaranteed to be an even network |
| uint32_t const net_even = merge->levels[0].networks[0]; |
| |
| // min of warps in block and remaining horizontal rows |
| uint32_t const active = MIN_MACRO(s_count, net_even); |
| |
| // conditional on blockIdx.x |
| if (active < merge->warps) |
| ops = hsg_op(ops,BX_MERGE_H_PRED(active)); // FIXME BX_MERGE |
| |
| // body begin |
| ops = hsg_begin(ops); |
| |
| // scale for min block |
| uint32_t const scale = net_even >= hsg_config.block.warps_min ? 1 : hsg_config.block.warps_min / net_even; |
| |
| // loop if more smem rows than warps |
| for (uint32_t rr=0; rr<s_count; rr+=active) |
| { |
| // body begin |
| ops = hsg_begin(ops); |
| |
| // skip down slab |
| uint32_t const gmem_base = r_lo - 1 + rr; |
| |
| // load registers horizontally -- striding across slabs |
| for (uint32_t ll=1; ll<=net_even; ll++) |
| ops = hsg_op(ops,BC_REG_GLOBAL_LOAD_LEFT(ll,gmem_base+(ll-1)*hsg_config.thread.regs,0)); |
| |
| // merge all registers |
| ops = hsg_thread_merge_prefix(ops,net_even,0); |
| |
| // if we're looping then there is a base |
| uint32_t const smem_base = rr * net_even * scale; |
| |
| // store all registers |
| for (uint32_t ll=1; ll<=net_even; ll++) |
| ops = hsg_op(ops,BX_REG_SHARED_STORE_LEFT(ll,smem_base+ll-1,0)); |
| |
| // body end |
| ops = hsg_end(ops); |
| } |
| |
| // body end |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_bc_half_merge(struct hsg_op * ops, struct hsg_merge const * const merge) |
| { |
| // |
| // will only be called with merge->warps >= 2 |
| // |
| uint32_t const warps = MAX_MACRO(merge->warps,hsg_config.block.warps_min); |
| |
| // guaranteed to be an even network |
| uint32_t const net_even = merge->levels[0].networks[0]; |
| |
| // set up left SMEM pointer |
| ops = hsg_op(ops,BC_MERGE_H_PREAMBLE(merge->index)); |
| |
| // trim to number of warps in block -- FIXME -- try make this a |
| // multiple of local processor count (Intel = 8, NVIDIA = 4) |
| uint32_t const s_max = merge->rows_bc; |
| |
| // for all the registers |
| for (uint32_t r_lo = 1; r_lo <= hsg_config.thread.regs; r_lo += s_max) |
| { |
| // compute store count |
| uint32_t const r_rem = hsg_config.thread.regs + 1 - r_lo; |
| uint32_t const s_count = MIN_MACRO(s_max,r_rem); |
| |
| // block sync -- can skip if first |
| if (r_lo > 1) |
| ops = hsg_op(ops,BLOCK_SYNC()); |
| |
| // merge loop |
| ops = hsg_bc_half_merge_level(ops,merge,r_lo,s_count); |
| |
| // block sync |
| ops = hsg_op(ops,BLOCK_SYNC()); |
| |
| // load rows from shared |
| for (uint32_t c=0; c<s_count; c++) |
| ops = hsg_op(ops,BC_REG_SHARED_LOAD_V(warps,r_lo+c,c)); |
| } |
| |
| return ops; |
| } |
| |
| // |
| // |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bs_flip_merge_level(struct hsg_op * ops, |
| struct hsg_merge const * const merge, |
| uint32_t const level, |
| uint32_t const s_pairs) |
| { |
| // |
| // Note there are a number of ways to flip merge these warps. There |
| // is a magic number in the merge structure that indicates which |
| // warp to activate as well as what network size to invoke. |
| // |
| // This more complex scheme was used in the past. |
| // |
| // The newest scheme is far dumber/simpler and simply directs a warp |
| // to gather up the network associated with a row and merge them. |
| // |
| // This scheme may use more registers per thread but not all |
| // compilers are high quality. |
| // |
| // If there are more warps than smem row pairs to merge then we |
| // disable the spare warps. |
| // |
| // If there are more row pairs than warps then each warp works on |
| // an equal number of rows. |
| // |
| // Note that it takes two warps to flip merge two smem rows. |
| // |
| // FIXME -- We may want to apply the warp smem "mod" value here to |
| // attempt to balance the load>merge>store operations across the |
| // multiprocessor cores. |
| // |
| // FIXME -- the old scheme attempted to keep all the warps active |
| // but the iteration logic was more complex. See 2016 checkins. |
| // |
| |
| // where are we in computed merge? |
| uint32_t const count = merge->levels[level].count; |
| uint32_t const index = 1 << level; |
| |
| uint32_t s_rows = s_pairs * 2; |
| uint32_t base = 0; |
| |
| while (s_rows > 0) |
| { |
| uint32_t active = merge->warps; |
| |
| // disable warps if necessary |
| if (merge->warps > s_rows) { |
| active = s_rows; |
| ops = hsg_op(ops,BX_MERGE_H_PRED(active)); |
| } |
| |
| // body begin |
| ops = hsg_begin(ops); |
| |
| // how many equal number of rows to merge? |
| uint32_t loops = s_rows / active; |
| |
| // decrement |
| s_rows -= loops * active; |
| |
| for (uint32_t ss=0; ss<loops; ss++) |
| { |
| // load all registers |
| for (uint32_t ii=0; ii<count; ii++) |
| { |
| // body begin |
| ops = hsg_begin(ops); |
| |
| uint32_t const offset = merge->offsets [index+ii]; |
| uint32_t const network = merge->networks[index+ii]; |
| uint32_t const lo = (network + 1) / 2; |
| |
| for (uint32_t ll=1; ll<=lo; ll++) |
| ops = hsg_op(ops,BS_REG_SHARED_LOAD_LEFT(ll,base+offset+ll-1,ii)); |
| |
| for (uint32_t rr=lo+1; rr<=network; rr++) |
| ops = hsg_op(ops,BS_REG_SHARED_LOAD_RIGHT(rr,base+offset+rr-1,ii)); |
| |
| // compare left and right |
| ops = hsg_thread_merge_left_right_prefix(ops,lo,network-lo,ii); |
| |
| // right merging network |
| ops = hsg_thread_merge_offset_prefix(ops,lo,network-lo,ii); |
| |
| // left merging network |
| ops = hsg_thread_merge_prefix(ops,lo,ii); |
| |
| for (uint32_t ll=1; ll<=lo; ll++) |
| ops = hsg_op(ops,BX_REG_SHARED_STORE_LEFT(ll,base+offset+ll-1,ii)); |
| |
| for (uint32_t rr=lo+1; rr<=network; rr++) |
| ops = hsg_op(ops,BS_REG_SHARED_STORE_RIGHT(rr,base+offset+rr-1,ii)); |
| |
| // body end |
| ops = hsg_end(ops); |
| } |
| |
| base += active * merge->warps; |
| } |
| |
| // body end |
| ops = hsg_end(ops); |
| } |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_bs_flip_merge(struct hsg_op * ops, struct hsg_merge const * const merge) |
| { |
| // set up horizontal smem pointer |
| ops = hsg_op(ops,BS_MERGE_H_PREAMBLE(merge->index)); |
| |
| // begin merge |
| uint32_t level = MERGE_LEVELS_MAX_LOG2; |
| |
| while (level-- > 0) |
| { |
| uint32_t const count = merge->levels[level].count; |
| |
| if (count == 0) |
| continue; |
| |
| uint32_t const r_mid = hsg_config.thread.regs/2 + 1; |
| uint32_t const s_pairs_max = merge->rows_bs/2; // this is warp mod |
| |
| // for all the registers |
| for (uint32_t r_lo=1; r_lo<r_mid; r_lo+=s_pairs_max) |
| { |
| uint32_t r_hi = hsg_config.thread.regs + 1 - r_lo; |
| |
| // compute store count |
| uint32_t const s_pairs = MIN_MACRO(s_pairs_max,r_mid - r_lo); |
| |
| // store rows to shared |
| for (uint32_t c=0; c<s_pairs; c++) |
| { |
| ops = hsg_op(ops,BS_REG_SHARED_STORE_V(merge->index,r_lo+c,c*2+0)); |
| ops = hsg_op(ops,BS_REG_SHARED_STORE_V(merge->index,r_hi-c,c*2+1)); |
| } |
| |
| // block sync |
| ops = hsg_op(ops,BLOCK_SYNC()); |
| |
| // merge loop |
| ops = hsg_bs_flip_merge_level(ops,merge,level,s_pairs); |
| |
| // block sync |
| ops = hsg_op(ops,BLOCK_SYNC()); |
| |
| // load rows from shared |
| for (uint32_t c=0; c<s_pairs; c++) |
| { |
| ops = hsg_op(ops,BS_REG_SHARED_LOAD_V(merge->index,r_lo+c,c*2+0)); |
| ops = hsg_op(ops,BS_REG_SHARED_LOAD_V(merge->index,r_hi-c,c*2+1)); |
| } |
| } |
| |
| // conditionally clean -- no-op if equal to number of warps/block |
| if (merge->levels[level].active.b64 != BITS_TO_MASK_64(merge->warps)) |
| ops = hsg_op(ops,BS_ACTIVE_PRED(merge->index,level)); |
| |
| // clean warp |
| ops = hsg_begin(ops); |
| ops = hsg_warp_half(ops,hsg_config.warp.lanes); |
| ops = hsg_end(ops); |
| } |
| |
| return ops; |
| } |
| |
| /* |
| |
| // |
| // DELETE ME WHEN READY |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bs_flip_merge_all(struct hsg_op * ops, const struct hsg_merge * const merge) |
| { |
| for (uint32_t merge_idx=0; merge_idx<MERGE_LEVELS_MAX_LOG2; merge_idx++) |
| { |
| const struct hsg_merge* const m = merge + merge_idx; |
| |
| if (m->warps < 2) |
| break; |
| |
| ops = hsg_op(ops,BS_FRAC_PRED(merge_idx,m->warps)); |
| ops = hsg_begin(ops); |
| ops = hsg_bs_flip_merge(ops,m); |
| ops = hsg_end(ops); |
| } |
| |
| return ops; |
| } |
| */ |
| |
| // |
| // GENERATE SORT KERNEL |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bs_sort(struct hsg_op * ops, struct hsg_merge const * const merge) |
| { |
| // func proto |
| ops = hsg_op(ops,BS_KERNEL_PROTO(merge->index)); |
| |
| // begin |
| ops = hsg_begin(ops); |
| |
| // shared declare |
| ops = hsg_op(ops,BS_KERNEL_PREAMBLE(merge->index)); |
| |
| // load |
| ops = hsg_bx_warp_load(ops,0); // 0 = load from vin[] |
| |
| // thread sorting network |
| ops = hsg_thread_sort(ops); |
| |
| // warp merging network |
| ops = hsg_warp_merge(ops); |
| |
| // slab merging network |
| if (merge->warps > 1) |
| ops = hsg_bs_flip_merge(ops,merge); |
| |
| // store |
| ops = hsg_bx_warp_store(ops); |
| |
| // end |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| // |
| // GENERATE SORT KERNELS |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bs_sort_all(struct hsg_op * ops) |
| { |
| uint32_t merge_idx = MERGE_LEVELS_MAX_LOG2; |
| |
| while (merge_idx-- > 0) |
| { |
| struct hsg_merge const * const m = hsg_merge + merge_idx; |
| |
| if (m->warps == 0) |
| continue; |
| |
| ops = hsg_bs_sort(ops,m); |
| } |
| |
| return ops; |
| } |
| |
| // |
| // GENERATE CLEAN KERNEL FOR A POWER-OF-TWO |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bc_clean(struct hsg_op * ops, struct hsg_merge const * const merge) |
| { |
| // func proto |
| ops = hsg_op(ops,BC_KERNEL_PROTO(merge->index)); |
| |
| // begin |
| ops = hsg_begin(ops); |
| |
| // shared declare |
| ops = hsg_op(ops,BC_KERNEL_PREAMBLE(merge->index)); |
| |
| // if warps == 1 then smem isn't used for merging |
| if (merge->warps == 1) |
| { |
| // load slab directly |
| ops = hsg_bx_warp_load(ops,1); // load from vout[] |
| } |
| else |
| { |
| // block merging network -- strided load of slabs |
| ops = hsg_bc_half_merge(ops,merge); |
| } |
| |
| // clean warp |
| ops = hsg_begin(ops); |
| ops = hsg_warp_half(ops,hsg_config.warp.lanes); |
| ops = hsg_end(ops); |
| |
| // store |
| ops = hsg_bx_warp_store(ops); |
| |
| // end |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| // |
| // GENERATE CLEAN KERNELS |
| // |
| |
| static |
| struct hsg_op * |
| hsg_bc_clean_all(struct hsg_op * ops) |
| { |
| uint32_t merge_idx = MERGE_LEVELS_MAX_LOG2; |
| |
| while (merge_idx-- > 0) |
| { |
| struct hsg_merge const * const m = hsg_merge + merge_idx; |
| |
| if (m->warps == 0) |
| continue; |
| |
| // only generate pow2 clean kernels less than or equal to max |
| // warps in block with the assumption that we would've generated |
| // a wider sort kernel if we could've so a wider clean kernel |
| // isn't a feasible size |
| if (!is_pow2_u32(m->warps)) |
| continue; |
| |
| ops = hsg_bc_clean(ops,m); |
| } |
| |
| return ops; |
| } |
| |
| // |
| // GENERATE FLIP MERGE KERNEL |
| // |
| |
| static |
| struct hsg_op * |
| hsg_fm_thread_load_left(struct hsg_op * ops, uint32_t const n) |
| { |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,FM_REG_GLOBAL_LOAD_LEFT(r,r-1)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_fm_thread_store_left(struct hsg_op * ops, uint32_t const n) |
| { |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,FM_REG_GLOBAL_STORE_LEFT(r,r-1)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_fm_thread_load_right(struct hsg_op * ops, uint32_t const half_span, uint32_t const half_case) |
| { |
| for (uint32_t r=0; r<half_case; r++) |
| ops = hsg_op(ops,FM_REG_GLOBAL_LOAD_RIGHT(r,half_span+1+r)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_fm_thread_store_right(struct hsg_op * ops, uint32_t const half_span, uint32_t const half_case) |
| { |
| for (uint32_t r=0; r<half_case; r++) |
| ops = hsg_op(ops,FM_REG_GLOBAL_STORE_RIGHT(r,half_span+1+r)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_fm_merge(struct hsg_op * ops, |
| uint32_t const scale_log2, |
| uint32_t const span_left, |
| uint32_t const span_right) |
| { |
| // func proto |
| ops = hsg_op(ops,FM_KERNEL_PROTO(scale_log2,msb_idx_u32(pow2_ru_u32(span_right)))); |
| |
| // begin |
| ops = hsg_begin(ops); |
| |
| // preamble for loading/storing |
| ops = hsg_op(ops,FM_KERNEL_PREAMBLE(span_left,span_right)); |
| |
| // load left span |
| ops = hsg_fm_thread_load_left(ops,span_left); |
| |
| // load right span |
| ops = hsg_fm_thread_load_right(ops,span_left,span_right); |
| |
| // compare left and right |
| ops = hsg_thread_merge_left_right(ops,span_left,span_right); |
| |
| // left merging network |
| ops = hsg_thread_merge(ops,span_left); |
| |
| // right merging network |
| ops = hsg_thread_merge_offset(ops,span_left,span_right); |
| |
| // store |
| ops = hsg_fm_thread_store_left(ops,span_left); |
| |
| // store |
| ops = hsg_fm_thread_store_right(ops,span_left,span_right); |
| |
| // end |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_fm_merge_all(struct hsg_op * ops, uint32_t const scale_log2, uint32_t const warps) |
| { |
| uint32_t const span_left = (warps << scale_log2) / 2; |
| uint32_t const span_left_ru = pow2_ru_u32(span_left); |
| |
| for (uint32_t span_right=1; span_right<=span_left_ru; span_right*=2) |
| ops = hsg_fm_merge(ops,scale_log2,span_left,MIN_MACRO(span_left,span_right)); |
| |
| return ops; |
| } |
| |
| // |
| // GENERATE HALF MERGE KERNELS |
| // |
| |
| static |
| struct hsg_op * |
| hsg_hm_thread_load(struct hsg_op * ops, uint32_t const n) |
| { |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,HM_REG_GLOBAL_LOAD(r,r-1)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_hm_thread_store(struct hsg_op * ops, uint32_t const n) |
| { |
| for (uint32_t r=1; r<=n; r++) |
| ops = hsg_op(ops,HM_REG_GLOBAL_STORE(r,r-1)); |
| |
| return ops; |
| } |
| |
| static |
| struct hsg_op * |
| hsg_hm_merge(struct hsg_op * ops, uint32_t const scale_log2, uint32_t const warps_pow2) |
| { |
| uint32_t const span = warps_pow2 << scale_log2; |
| |
| // func proto |
| ops = hsg_op(ops,HM_KERNEL_PROTO(scale_log2)); |
| |
| // begin |
| ops = hsg_begin(ops); |
| |
| // preamble for loading/storing |
| ops = hsg_op(ops,HM_KERNEL_PREAMBLE(span/2)); |
| |
| // load |
| ops = hsg_hm_thread_load(ops,span); |
| |
| // thread merging network |
| ops = hsg_thread_merge(ops,span); |
| |
| // store |
| ops = hsg_hm_thread_store(ops,span); |
| |
| // end |
| ops = hsg_end(ops); |
| |
| return ops; |
| } |
| |
| // |
| // GENERATE MERGE KERNELS |
| // |
| |
| static |
| struct hsg_op * |
| hsg_xm_merge_all(struct hsg_op * ops) |
| { |
| uint32_t const warps = hsg_merge[0].warps; |
| uint32_t const warps_pow2 = pow2_rd_u32(warps); |
| |
| // |
| // GENERATE FLIP MERGE KERNELS |
| // |
| for (uint32_t scale_log2=hsg_config.merge.flip.lo; scale_log2<=hsg_config.merge.flip.hi; scale_log2++) |
| ops = hsg_fm_merge_all(ops,scale_log2,warps); |
| |
| // |
| // GENERATE HALF MERGE KERNELS |
| // |
| for (uint32_t scale_log2=hsg_config.merge.half.lo; scale_log2<=hsg_config.merge.half.hi; scale_log2++) |
| ops = hsg_hm_merge(ops,scale_log2,warps_pow2); |
| |
| return ops; |
| } |
| |
| // |
| // |
| // |
| |
| static |
| struct hsg_op const * |
| hsg_op_translate_depth(hsg_target_pfn target_pfn, |
| struct hsg_target * const target, |
| struct hsg_config const * const config, |
| struct hsg_merge const * const merge, |
| struct hsg_op const * ops, |
| uint32_t const depth) |
| { |
| while (ops->type != HSG_OP_TYPE_EXIT) |
| { |
| switch (ops->type) |
| { |
| case HSG_OP_TYPE_END: |
| target_pfn(target,config,merge,ops,depth-1); |
| return ops + 1; |
| |
| case HSG_OP_TYPE_BEGIN: |
| target_pfn(target,config,merge,ops,depth); |
| ops = hsg_op_translate_depth(target_pfn,target,config,merge,ops+1,depth+1); |
| break; |
| |
| default: |
| target_pfn(target,config,merge,ops++,depth); |
| } |
| } |
| |
| return ops; |
| } |
| |
| static |
| void |
| hsg_op_translate(hsg_target_pfn target_pfn, |
| struct hsg_target * const target, |
| struct hsg_config const * const config, |
| struct hsg_merge const * const merge, |
| struct hsg_op const * ops) |
| { |
| hsg_op_translate_depth(target_pfn,target,config,merge,ops,0); |
| } |
| |
| // |
| // |
| // |
| |
| int |
| main(int argc, char * argv[]) |
| { |
| // |
| // PROCESS OPTIONS |
| // |
| int32_t opt = 0; |
| bool verbose = false; |
| bool autotune = false; |
| char const * arch = "undefined"; |
| struct hsg_target target = { .define = NULL }; |
| |
| while ((opt = getopt(argc,argv,"hva:g:G:s:S:w:b:B:m:M:k:r:x:t:f:F:c:C:p:P:D:z")) != EOF) |
| { |
| switch (opt) |
| { |
| case 'h': |
| fprintf(stderr,"Help goes here...\n"); |
| return EXIT_FAILURE; |
| |
| case 'v': |
| verbose = true; |
| break; |
| |
| case 'a': |
| arch = optarg; |
| break; |
| |
| case 'g': |
| hsg_config.block.smem_min = atoi(optarg); |
| break; |
| |
| case 'G': |
| hsg_config.block.smem_quantum = atoi(optarg); |
| break; |
| |
| case 's': |
| hsg_config.block.smem_bs = atoi(optarg); |
| |
| // set smem_bc if not already set |
| if (hsg_config.block.smem_bc == UINT32_MAX) |
| hsg_config.block.smem_bc = hsg_config.block.smem_bs; |
| break; |
| |
| case 'S': |
| hsg_config.block.smem_bc = atoi(optarg); |
| break; |
| |
| case 'w': |
| hsg_config.warp.lanes = atoi(optarg); |
| hsg_config.warp.lanes_log2 = msb_idx_u32(hsg_config.warp.lanes); |
| break; |
| |
| case 'b': |
| // maximum warps in a workgroup / cta / thread block |
| { |
| uint32_t const warps = atoi(optarg); |
| |
| // must always be even |
| if ((warps & 1) != 0) |
| { |
| fprintf(stderr,"Error: -b must be even.\n"); |
| return EXIT_FAILURE; |
| } |
| |
| hsg_merge[0].index = 0; |
| hsg_merge[0].warps = warps; |
| |
| // set warps_max if not already set |
| if (hsg_config.block.warps_max == UINT32_MAX) |
| hsg_config.block.warps_max = pow2_ru_u32(warps); |
| } |
| break; |
| |
| case 'B': |
| // maximum warps that can fit in a multiprocessor |
| hsg_config.block.warps_max = atoi(optarg); |
| break; |
| |
| case 'm': |
| // blocks using smem barriers must have at least this many warps |
| hsg_config.block.warps_min = atoi(optarg); |
| break; |
| |
| case 'M': |
| // the number of warps necessary to load balance horizontal merging |
| hsg_config.block.warps_mod = atoi(optarg); |
| break; |
| |
| case 'r': |
| { |
| uint32_t const regs = atoi(optarg); |
| |
| if ((regs & 1) != 0) |
| { |
| fprintf(stderr,"Error: -r must be even.\n"); |
| return EXIT_FAILURE; |
| } |
| |
| hsg_config.thread.regs = regs; |
| } |
| break; |
| |
| case 'x': |
| hsg_config.thread.xtra = atoi(optarg); |
| break; |
| |
| case 't': |
| hsg_config.type.words = atoi(optarg); |
| break; |
| |
| case 'f': |
| hsg_config.merge.flip.lo = atoi(optarg); |
| break; |
| |
| case 'F': |
| hsg_config.merge.flip.hi = atoi(optarg); |
| break; |
| |
| case 'c': |
| hsg_config.merge.half.lo = atoi(optarg); |
| break; |
| |
| case 'C': |
| hsg_config.merge.half.hi = atoi(optarg); |
| break; |
| |
| case 'p': |
| hsg_config.merge.flip.warps = atoi(optarg); |
| break; |
| |
| case 'P': |
| hsg_config.merge.half.warps = atoi(optarg); |
| break; |
| |
| case 'D': |
| target.define = optarg; |
| break; |
| |
| case 'z': |
| autotune = true; |
| break; |
| } |
| } |
| |
| // |
| // INIT MERGE |
| // |
| uint32_t const warps_ru_pow2 = pow2_ru_u32(hsg_merge[0].warps); |
| |
| for (uint32_t ii=1; ii<MERGE_LEVELS_MAX_LOG2; ii++) |
| { |
| hsg_merge[ii].index = ii; |
| hsg_merge[ii].warps = warps_ru_pow2 >> ii; |
| } |
| |
| // |
| // WHICH ARCH TARGET? |
| // |
| hsg_target_pfn hsg_target_pfn; |
| |
| if (strcmp(arch,"debug") == 0) |
| hsg_target_pfn = hsg_target_debug; |
| else if (strcmp(arch,"cuda") == 0) |
| hsg_target_pfn = hsg_target_cuda; |
| else if (strcmp(arch,"opencl") == 0) |
| hsg_target_pfn = hsg_target_opencl; |
| else if (strcmp(arch,"glsl") == 0) |
| hsg_target_pfn = hsg_target_glsl; |
| else { |
| fprintf(stderr,"Invalid arch: %s\n",arch); |
| exit(EXIT_FAILURE); |
| } |
| |
| if (verbose) |
| fprintf(stderr,"Target: %s\n",arch); |
| |
| // |
| // INIT SMEM KEY ALLOCATION |
| // |
| hsg_config_init_shared(); |
| |
| // |
| // INIT MERGE MAGIC |
| // |
| for (uint32_t ii=0; ii<MERGE_LEVELS_MAX_LOG2; ii++) |
| { |
| struct hsg_merge * const merge = hsg_merge + ii; |
| |
| if (merge->warps == 0) |
| break; |
| |
| fprintf(stderr,">>> Generating: %1u %5u %5u %3u %3u ...\n", |
| hsg_config.type.words, |
| hsg_config.block.smem_bs, |
| hsg_config.block.smem_bc, |
| hsg_config.thread.regs, |
| merge->warps); |
| |
| hsg_merge_levels_init_shared(merge); |
| |
| hsg_merge_levels_init_1(merge,merge->warps,0,0); |
| |
| hsg_merge_levels_hint(merge,autotune); |
| |
| // |
| // THESE ARE FOR DEBUG/INSPECTION |
| // |
| if (verbose) |
| { |
| hsg_merge_levels_debug(merge); |
| } |
| } |
| |
| if (verbose) |
| fprintf(stderr,"\n\n"); |
| |
| // |
| // GENERATE THE OPCODES |
| // |
| uint32_t const op_count = 1<<17; |
| struct hsg_op * const ops_begin = malloc(sizeof(*ops_begin) * op_count); |
| struct hsg_op * ops = ops_begin; |
| |
| // |
| // OPEN INITIAL FILES AND APPEND HEADER |
| // |
| ops = hsg_op(ops,TARGET_BEGIN()); |
| |
| // |
| // GENERATE SORT KERNEL |
| // |
| ops = hsg_bs_sort_all(ops); |
| |
| // |
| // GENERATE CLEAN KERNELS |
| // |
| ops = hsg_bc_clean_all(ops); |
| |
| // |
| // GENERATE MERGE KERNELS |
| // |
| ops = hsg_xm_merge_all(ops); |
| |
| // |
| // GENERATE TRANSPOSE KERNEL |
| // |
| ops = hsg_warp_transpose(ops); |
| |
| // |
| // APPEND FOOTER AND CLOSE INITIAL FILES |
| // |
| ops = hsg_op(ops,TARGET_END()); |
| |
| // |
| // ... WE'RE DONE! |
| // |
| ops = hsg_exit(ops); |
| |
| // |
| // APPLY TARGET TRANSLATOR TO ACCUMULATED OPS |
| // |
| hsg_op_translate(hsg_target_pfn,&target,&hsg_config,hsg_merge,ops_begin); |
| |
| // |
| // DUMP INSTRUCTION COUNTS |
| // |
| if (verbose) |
| hsg_op_debug(); |
| |
| return EXIT_SUCCESS; |
| } |
| |
| // |
| // |
| // |
