Skip to content
This repository was archived by the owner on Jul 28, 2025. It is now read-only.

Conversation

@isVoid
Copy link
Contributor

@isVoid isVoid commented May 19, 2023

Description

Place holder for PRs to close during burndown period.

Closes #1061

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

@github-actions github-actions bot added the libcuspatial Relates to the cuSpatial C++ library label May 19, 2023
auto aggregate = BlockReduce(temp_storage).Reduce(partial, cub::Min());

// atmomic with leading thread
if (cooperative_groups::this_thread_block().thread_rank() == 0)
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there any advantage to this over if (threadIdx.x == 0)? What code does this produce? It's impossible for threadIdx.x == 0 to have exited before any other threads in its block in this code.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Cannot think of any advantage comparing to threadIdx.x == 0. I originally had hopped writing modularized kernel with cg but but later regressed. This is probably a remnant from prototyping.

std::size_t constexpr threads_per_block = 256;
std::size_t const num_blocks =
(multilinestrings1.num_points() + threads_per_block - 1) / threads_per_block;
std::size_t num_threads = 1e8;
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rather than hard coding, this may be a job for the occupancy API...
https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__OCCUPANCY.html

@isVoid isVoid changed the base branch from branch-23.06 to branch-23.08 June 2, 2023 18:56
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.

Labels

libcuspatial Relates to the cuSpatial C++ library

Projects

Status: Todo

Development

Successfully merging this pull request may close these issues.

Make ST_distance core kernels (point-point, point-linestring, linestring-linestring) perfectly load balanced.

2 participants