Skip to content

Commit

Permalink
Adressed comments uxlfoundation#7
Browse files Browse the repository at this point in the history
  • Loading branch information
amgrigoriev committed Mar 29, 2020
1 parent 0169cdb commit 65d11a2
Show file tree
Hide file tree
Showing 2 changed files with 15 additions and 15 deletions.
10 changes: 5 additions & 5 deletions algorithms/kernel/dbscan/oneapi/cl_kernels/dbscan_cl_kernels.cl
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -274,23 +274,23 @@ services::Status DBSCANBatchKernelUCAPI<algorithmFPType>::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));

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);
Expand Down

0 comments on commit 65d11a2

Please sign in to comment.