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

Jump to bottom

maximum number of threads per block for sm_86 is 1536 #45889

Closed
wants to merge 2 commits into from
Closed

maximum number of threads per block for sm_86 is 1536 #45889

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4129b89
__CUDA_ARCH__ for sm_86 is 1536
xwang233 Oct 6, 2020
a49811e
comments
xwang233 Oct 6, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
7 changes: 5 additions & 2 deletions c10/macros/Macros.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -193,11 +193,14 @@ namespace at { namespace cuda { using namespace c10::hip; }}
#define C10_DEVICE __device__
#define C10_HOST __host__
// constants from (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications)
// The maximum number of threads per multiprocessor is 1024 for Turing architecture (7.5)
// but 2048 for previous architectures. You'll get warnings if you exceed these constants.
// The maximum number of threads per multiprocessor is 1024 for Turing architecture (7.5),
// 1536 for Geforce Ampere (8.6),
// and 2048 for all other architectures. You'll get warnings if you exceed these constants.
// Hence, the following macros adjust the input values from the user to resolve potential warnings.
#if __CUDA_ARCH__ == 750
constexpr uint32_t CUDA_MAX_THREADS_PER_SM = 1024;
#elif __CUDA_ARCH__ == 860
constexpr uint32_t CUDA_MAX_THREADS_PER_SM = 1536;
#else
constexpr uint32_t CUDA_MAX_THREADS_PER_SM = 2048;
#endif
Expand Down