From 65d11a2120b722b60f438cf8a1954cb09c4ff369 Mon Sep 17 00:00:00 2001 From: agrigori Date: Thu, 26 Mar 2020 21:07:23 +0300 Subject: [PATCH] Adressed comments #7 --- .../oneapi/cl_kernels/dbscan_cl_kernels.cl | 10 +++++----- .../dbscan_dense_default_batch_ucapi_impl.i | 20 +++++++++---------- 2 files changed, 15 insertions(+), 15 deletions(-) diff --git a/algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl b/algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl index 26c4f133b7c..926d60e9c63 100644 --- a/algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl +++ b/algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl @@ -142,14 +142,14 @@ DECLARE_SOURCE( } } - __kernel void push_to_queue(__global const algorithmFPType * distances, __global const int * offsets, __global int * assignments, - __global int * queue, int queue_end, int point_id, int cluster_id, int chunk_offset, int chunk_size, - algorithmFPType eps, int num_points) + __kernel void push_points_to_queue(__global const algorithmFPType * distances, __global const int * offsets, + int queue_end, int point_id, int cluster_id, int first_chunk_offset, int chunk_size, + algorithmFPType eps, int num_points, __global int * assignments, __global int * queue) { const int subgroup_index = get_global_id(0) * get_num_sub_groups() + get_sub_group_id(); const int offset = subgroup_index * chunk_size; - point_id = chunk_offset < 0 ? point_id : queue[point_id]; - const int dist_offset = offset + (chunk_offset < 0 ? 0 : chunk_offset); + point_id = first_chunk_offset < 0 ? point_id : queue[point_id]; + const int dist_offset = offset + (first_chunk_offset < 0 ? 0 : first_chunk_offset); if (offset >= num_points) return; const int subgroup_point_number = num_points - offset < chunk_size ? num_points - offset : chunk_size; diff --git a/algorithms/kernel/dbscan/oneapi/dbscan_dense_default_batch_ucapi_impl.i b/algorithms/kernel/dbscan/oneapi/dbscan_dense_default_batch_ucapi_impl.i index 873d8588941..c835a569fd3 100644 --- a/algorithms/kernel/dbscan/oneapi/dbscan_dense_default_batch_ucapi_impl.i +++ b/algorithms/kernel/dbscan/oneapi/dbscan_dense_default_batch_ucapi_impl.i @@ -274,7 +274,7 @@ services::Status DBSCANBatchKernelUCAPI::pushNeighborsToQueue(c auto & context = Environment::getInstance()->getDefaultExecutionContext(); auto & kernel_factory = context.getClKernelFactory(); DAAL_CHECK_STATUS_VAR(buildProgram(kernel_factory)); - auto kernel = kernel_factory.getKernel("push_to_queue", &st); + auto kernel = kernel_factory.getKernel("push_points_to_queue", &st); DAAL_CHECK_STATUS_VAR(st); uint32_t chunkSize = nRows / numberOfChunks + uint32_t(bool(nRows % numberOfChunks)); @@ -282,15 +282,15 @@ services::Status DBSCANBatchKernelUCAPI::pushNeighborsToQueue(c KernelArguments args(11); args.set(0, distances, AccessModeIds::read); args.set(1, chunkOffests, AccessModeIds::read); - args.set(2, assignments, AccessModeIds::readwrite); - args.set(3, queue, AccessModeIds::readwrite); - args.set(4, queueEnd); - args.set(5, rowId); - args.set(6, clusterId); - args.set(7, chunkOffset); - args.set(8, chunkSize); - args.set(9, epsP); - args.set(10, nRows); + args.set(2, queueEnd); + args.set(3, rowId); + args.set(4, clusterId); + args.set(5, chunkOffset); + args.set(6, chunkSize); + args.set(7, epsP); + args.set(8, nRows); + args.set(9, assignments, AccessModeIds::readwrite); + args.set(10, queue, AccessModeIds::readwrite); KernelRange local_range(1, _maxWorkgroupSize); KernelRange global_range(numberOfChunks * _minSubgroupSize / _maxWorkgroupSize + 1, _maxWorkgroupSize);