From 9b2e5fa283e30f2cda7bd5d7e48764e0f598683e Mon Sep 17 00:00:00 2001 From: agrigori Date: Thu, 26 Mar 2020 19:22:45 +0300 Subject: [PATCH] Adressed comments #3 --- .../oneapi/cl_kernels/dbscan_cl_kernels.cl | 18 +++++++----------- .../dbscan_dense_default_batch_ucapi_impl.i | 2 +- 2 files changed, 8 insertions(+), 12 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 59513c0b62a..b2e6581ce07 100644 --- a/algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl +++ b/algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl @@ -29,27 +29,23 @@ DECLARE_SOURCE( dbscan_cl_kernels, - __kernel void point_distances(__global const algorithmFPType * points, int point_id, int power, int dim, + __kernel void compute_point_distances(__global const algorithmFPType * points, int point_id, int power, int num_features, int num_points, __global algorithmFPType * dist) { const int subgroup_index = get_global_id(0) * get_max_sub_group_size() + get_sub_group_id(); if (subgroup_index >= num_points) return; + const int subgroup_size = get_sub_group_size(); const int local_id = get_sub_group_local_id(); algorithmFPType sum = 0.0; - for (int i = local_id; i < dim; i += subgroup_size) + for (int i = local_id; i < num_features; i += subgroup_size) { - algorithmFPType val = fabs(points[point_id * dim + i] - points[subgroup_index * dim + i]); - algorithmFPType dm = 1.0; - for (int j = 0; j < power; j++) - { - dm *= val; - } - sum += dm; + algorithmFPType val = fabs(points[point_id * num_features + i] - points[subgroup_index * num_features + i]); + sum += pown(val, power); } - algorithmFPType ret = sub_group_reduce_add(sum); + algorithmFPType cur_nbr_distance = sub_group_reduce_add(sum); if (local_id == 0) { - dist[subgroup_index] = ret; + dist[subgroup_index] = cur_nbr_distance; } } 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 23c21173ff1..232effad6d4 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 @@ -379,7 +379,7 @@ services::Status DBSCANBatchKernelUCAPI::getPointDistances(cons auto & context = Environment::getInstance()->getDefaultExecutionContext(); auto & kernel_factory = context.getClKernelFactory(); DAAL_CHECK_STATUS_VAR(buildProgram(kernel_factory)); - auto kernel = kernel_factory.getKernel("point_distances", &st); + auto kernel = kernel_factory.getKernel("compute_point_distances", &st); DAAL_CHECK_STATUS_VAR(st); KernelArguments args(6);