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

Set __launch_bounds__ in kernel whenever we are able #3794

Open
wants to merge 8 commits into
base: main
Choose a base branch
from
Open

Set __launch_bounds__ in kernel whenever we are able #3794

Changes from 1 commit
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
b98f522
Set __launch_bounds__ in kernel whenever we are able
jacobhinkle Jan 29, 2025
eb80538
Compute num threads per block properly
jacobhinkle Jan 30, 2025
b78ae42
Fix error caused by using evaluate() without checking if const
jacobhinkle Jan 31, 2025
254a8a9
Fix failing tests with expected cuda code
jacobhinkle Jan 31, 2025
4495654
Just use a const_cast
jacobhinkle Jan 31, 2025
508d7ae
Disable clang-tidy warning on const_cast
jacobhinkle Jan 31, 2025
abebb2a
Merge branch 'main' into jh/set_launch_bounds_more_often
jacobhinkle Feb 4, 2025
1e6e6b9
Merge branch 'main' into jh/set_launch_bounds_more_often
jacobhinkle Feb 4, 2025
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: 4 additions & 2 deletions csrc/codegen.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -279,13 +279,15 @@ class CudaKernelGenerator : private kir::ConstIrVisitor {
const std::string& kernel_name,
std::optional<int64_t> num_threads_per_cta) {
code_ << "__global__ void ";
if (num_threads_per_cta.has_value()) {
code_ << "__launch_bounds__(/*MAX_THREADS_PER_BLOCK=*/"
<< num_threads_per_cta.value() << ") ";
}
if (kernel_->hasManaged("enable_register_sharing") &&
kernel_->getManaged<bool>("enable_register_sharing")) {
NVF_ERROR(
num_threads_per_cta.has_value(),
"__launch_bounds__ must be set for register sharing warp specialization");
code_ << "__launch_bounds__(/*MAX_THREADS_PER_BLOCK=*/"
<< num_threads_per_cta.value() << ") ";
}
if (kernel_->hasManaged("cluster_dims")) {
auto cluster_dims =
Expand Down