| /* |
| * 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 "kernel_cl_12.h" |
| |
| // |
| // BEST TO RUN THESE ON AN OUT-OF-ORDER CQ |
| // |
| |
| __kernel |
| SKC_BP_INIT_IDS_KERNEL_ATTRIBS |
| void |
| skc_kernel_block_pool_init_ids(__global uint * const ids, uint const bp_size) |
| { |
| uint const gid = get_global_id(0); |
| |
| // |
| // FIXME -- TUNE FOR ARCH -- evaluate if it's much faster to |
| // accomplish this with fewer threads and using either IPC and/or |
| // vector stores -- it should be on certain architectures! |
| // |
| |
| // |
| // initialize pool with sequence |
| // |
| if (gid < bp_size) |
| ids[gid] = gid * SKC_DEVICE_SUBBLOCKS_PER_BLOCK; |
| } |
| |
| // |
| // |
| // |
| |
| __kernel |
| SKC_BP_INIT_ATOMICS_KERNEL_ATTRIBS |
| void |
| skc_kernel_block_pool_init_atomics(__global uint * const bp_atomics, uint const bp_size) |
| { |
| // the version test is to squelch a bug with the Intel OpenCL CPU |
| // compiler declaring it supports the cl_intel_subgroups extension |
| #if defined(cl_intel_subgroups) || defined (cl_khr_subgroups) |
| uint const tid = get_sub_group_local_id(); |
| #else |
| uint const tid = get_local_id(0); |
| #endif |
| |
| // |
| // launch two threads and store [ 0, bp_size ] |
| // |
| bp_atomics[tid] = tid * bp_size; |
| } |
| |
| // |
| // |
| // |
