Spaces:
Running
Running
File size: 2,061 Bytes
81efcf0 |
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 |
// Copyright (C) 2018-2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
//
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#pragma OPENCL EXTENSION cl_khr_extended_async_copies : enable
__kernel void cvtu8f16(__global const uchar *restrict src, __global half *restrict dst, float scale, float bias)
{
__local uchar local_src[8 * 1024];
__local half local_dst[8 * 1024];
event_t e1 = async_work_group_copy_3D3D(
local_src, // dst
src + get_group_id(0) * get_local_size(0) + get_group_id(1) * get_local_size(1) * get_global_size(0)
+ get_group_id(2) * get_local_size(2) * get_global_size(0) * get_global_size(1), // src
get_local_size(0), // num_elements_per_line
get_local_size(0) * get_local_size(1) / (get_local_size(0)), // num_lines
get_global_size(0) - get_local_size(0), // src_line_stride
0, // dst_line_stride
get_local_size(2), // num planes
get_global_size(0) * (get_global_size(1) - get_local_size(1)), // src plane stride
0, // dst plane stride
0);
wait_group_events(1, &e1);
size_t idx = get_local_id(0)
+ get_local_id(1) * get_local_size(0)
+ get_local_id(2) * get_local_size(0) * get_local_size(1);
local_dst[idx] = convert_half(local_src[idx]) * (half)scale + (half)bias;
barrier(CLK_LOCAL_MEM_FENCE);
event_t e2 = async_work_group_copy_3D3D(
dst + get_group_id(0) * get_local_size(0) + get_group_id(1) * get_local_size(1) * get_global_size(0)
+ get_group_id(2) * get_local_size(2) * get_global_size(0) * get_global_size(1), // dst
local_dst, // src
get_local_size(0), // num_elements_per_line
get_local_size(1), // num_lines
0, // src_line_stride
get_global_size(0) - get_local_size(0), // dst_line_stride
get_local_size(2), // num_planes
0, // src_plane_stride
get_global_size(0) * (get_global_size(1) - get_local_size(1)), // dst_plane_stride
0);
wait_group_events(1, &e2);
}
|