Skip to content

Avoid use of class static variable in device function#7776

Merged
miscco merged 1 commit intoNVIDIA:mainfrom
miscco:avoid_device_items_per_tile
Feb 27, 2026
Merged

Avoid use of class static variable in device function#7776
miscco merged 1 commit intoNVIDIA:mainfrom
miscco:avoid_device_items_per_tile

Conversation

@miscco
Copy link
Copy Markdown
Contributor

@miscco miscco commented Feb 24, 2026

We should not use a class static member in a device function, because it is not defined there.

Luckily those tests do not need to be defined as __device__ anyway

Fixes nvbug5925530

@miscco miscco requested a review from a team as a code owner February 24, 2026 10:55
@miscco miscco requested a review from gevtushenko February 24, 2026 10:55
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Feb 24, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Feb 24, 2026
@github-actions

This comment has been minimized.

@miscco miscco force-pushed the avoid_device_items_per_tile branch from 1f7119e to 5a7984d Compare February 26, 2026 07:57
@bernhardmgruber bernhardmgruber changed the title Avoid use of class stati variable in device function Avoid use of class static variable in device function Feb 26, 2026
@github-actions

This comment has been minimized.

Comment on lines 197 to 199
thresholds = {ActivePolicyT::SmallReducePolicy::ITEMS_PER_TILE,
ActivePolicyT::MediumReducePolicy::ITEMS_PER_TILE,
ActivePolicyT::ReducePolicy::BLOCK_THREADS * ActivePolicyT::ReducePolicy::ITEMS_PER_THREAD};
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Reading through the nvbug, maybe it's enough to just avoid ODR use:

Suggested change
thresholds = {ActivePolicyT::SmallReducePolicy::ITEMS_PER_TILE,
ActivePolicyT::MediumReducePolicy::ITEMS_PER_TILE,
ActivePolicyT::ReducePolicy::BLOCK_THREADS * ActivePolicyT::ReducePolicy::ITEMS_PER_THREAD};
thresholds = {+ActivePolicyT::SmallReducePolicy::ITEMS_PER_TILE,
+ActivePolicyT::MediumReducePolicy::ITEMS_PER_TILE,
+ActivePolicyT::ReducePolicy::BLOCK_THREADS * +ActivePolicyT::ReducePolicy::ITEMS_PER_THREAD};

@miscco miscco force-pushed the avoid_device_items_per_tile branch from 5a7984d to 493d5a1 Compare February 26, 2026 16:51
We should not use a class static member in a device function, because it is not defined there.

Luckily those tests do not need to be defined as `__device__` anyway

Fixes nvbug5925530
@miscco miscco force-pushed the avoid_device_items_per_tile branch from 493d5a1 to df1e5a9 Compare February 26, 2026 16:59
@github-actions

This comment has been minimized.

@miscco miscco enabled auto-merge (squash) February 27, 2026 06:56
@github-actions
Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 14h 26m: Pass: 100%/213 | Total: 8d 11h | Max: 3h 56m | Hits: 71%/119936

See results here.

@miscco miscco merged commit b23597b into NVIDIA:main Feb 27, 2026
459 of 462 checks passed
@miscco miscco deleted the avoid_device_items_per_tile branch February 27, 2026 07:29
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Feb 27, 2026
@github-actions
Copy link
Copy Markdown
Contributor

Backport failed for branch/3.3.x, because it was unable to cherry-pick the commit(s).

Please cherry-pick the changes locally and resolve any conflicts.

git fetch origin branch/3.3.x
git worktree add -d .worktree/backport-7776-to-branch/3.3.x origin/branch/3.3.x
cd .worktree/backport-7776-to-branch/3.3.x
git switch --create backport-7776-to-branch/3.3.x
git cherry-pick -x b23597bc8d1565c81bf9f901af69b5b8816b463d

miscco added a commit to miscco/cccl that referenced this pull request Feb 27, 2026
We should not use a class static member in a device function, because it is not defined there.

Luckily those tests do not need to be defined as `__device__` anyway

Fixes nvbug5925530
wmaxey pushed a commit that referenced this pull request Feb 27, 2026
We should not use a class static member in a device function, because it is not defined there.

Luckily those tests do not need to be defined as `__device__` anyway

Fixes nvbug5925530
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

4 participants