diff options
Diffstat (limited to 'src/compute/skc/platforms/cl_12/kernels/paths_copy.cl')
| -rw-r--r-- | src/compute/skc/platforms/cl_12/kernels/paths_copy.cl | 543 |
1 files changed, 543 insertions, 0 deletions
diff --git a/src/compute/skc/platforms/cl_12/kernels/paths_copy.cl b/src/compute/skc/platforms/cl_12/kernels/paths_copy.cl new file mode 100644 index 0000000000..302ea14af2 --- /dev/null +++ b/src/compute/skc/platforms/cl_12/kernels/paths_copy.cl @@ -0,0 +1,543 @@ +/*
+ * Copyright 2017 Google Inc.
+ *
+ * Use of this source code is governed by a BSD-style license that can
+ * be found in the LICENSE file.
+ *
+ */
+
+//
+//
+//
+
+#include "path.h"
+#include "block_pool_cl.h"
+#include "path_builder_cl_12.h"
+#include "device_cl_12.h"
+
+//
+//
+//
+
+#if 0
+
+//
+// SIMD AVX2
+//
+
+#define SKC_PATHS_COPY_WORDS_PER_ELEM 8
+#define SKC_PATHS_COPY_SUBGROUP_SIZE 1
+#define SKC_PATHS_COPY_KERNEL_ATTRIBUTES
+
+typedef skc_uint8 skc_paths_copy_elem;
+typedef skc_uint8 skc_pb_idx_v;
+
+#define SKC_PATHS_COPY_ELEM_EXPAND() SKC_EXPAND_8()
+
+#define SKC_IS_NOT_PATH_HEAD(sg,I) ((sg) + I >= SKC_PATH_HEAD_WORDS)
+
+#endif
+
+//
+//
+//
+
+#define SKC_PATHS_COPY_SUBGROUP_SIZE_MASK (SKC_PATHS_COPY_SUBGROUP_SIZE - 1)
+#define SKC_PATHS_COPY_ELEMS_PER_BLOCK (SKC_DEVICE_BLOCK_WORDS / SKC_PATHS_COPY_ELEM_WORDS)
+#define SKC_PATHS_COPY_ELEMS_PER_SUBBLOCK (SKC_DEVICE_SUBBLOCK_WORDS / SKC_PATHS_COPY_ELEM_WORDS)
+#define SKC_PATHS_COPY_ELEMS_PER_THREAD (SKC_PATHS_COPY_ELEMS_PER_BLOCK / SKC_PATHS_COPY_SUBGROUP_SIZE)
+
+// FIXME -- use SUBGROUP terminology everywhere
+#define SKC_PATHS_COPY_SUBGROUP_WORDS (SKC_PATHS_COPY_SUBGROUP_SIZE * SKC_PATHS_COPY_ELEM_WORDS)
+
+//
+//
+//
+
+#define SKC_PATHS_COPY_ELEMS_BEFORE_HEADER \
+ (SKC_PATHS_COPY_SUBGROUP_SIZE * ((SKC_PATH_HEAD_WORDS / SKC_PATHS_COPY_ELEM_WORDS) / SKC_PATHS_COPY_SUBGROUP_WORDS))
+
+#define SKC_PATHS_COPY_ELEMS_INCLUDING_HEADER \
+ (SKC_PATHS_COPY_SUBGROUP_SIZE * ((SKC_PATH_HEAD_WORDS + SKC_PATHS_COPY_SUBGROUP_WORDS - 1) / SKC_PATHS_COPY_SUBGROUP_WORDS))
+
+// #define SKC_PATHS_COPY_HEAD_ELEMS ((SKC_PATH_HEAD_WORDS + SKC_PATHS_COPY_ELEM_WORDS - 1) / SKC_PATHS_COPY_ELEM_WORDS)
+
+//
+//
+//
+
+//
+// BIT-FIELD EXTRACT/INSERT ARE NOT AVAILABLE IN OPENCL
+//
+
+#define SKC_CMD_PATHS_COPY_ONE_BITS (SKC_TAGGED_BLOCK_ID_BITS_TAG + SKC_DEVICE_SUBBLOCK_WORDS_LOG2)
+
+#define SKC_CMD_PATHS_COPY_ONE_MASK SKC_BITS_TO_MASK(SKC_CMD_PATHS_COPY_ONE_BITS)
+
+#define SKC_CMD_PATHS_COPY_ONE (1u << SKC_CMD_PATHS_COPY_ONE_BITS)
+
+#define SKC_CMD_PATHS_COPY_GET_TAG(ti) SKC_TAGGED_BLOCK_ID_GET_TAG(ti)
+
+#define SKC_CMD_PATHS_COPY_GET_ROLLING(ti) ((ti) >> SKC_CMD_PATHS_COPY_ONE_BITS)
+
+#define SKC_CMD_PATHS_COPY_UPDATE_ROLLING(ti,b) (((ti) & SKC_CMD_PATHS_COPY_ONE_MASK) | ((b) << SKC_TAGGED_BLOCK_ID_BITS_TAG))
+
+//
+//
+//
+
+skc_uint
+skc_sub_group_local_id()
+{
+#if SKC_PATHS_COPY_SUBGROUP_SIZE > 1
+ return get_sub_group_local_id();
+#else
+ return 0;
+#endif
+}
+
+//
+// convert an atomic read counter offset to a block id
+//
+
+skc_block_id_t
+skc_bp_off_to_id(__global skc_block_id_t const * const bp_ids,
+ skc_uint const bp_idx_mask,
+ skc_uint const bp_reads,
+ skc_uint const bp_off)
+{
+ skc_uint const bp_idx = (bp_reads + bp_off) & bp_idx_mask;
+
+ return bp_ids[bp_idx];
+}
+
+//
+//
+//
+
+void
+skc_copy_segs(__global skc_paths_copy_elem * const bp_elems, // to
+ skc_uint const bp_elems_idx,
+ __global skc_paths_copy_elem const * const pb_elems, // from
+ skc_uint const pb_elems_idx)
+{
+ for (skc_uint ii=0; ii<SKC_PATHS_COPY_ELEMS_PER_BLOCK; ii+=SKC_PATHS_COPY_SUBGROUP_SIZE)
+ {
+ (bp_elems+bp_elems_idx)[ii] = (pb_elems+pb_elems_idx)[ii];
+ }
+
+#if 0
+ //
+ // NOTE THIS IS PRINTING 8 ROWS
+ //
+ printf("%5u : (%8u) : { { %5.0f, %5.0f }, { %5.0f, %5.0f } },\n",
+ (skc_uint)get_global_id(0),pb_elems_idx,
+ as_float((pb_elems+pb_elems_idx)[0*SKC_PATHS_COPY_SUBGROUP_SIZE]),
+ as_float((pb_elems+pb_elems_idx)[1*SKC_PATHS_COPY_SUBGROUP_SIZE]),
+ as_float((pb_elems+pb_elems_idx)[2*SKC_PATHS_COPY_SUBGROUP_SIZE]),
+ as_float((pb_elems+pb_elems_idx)[3*SKC_PATHS_COPY_SUBGROUP_SIZE]));
+ printf("%5u : (%8u) : { { %5.0f, %5.0f }, { %5.0f, %5.0f } },\n",
+ (skc_uint)get_global_id(0),pb_elems_idx,
+ as_float((pb_elems+pb_elems_idx)[4*SKC_PATHS_COPY_SUBGROUP_SIZE]),
+ as_float((pb_elems+pb_elems_idx)[5*SKC_PATHS_COPY_SUBGROUP_SIZE]),
+ as_float((pb_elems+pb_elems_idx)[6*SKC_PATHS_COPY_SUBGROUP_SIZE]),
+ as_float((pb_elems+pb_elems_idx)[7*SKC_PATHS_COPY_SUBGROUP_SIZE]));
+#endif
+}
+
+//
+//
+//
+
+void
+skc_copy_node(__global skc_paths_copy_elem * const bp_elems, // to
+ skc_uint const bp_elems_idx,
+ __global skc_block_id_t const * const bp_ids,
+ skc_uint const bp_reads,
+ skc_uint const bp_idx_mask,
+ __global skc_paths_copy_elem const * const pb_elems, // from
+ skc_uint const pb_elems_idx,
+ skc_uint const pb_rolling)
+{
+ //
+ // remap block id tags bp_elems the host-side rolling counter pb_elems a
+ // device-side block pool id
+ //
+ for (skc_uint ii=0; ii<SKC_PATHS_COPY_ELEMS_PER_BLOCK; ii+=SKC_PATHS_COPY_SUBGROUP_SIZE)
+ {
+ // load block_id_tag words
+ skc_paths_copy_elem elem = (pb_elems + pb_elems_idx)[ii];
+
+ // calculate ahead of time -- if elem was invalid then bp_idx is definitely invalid
+ skc_pb_idx_v const bp_idx = (bp_reads + SKC_CMD_PATHS_COPY_GET_ROLLING(elem - pb_rolling)) & bp_idx_mask;
+
+ // FIXME ^^^^^ THE IDX PROBABLY DOESN'T NEED TO BE SHIFTED TWICE AND WE CAN SAVE A FEW INSTRUCTIONS
+
+ //
+ // FIXME -- SIMD can be fully parallelized since a bp_ids[] load
+ // will _always_ be safe as long as we don't use the loaded
+ // value! So... fix UPDATE_ROLLING to be SIMD-friendly instead
+ // of iterating over the vector components.
+ //
+
+ // only convert if original elem is not invalid
+
+#undef SKC_EXPAND_X
+#define SKC_EXPAND_X(I,S,C,P,R) \
+ if (elem C != SKC_TAGGED_BLOCK_ID_INVALID) { \
+ skc_block_id_t const b = bp_ids[bp_idx C]; \
+ elem C = SKC_CMD_PATHS_COPY_UPDATE_ROLLING(elem C,b); \
+ }
+
+ // printf("%2u: < %8X, %8X, %8X >\n",ii,bp_idx,b,elem C);
+
+ SKC_PATHS_COPY_ELEM_EXPAND();
+
+ // store the elem back
+ (bp_elems+bp_elems_idx)[ii] = elem;
+ }
+}
+
+//
+//
+//
+
+void
+skc_host_map_update(__global skc_uint * const host_map,
+ skc_uint const block,
+ skc_paths_copy_elem const elem)
+{
+ //
+ // write first elem to map -- FIXME -- this is a little nasty
+ // because it relies on the the host handle always being the first
+ // word in the path header.
+ //
+ // OTOH, this is not unreasonable. The alternative is to have a
+ // separate kernel initializing the map.
+ //
+#if SKC_PATHS_COPY_SUBGROUP_SIZE > 1
+ if (get_sub_group_local_id() == SKC_PATH_HEAD_OFFSET_HANDLE)
+#endif
+ {
+#if SKC_PATHS_COPY_ELEM_WORDS == 1
+ host_map[elem] = block;
+#if 0
+ printf("[%u] = %u\n",elem,block);
+#endif
+#else
+ host_map[elem.SKC_CONCAT(s,SKC_PATH_HEAD_OFFSET_HANDLE)] = block;
+#endif
+ }
+}
+
+//
+//
+//
+
+void
+skc_copy_head(__global skc_uint * const host_map,
+ skc_uint const block,
+ __global skc_paths_copy_elem * const bp_elems, // to
+ skc_uint const bp_elems_idx,
+ __global skc_block_id_t const * const bp_ids,
+ skc_uint const bp_reads,
+ skc_uint const bp_idx_mask,
+ __global skc_paths_copy_elem const * const pb_elems, // from
+ skc_uint const pb_elems_idx,
+ skc_uint const pb_rolling)
+{
+ //
+ // if there are more path header words than there are
+ // threads-per-block then we can just copy the initial header words
+ //
+#if ( SKC_PATHS_COPY_ELEMS_BEFORE_HEADER > 0 )
+ for (skc_uint ii=0; ii<SKC_PATHS_COPY_ELEMS_BEFORE_HEADER; ii+=SKC_PATHS_COPY_SUBGROUP_SIZE)
+ {
+ skc_paths_copy_elem const elem = (pb_elems+pb_elems_idx)[ii];
+
+ (bp_elems+bp_elems_idx)[ii] = elem;
+
+ if (ii == 0) {
+ skc_host_map_update(host_map,block,elem);
+ }
+ }
+#endif
+
+ //
+ // this is similar to copy node but the first H words of the path
+ // header are not modified and simply copied
+ //
+ for (skc_uint ii=SKC_PATHS_COPY_ELEMS_BEFORE_HEADER; ii<SKC_PATHS_COPY_ELEMS_INCLUDING_HEADER; ii+=SKC_PATHS_COPY_SUBGROUP_SIZE)
+ {
+ skc_paths_copy_elem elem = (pb_elems+pb_elems_idx)[ii];
+
+#if ( SKC_PATHS_COPY_ELEMS_BEFORE_HEADER == 0 )
+ if (ii == 0) {
+ skc_host_map_update(host_map,block,elem);
+ }
+#endif
+ // calculate ahead of time -- if elem was invalid then bp_idx is definitely invalid
+ skc_pb_idx_v const bp_idx = (bp_reads + SKC_CMD_PATHS_COPY_GET_ROLLING(elem - pb_rolling)) & bp_idx_mask;
+
+ //
+ // FIXME -- SIMD can be fully parallelized since a bp_ids[] load
+ // will _always_ be safe as long as we don't use the loaded
+ // value! So... fix UPDATE_ROLLING to be SIMD-friendly instead
+ // of iterating over the vector components.
+ //
+
+ // FIXME ^^^^^ THE IDX PROBABLY DOESN'T NEED TO BE SHIFTED TWICE AND WE CAN SAVE A FEW INSTRUCTIONS
+
+ // FIXME -- MIX MIX MIX MIX / SELECT
+
+ // only convert if original elem is not invalid
+#undef SKC_EXPAND_X
+#define SKC_EXPAND_X(I,S,C,P,R) \
+ if (SKC_IS_NOT_PATH_HEAD(ii,I) && (elem C != SKC_TAGGED_BLOCK_ID_INVALID)) { \
+ skc_block_id_t const b = bp_ids[bp_idx C]; \
+ elem C = SKC_CMD_PATHS_COPY_UPDATE_ROLLING(elem C,b); \
+ }
+
+ // printf("%2u: ( %8X, %8X, %8X )\n",ii,bp_idx,b,elem C);
+
+ SKC_PATHS_COPY_ELEM_EXPAND();
+
+ // store the elem back
+ (bp_elems+bp_elems_idx)[ii] = elem;
+ }
+
+ //
+ // the remaining words are treated like a node
+ //
+ for (skc_uint ii=SKC_PATHS_COPY_ELEMS_INCLUDING_HEADER; ii<SKC_PATHS_COPY_ELEMS_PER_BLOCK; ii+=SKC_PATHS_COPY_SUBGROUP_SIZE)
+ {
+ // load block_id_tag words
+ skc_paths_copy_elem elem = (pb_elems+pb_elems_idx)[ii];
+
+ // calculate ahead of time
+ skc_pb_idx_v const bp_idx = (bp_reads + SKC_CMD_PATHS_COPY_GET_ROLLING(elem - pb_rolling)) & bp_idx_mask;
+
+ //
+ // FIXME -- SIMD can be fully parallelized since a bp_ids[] load
+ // will _always_ be safe as long as we don't use the loaded
+ // value! So... fix UPDATE_ROLLING to be SIMD-friendly instead
+ // of iterating over the vector components.
+ //
+
+ // FIXME ^^^^^ THE IDX PROBABLY DOESN'T NEED TO BE SHIFTED TWICE AND WE CAN SAVE A FEW INSTRUCTIONS
+
+ // only convert if original elem is not invalid
+#undef SKC_EXPAND_X
+#define SKC_EXPAND_X(I,S,C,P,R) \
+ if (elem C != SKC_TAGGED_BLOCK_ID_INVALID) { \
+ skc_block_id_t const b = bp_ids[bp_idx C]; \
+ elem C = SKC_CMD_PATHS_COPY_UPDATE_ROLLING(elem C,b); \
+ }
+
+ // printf("%2u: [ %8X, %8X, %8X ]\n",ii,bp_idx,b,elem C);
+
+ SKC_PATHS_COPY_ELEM_EXPAND();
+
+ // store the elem
+ (bp_elems+bp_elems_idx)[ii] = elem;
+ }
+}
+
+//
+// FIXME -- pack some of these constant integer args in a vec or struct
+//
+
+__kernel
+SKC_PATHS_COPY_KERNEL_ATTRIBS
+void
+skc_kernel_paths_copy
+(__global skc_uint * const host_map,
+
+ __global skc_block_id_t const * const bp_ids,
+ __global skc_paths_copy_elem * const bp_elems,
+ skc_uint const bp_idx_mask, // pow2 modulo mask for block pool ring
+
+ __global skc_uint const * const bp_alloc, // block pool ring base
+ skc_uint const bp_alloc_idx,// which subbuf
+
+ __global union skc_tagged_block_id const * const pb_cmds,
+ __global skc_paths_copy_elem const * const pb_elems,
+
+ skc_uint const pb_size, // # of commands/blocks in buffer
+ skc_uint const pb_rolling, // shifted rolling counter base
+
+ skc_uint const pb_prev_from,
+ skc_uint const pb_prev_span,
+ skc_uint const pb_curr_from)
+{
+ //
+ // THERE ARE 3 TYPES OF PATH COPYING COMMANDS:
+ //
+ // - HEAD
+ // - NODE
+ // - SEGS
+ //
+ // THESE ARE SUBGROUP ORIENTED KERNELS
+ //
+ // A SUBGROUP CAN OPERATE ON [1,N] BLOCKS
+ //
+
+ //
+ // It's likely that peak bandwidth is achievable with a single
+ // workgroup.
+ //
+ // So let's keep the grids modestly sized and for simplicity and
+ // portability, let's assume that a single workgroup can perform all
+ // steps in the copy.
+ //
+ // Launch as large of a workgroup as possiblex
+ //
+ // 1. ATOMICALLY ALLOCATE BLOCKS BP_ELEMS POOL
+ // 2. CONVERT COMMANDS IN PB_ELEMS BLOCK OFFSETS
+ // 3. FOR EACH COMMAND:
+ // - HEAD: SAVED HEAD ID PB_ELEMS MAP. CONVERT AND COPY H INDICES.
+ // - NODE: CONVERT AND COPY B INDICES
+ // - SEGS: BULK COPY
+ //
+ // B : number of words in block -- always pow2
+ // W : intelligently/arbitrarily chosen factor of B -- always pow2
+ //
+
+ //
+ // There are several approaches to processing the commands:
+ //
+ // 1. B threads are responsible for one block. All threads broadcast
+ // load a single command word. Workgroup size must be a facpb_elemsr of
+ // B.
+ //
+ // 2. W threads process an entire block. W will typically be the
+ // device's subgroup/warp/wave width. W threads broadcast load a
+ // single command word.
+ //
+ // 3. W threads process W blocks. W threads load W command words and
+ // process W blocks.
+ //
+ // Clearly (1) has low I/O intensity but will achieve high
+ // parallelism by activating the most possible threads. The downside
+ // of this kind of approach is that the kernel will occupy even a
+ // large GPU with low intensity work and reduce opportunities for
+ // concurrent kernel execution (of other kernels).
+ //
+ // See Vasily Volkov's CUDA presentation describing these tradeoffs.
+ //
+ // Note that there are many other approaches. For example, similar
+ // pb_elems (1) but each thread loads a pow2 vector of block data.
+ //
+
+ // load the copied atomic read "base" from gmem
+ skc_uint const bp_reads = bp_alloc[bp_alloc_idx];
+ // will always be less than 2^32
+ skc_uint const gid = get_global_id(0);
+ // every subgroup/simd that will work on the block loads the same command
+ skc_uint const sg_idx = gid / SKC_PATHS_COPY_SUBGROUP_SIZE;
+ // path builder data can be spread across two spans
+ skc_uint pb_idx = sg_idx + ((sg_idx < pb_prev_span) ? pb_prev_from : pb_curr_from);
+
+ // no need pb_elems make this branchless
+ if (pb_idx >= pb_size)
+ pb_idx -= pb_size;
+
+ // broadcast load the command
+ union skc_tagged_block_id const pb_cmd = pb_cmds[pb_idx];
+
+ // what do we want pb_elems do with this block?
+ skc_cmd_paths_copy_tag const tag = SKC_CMD_PATHS_COPY_GET_TAG(pb_cmd.u32);
+
+ // compute offset from rolling base to get index into block pool ring allocation
+ skc_uint const bp_off = SKC_CMD_PATHS_COPY_GET_ROLLING(pb_cmd.u32 - pb_rolling);
+
+ // convert the pb_cmd's offset counter pb_elems a block id
+ skc_block_id_t const block = skc_bp_off_to_id(bp_ids,bp_idx_mask,bp_reads,bp_off);
+
+#if 0
+ if (get_sub_group_local_id() == 0) {
+ printf("bp_off/reads = %u / %u\n",bp_off,bp_reads);
+ printf("< %8u >\n",block);
+ }
+#endif
+
+ // FIXME -- could make this 0 for SIMD, gid&mask or get_sub_group_local_id()
+ skc_uint const tid = gid & SKC_PATHS_COPY_SUBGROUP_SIZE_MASK;
+
+ // calculate bp_elems (to) / pb_elems (from)
+ skc_uint const bp_elems_idx = block * SKC_PATHS_COPY_ELEMS_PER_SUBBLOCK + tid;
+ skc_uint const pb_elems_idx = pb_idx * SKC_PATHS_COPY_ELEMS_PER_BLOCK + tid;
+
+ if (tag == SKC_CMD_PATHS_COPY_TAG_SEGS)
+ {
+#if 0
+ if (tid == 0)
+ printf("%3u, segs\n",bp_off);
+#endif
+ skc_copy_segs(bp_elems,
+ bp_elems_idx,
+ pb_elems,
+ pb_elems_idx);
+ }
+ else if (tag == SKC_CMD_PATHS_COPY_TAG_NODE)
+ {
+#if 0
+ if (tid == 0)
+ printf("%3u, NODE\n",bp_off);
+#endif
+ skc_copy_node(bp_elems, // to
+ bp_elems_idx,
+ bp_ids,
+ bp_reads,
+ bp_idx_mask,
+ pb_elems, // from
+ pb_elems_idx,
+ pb_rolling);
+ }
+ else // ( tag == SKC_CMD_PATHS_COPY_TAG_HEAD)
+ {
+#if 0
+ if (tid == 0)
+ printf("%3u, HEAD\n",bp_off);
+#endif
+ skc_copy_head(host_map,
+ block,
+ bp_elems, // to
+ bp_elems_idx,
+ bp_ids,
+ bp_reads,
+ bp_idx_mask,
+ pb_elems, // from
+ pb_elems_idx,
+ pb_rolling);
+ }
+}
+
+//
+//
+//
+
+__kernel
+SKC_PATHS_ALLOC_KERNEL_ATTRIBS
+void
+skc_kernel_paths_alloc(__global skc_uint volatile * const bp_atomics,
+ __global skc_uint * const bp_alloc,
+ skc_uint const bp_alloc_idx,
+ skc_uint const pb_cmd_count)
+{
+ //
+ // allocate blocks in block pool
+ //
+ skc_uint const reads = atomic_add(bp_atomics+SKC_BP_ATOMIC_OFFSET_READS,pb_cmd_count);
+
+ // store in slot
+ bp_alloc[bp_alloc_idx] = reads;
+
+#if 0
+ printf("pc: %8u + %u\n",reads,pb_cmd_count);
+#endif
+}
+
+//
+//
+//
|