Don't create CUDA atomics for basic indices - #1441
Conversation
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
Codecov ReportAll modified and coverable lines are covered by tests ✅
Additional details and impacted files@@ Coverage Diff @@
## master #1441 +/- ##
==========================================
+ Coverage 94.71% 94.78% +0.07%
==========================================
Files 56 56
Lines 9762 9874 +112
==========================================
+ Hits 9246 9359 +113
+ Misses 516 515 -1
🚀 New features to boost your workflow:
|
3de551c to
e50e165
Compare
kchristin22
left a comment
kchristin22
left a comment
There was a problem hiding this comment.
I think this approach is great, we can cover the const cases now and extend it to re-assignments later. This could be done by adding the decl refs of indexes that are injective to a list as we traverse them in the Inject checker and then update these decl refs when we ecnounter an assignment in the original visitor (reverse mode visitor). We don't have to store them in this PR, it could be left for the re-assignment one.
|
|
||
| bool isInjectiveIdx(const clang::Expr* E) { | ||
| // NOLINTNEXTLINE(cppcoreguidelines-pro-type-const-cast) | ||
| TraverseStmt(const_cast<clang::Expr*>(E)); |
There was a problem hiding this comment.
Can we avoid const_casts its UB.
| } else if (const auto* ASE = dyn_cast<ArraySubscriptExpr>(E)) { | ||
| const auto* base = | ||
| dyn_cast<DeclRefExpr>(ASE->getBase()->IgnoreImpCasts()); | ||
| if (const auto* PVD = dyn_cast<ParmVarDecl>(base->getDecl())) { |
There was a problem hiding this comment.
Checking if it's a parameter is common in both branches (DRE and ASE), so we should have that as the outer if-condition
|
|
||
| static bool isInjectiveE(const clang::Expr* E) { | ||
| class InjectiveCheckerExpr | ||
| : public clang::RecursiveASTVisitor<InjectiveCheckerExpr> { |
There was a problem hiding this comment.
warning: no header providing "clang::RecursiveASTVisitor" is directly included [misc-include-cleaner]
lib/Differentiator/ReverseModeVisitor.cpp:52:
- #include <cstddef>
+ #include <clang/AST/RecursiveASTVisitor.h>
+ #include <cstddef>| return true; | ||
| if (const auto* DRE = dyn_cast<DeclRefExpr>(E)) { | ||
| if (const auto* PVD = dyn_cast<ParmVarDecl>(DRE->getDecl())) { | ||
| if (m_DiffReq->hasAttr<clang::CUDAGlobalAttr>()) |
There was a problem hiding this comment.
warning: no header providing "clang::CUDAGlobalAttr" is directly included [misc-include-cleaner]
lib/Differentiator/ReverseModeVisitor.cpp:52:
- #include <cstddef>
+ #include <clang/AST/Attrs.inc>
+ #include <cstddef>
be9b58a to
72a6131
Compare
|
@kchristin22, is this good to go? |
| public: | ||
| InjectiveCheckerExpr() = default; | ||
|
|
||
| bool comparePatternToTree(const IdxNode* current, |
There was a problem hiding this comment.
We need some more documentation about what this routine does.
| const std::vector<std::string>& patternIdx, | ||
| size_t i = 1) { | ||
|
|
||
| if (!current && (i >= patternIdx.size() || patternIdx[i].empty())) |
There was a problem hiding this comment.
Under what conditions current can be a nullptr?
There was a problem hiding this comment.
Each node has a left and right children whatsoever, those might be not initialized.
There was a problem hiding this comment.
In that case, I'd move the
return false;
before that statement and remove the !current checks from the rest.
| return true; | ||
| } | ||
|
|
||
| bool TraverseDeclRefExpr(clang::DeclRefExpr* DRE) { |
There was a problem hiding this comment.
Why not VisitDeclRefExpr but TraverseDeclRefExpr?
There was a problem hiding this comment.
I think we have converged that traverse methods are better, since they do not visit subexpressions automatically, meaning if x=y is visited and not traversed, we would visit LHS, RHS and then x, y as subexpressions of x=y.
There was a problem hiding this comment.
Can we capture that as a comment in the appropriate place..
|
|
||
| [[nodiscard]] bool isExpr() const { return m_Kind == Kind::Expr; } | ||
| [[nodiscard]] bool isOpcode() const { return m_Kind == Kind::Opcode; } | ||
| }; |
There was a problem hiding this comment.
warning: enum 'side' uses a larger base type ('int', size: 4 bytes) than necessary for its value set, consider using 'std::uint8_t' (1 byte) as the base type to reduce its size [performance-enum-size]
enum class side { left, right } m_Side;
^| }; | ||
|
|
||
| ExprOrBinOp Node; | ||
| std::unique_ptr<IdxNode> left; |
There was a problem hiding this comment.
warning: constructor does not initialize these fields: m_Side [cppcoreguidelines-pro-type-member-init]
InjectiveCheckerExpr() = default;
^|
|
||
| if (m_Side == side::right) | ||
| m_ParentNode->right = std::move(curr); | ||
| } |
There was a problem hiding this comment.
warning: member 'm_Context' of type 'clang::ASTContext &' is a reference [cppcoreguidelines-avoid-const-or-ref-data-members]
clang::ASTContext& m_Context;
^| return false; | ||
|
|
||
| if (!isConstL && isConstR) | ||
| return TraverseStmt(L); |
There was a problem hiding this comment.
warning: method 'TraverseDeclRefExpr' can be made static [readability-convert-member-functions-to-static]
| return TraverseStmt(L); | |
| static bool TraverseDeclRefExpr(clang::DeclRefExpr* DRE) { |
|
|
||
| bool TraverseBinaryOperator(clang::BinaryOperator* BinOp) { | ||
| const auto opCode = BinOp->getOpcode(); | ||
| if (opCode == BO_Add || opCode == BO_Mul) { |
There was a problem hiding this comment.
warning: no header providing "clang::BO_Add" is directly included [misc-include-cleaner]
if (opCode == BO_Add || opCode == BO_Mul) {
^|
|
||
| bool TraverseBinaryOperator(clang::BinaryOperator* BinOp) { | ||
| const auto opCode = BinOp->getOpcode(); | ||
| if (opCode == BO_Add || opCode == BO_Mul) { |
There was a problem hiding this comment.
warning: no header providing "clang::BO_Mul" is directly included [misc-include-cleaner]
if (opCode == BO_Add || opCode == BO_Mul) {
^
| IdxNode* currPtr = nullptr; | ||
|
|
||
| if (!m_Root) { | ||
| m_Root = std::move(curr); |
There was a problem hiding this comment.
warning: no header providing "std::move" is directly included [misc-include-cleaner]
lib/Differentiator/CladUtils.cpp:21:
- #include <vector>
+ #include <utility>
+ #include <vector>| bool isInjective(const clang::Expr* E, clang::ASTContext& ctx) { | ||
| class InjectiveChecker | ||
| : public clang::RecursiveASTVisitor<InjectiveChecker> { | ||
| clang::ASTContext& m_Context; |
There was a problem hiding this comment.
warning: member 'm_Context' of type 'clang::ASTContext &' is a reference [cppcoreguidelines-avoid-const-or-ref-data-members]
clang::ASTContext& m_Context;
^| return false; | ||
| } | ||
|
|
||
| bool TraverseDeclRefExpr(clang::DeclRefExpr* DRE) { |
There was a problem hiding this comment.
warning: method 'TraverseDeclRefExpr' can be made static [readability-convert-member-functions-to-static]
| bool TraverseDeclRefExpr(clang::DeclRefExpr* DRE) { | |
| static bool TraverseDeclRefExpr(clang::DeclRefExpr* DRE) { |
| return false; | ||
| } | ||
|
|
||
| bool TraverseIntegerLiteral(clang::IntegerLiteral* IL) { return false; } |
There was a problem hiding this comment.
warning: method 'TraverseIntegerLiteral' can be made static [readability-convert-member-functions-to-static]
| bool TraverseIntegerLiteral(clang::IntegerLiteral* IL) { return false; } | |
| static bool TraverseIntegerLiteral(clang::IntegerLiteral* IL) { return false; } |
a716646 to
a827622
Compare
Before this PR, we unconditionally produced atomic operations for writes to global gpu memory. For example, the gradient of the kernel below:
```cpp
__global__ void indices_perm(int *out, int *in) {
int index = threadIdx.x + blockIdx.x * blockDim.x;
out[index] += in[index];
}
```
used to have an `atomicAdd` in the reverse pass corresponding to the `out[index] += in[index]` statement. The way `index` is defined makes it injective, meaning no two threads in any two blocks would share the same `index`. In cases like this no data races occur, so we can avoid using atomic operations.
3 participants
Before this PR, we unconditionally produced atomic operations for writes to global gpu memory. For example, the gradient of the kernel below:
used to have an
atomicAddin the reverse pass corresponding to theout[index] += in[index]statement. The wayindexis defined makes it injective, meaning no two threads in any two blocks would share the sameindex. In cases like this no data races occur, so we can avoid using atomic operations.