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

Refactor to use FastDivmod for predicated strided dgrad iterators. #1453

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

Refactor to use FastDivmod for predicated strided dgrad iterators. #1453

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
91fce1f
refactor to use fast div mod for predicated tiler iterator strided dg…
ZelboK Apr 3, 2024
be88f24
clean up
ZelboK Apr 3, 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
1 change: 1 addition & 0 deletions .gitignore
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,4 @@
build/
# PyCache files
__pycache__/
cutlass_library.egg-info/
43 changes: 32 additions & 11 deletions include/cutlass/epilogue/threadblock/predicated_tile_iterator_strided_dgrad.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,6 +116,9 @@ class PredicatedTileIteratorStridedDgrad {
/// Convolution problem size
cutlass::conv::Conv2dProblemSize problem_size;
int tiled_rows_per_filter;

FastDivmod pq_divmod;
FastDivmod q_divmod;

CUTLASS_HOST_DEVICE
Params() { }
Expand Down Expand Up @@ -234,7 +237,8 @@ class PredicatedTileIteratorStridedDgrad {
):
params_(params)
{



TensorCoord thread_offset = ThreadMap::initial_offset(thread_idx) + threadblock_offset;

int r = start_r;
Expand All @@ -254,6 +258,8 @@ class PredicatedTileIteratorStridedDgrad {

p_ = (params_.problem_size.H - start_h_ + params_.problem_size.stride_h - 1) / params_.problem_size.stride_h;
q_ = (params_.problem_size.W - start_w_ + params_.problem_size.stride_w - 1) / params_.problem_size.stride_w;
params_.pq_divmod = FastDivmod(p_*q_);
params_.q_divmod = FastDivmod(q_);

extent_row_ = extent.row();
thread_start_row_ = thread_offset.row();
Expand Down Expand Up @@ -312,11 +318,19 @@ class PredicatedTileIteratorStridedDgrad {
int npq_offset = (row_offset + thread_start_row_) % params_.tiled_rows_per_filter;

// (STEP 4.a) [order NHW rows to be loaded and stored in output Dx NHWxC layout]
int n = npq_offset / (p_ * q_);
int residual = npq_offset % (p_ * q_);
int p = residual / q_;
int q = residual % q_;


// The subsequent fast_divmod() operations are equivalent to the following logical computation:
// int nzpq = npq_offset;
// int n = nzpq / (p_ * q_);
// int residual = nzpq % (p_ * q_);
// int p = residual1 / q_;
// int q = residual1 % q_;

int p, q, residual, n;

params_.pq_divmod(n, residual, npq_offset);
params_.q_divmod(p, q, residual);

int mapped_row_offset = n * (params_.problem_size.H * params_.problem_size.W) +
(start_h_ + p * params_.problem_size.stride_h) * params_.problem_size.W +
(start_w_ + q * params_.problem_size.stride_w);
Expand Down Expand Up @@ -379,11 +393,18 @@ class PredicatedTileIteratorStridedDgrad {
int npq_offset = (row_offset + thread_start_row_) % params_.tiled_rows_per_filter;

// (STEP 4.a) [order NHW rows to be loaded and stored in output Dx NHWxC layout]
int n = npq_offset / (p_ * q_);
int residual = npq_offset % (p_ * q_);
int p = residual / q_;
int q = residual % q_;


// The subsequent fast_divmod() operations are equivalent to the following logical computation:

// int n = npq_offset / (p_ * q_);
// int residual = npq_offset % (p_ * q_);
// int p = residual / q_;
// int q = residual % q_;

int n, residual, p, q;
params_.pq_divmod(n, residual, npq_offset);
params_.q_divmod(p, q, residual);

int mapped_row_offset = n * (params_.problem_size.H * params_.problem_size.W) +
(start_h_ + p * params_.problem_size.stride_h) * params_.problem_size.W +
(start_w_ + q * params_.problem_size.stride_w);
Expand Down