Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix compile failure on RTX 4090 #2076

Merged
merged 6 commits into from
Jan 17, 2024
Merged

Fix compile failure on RTX 4090 #2076

Changes from 2 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
bb6aca8
Fix compile failure on RTX 4090 and update Copyright year. related is…
JieFengWang Jan 4, 2024
6cbc24e
Merge branch 'branch-24.02' into branch-24.02
wphicks Jan 4, 2024
661f637
Update cpp/include/raft/neighbors/detail/nn_descent.cuh
JieFengWang Jan 10, 2024
068f546
Merge branch 'branch-24.02' into branch-24.02
cjnolet Jan 16, 2024
56efa5d
Merge branch 'branch-24.02' into branch-24.02
cjnolet Jan 17, 2024
10aaa80
Merge branch 'branch-24.02' into branch-24.02
cjnolet Jan 17, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions cpp/include/raft/neighbors/detail/nn_descent.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2023, NVIDIA CORPORATION.
* Copyright (c) 2023-2024, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -686,12 +686,12 @@ __device__ __forceinline__ void remove_duplicates(
// Per
// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications,
// MAX_RESIDENT_THREAD_PER_SM = BLOCK_SIZE * BLOCKS_PER_SM = 2048
// For architectures 750 and 860, the values for MAX_RESIDENT_THREAD_PER_SM
// For architectures 750 and 860 (890), the values for MAX_RESIDENT_THREAD_PER_SM
// is 1024 and 1536 respectively, which means the bounds don't work anymore
template <typename Index_t, typename ID_t = InternalID_t<Index_t>>
RAFT_KERNEL
#ifdef __CUDA_ARCH__
#if (__CUDA_ARCH__) == 750 || (__CUDA_ARCH__) == 860
#if (__CUDA_ARCH__) == 750 || (__CUDA_ARCH__) == 860 || (__CUDA_ARCH__) == 890
__launch_bounds__(BLOCK_SIZE)
#else
__launch_bounds__(BLOCK_SIZE, 4)
Expand Down
Loading