Skip to content

Add reverse-mode support for CUDA shared memory attributes - #1826

Open
Vedant2005goyal wants to merge 1 commit into
vgvassilev:masterfrom
Vedant2005goyal:add-shared-memory
Open

Add reverse-mode support for CUDA shared memory attributes#1826
Vedant2005goyal wants to merge 1 commit into
vgvassilev:masterfrom
Vedant2005goyal:add-shared-memory

Conversation

@Vedant2005goyal

Copy link
Copy Markdown
Contributor

This PR enables differentiation of CUDA kernels having __shared__ attribute. It also enable differentiation support of dynamic shared memory extern __shared__. __syncthreads and atomicAdd is also enabled during the reverse sweep of the reverse mode to ensure no race condition.

@github-actions

github-actions Bot commented Jun 1, 2026

Copy link
Copy Markdown
Contributor

clang-tidy review says "All clean, LGTM! 👍"

@codecov

codecov Bot commented Jun 1, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 97.72727% with 2 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
lib/Differentiator/VisitorBase.cpp 60.00% 2 Missing ⚠️

📢 Thoughts on this report? Let us know!

@github-actions

github-actions Bot commented Jun 1, 2026

Copy link
Copy Markdown
Contributor

clang-tidy review says "All clean, LGTM! 👍"

@github-actions

github-actions Bot commented Jun 1, 2026

Copy link
Copy Markdown
Contributor

clang-tidy review says "All clean, LGTM! 👍"

@Vedant2005goyal
Vedant2005goyal marked this pull request as draft June 3, 2026 14:04
@Vedant2005goyal Vedant2005goyal changed the title Add reverse-mode support for CUDA shared memory variables Add reverse-mode support for CUDA shared memory attributes Jun 3, 2026

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-tidy made some suggestions

if (VD->getType()->isIncompleteArrayType()) {
llvm::SmallVector<Expr*, 2> args = {
BuildDeclRef(VDDerived),
BuildDeclRef(m_SharedMemSizeParam)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: use of undeclared identifier 'm_SharedMemSizeParam' [clang-diagnostic-error]

                  BuildDeclRef(m_SharedMemSizeParam)
                               ^

BuildDeclRef(VDDerived),
BuildDeclRef(m_SharedMemSizeParam)
};
m_HasDynamicSharedMem = true;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: use of undeclared identifier 'm_HasDynamicSharedMem' [clang-diagnostic-error]

                m_HasDynamicSharedMem = true;
                ^


params.push_back(dPVD);
}
if (FD->hasAttr<CUDAGlobalAttr>() && m_HasDynamicSharedMem) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: use of undeclared identifier 'm_HasDynamicSharedMem' [clang-diagnostic-error]

    if (FD->hasAttr<CUDAGlobalAttr>() && m_HasDynamicSharedMem) {
                                         ^

II,
m_Context.getSizeType());
m_Sema.PushOnScopeChains( SharedBytesPVD, getCurrentScope(), /*AddToContext=*/false);
m_SharedMemSizeParam = SharedBytesPVD;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: use of undeclared identifier 'm_SharedMemSizeParam' [clang-diagnostic-error]

          m_SharedMemSizeParam = SharedBytesPVD;
          ^

@Vedant2005goyal
Vedant2005goyal marked this pull request as ready for review June 12, 2026 07:16

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-tidy made some suggestions

Comment thread lib/Differentiator/ReverseModeVisitor.cpp Outdated
Comment thread lib/Differentiator/ReverseModeVisitor.cpp Outdated
@github-actions

Copy link
Copy Markdown
Contributor

clang-tidy review says "All clean, LGTM! 👍"

1 similar comment
@github-actions

Copy link
Copy Markdown
Contributor

clang-tidy review says "All clean, LGTM! 👍"

if (!isDynamicSharedMem && VD->hasAttr<clang::CUDASharedAttr>()) {
auto* SharedAttr = clang::CUDASharedAttr::CreateImplicit(
m_Context, VD->getAttr<clang::CUDASharedAttr>()->getRange());
SharedAttr->setImplicit(false);

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why do we create implicit attribute and then change that -- can't we use the usual way of creating attributes?

@Vedant2005goyal Vedant2005goyal Jun 12, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the review. Older versions of Clang have a problem with the standard Create() method, so using an implicit attribute is a workaround for the problem

VDClone->setStorageClass(VD->getStorageClass() == clang::SC_Extern
? clang::SC_Extern
: clang::SC_Static);
}

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I believe that code should go in its own function.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-tidy made some suggestions

} else {
// Dynamic Shared Memory: Automate zero-init purely via AST
Expr* derivedRef = BuildDeclRef(VDDerived);
Expr* zeroIdx = ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: argument comment missing for literal argument 'val' [bugprone-argument-comment]

Suggested change
Expr* zeroIdx = ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);
Expr* zeroIdx = ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, /*val=*/0);

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-tidy made some suggestions

// Dynamic Shared Memory: Automate zero-init purely via AST
Expr* derivedRef = BuildDeclRef(VDDerived);
Expr* zeroIdx =
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: argument comment missing for literal argument 'val' [bugprone-argument-comment]

Suggested change
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, /*val=*/0);

initDiff.updateStmtDx(finalInit);
}

CloneCUDASharedAttr(VD, VDClone);

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That seems should be a part of BuildGlobalVarDecl.

Stmt* ReverseResult = utils::unwrapIfSingleStmt(RCS);
return {StmtDiff(ForwardResult, ReverseResult), EDiff};
}
void ReverseModeVisitor::CloneCUDASharedAttr(const clang::VarDecl* OriginalVD,

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This does not need to be a member function, a file static if at all, should be enough.

: clang::SC_Static);
}
}
void ReverseModeVisitor::HandleCUDASharedMemoryDecl(

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This should be inside a new file called ReverseModeVisitorCUDA.cpp

@Vedant2005goyal
Vedant2005goyal force-pushed the add-shared-memory branch 4 times, most recently from 429a5f0 to 303621d Compare June 16, 2026 17:39

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-tidy made some suggestions

if (!isDynamicSharedMem) {
CloneCUDASharedAttr(VD, VDDerived);
VDDerived->setStorageClass(clang::SC_Static);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: no header providing "clang::SC_Static" is directly included [misc-include-cleaner]

lib/Differentiator/ReverseModeVisitorCUDA.cpp:5:

+ #include <clang/Basic/Specifiers.h>

VDForward->setStorageClass(clang::SC_Static);

llvm::SmallVector<Expr*, 1> args = {BuildDeclRef(VDDerived)};
Stmt* initCall = GetCladZeroInit(args);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: no header providing "llvm::SmallVector" is directly included [misc-include-cleaner]

lib/Differentiator/ReverseModeVisitorCUDA.cpp:5:

+ #include <llvm/ADT/SmallVector.h>

Expr* derivedRef = BuildDeclRef(VDDerived);
Expr* zeroIdx =
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);
Expr* arraySub = BuildArraySubscript(derivedRef, {zeroIdx});

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: argument comment missing for literal argument 'val' [bugprone-argument-comment]

Suggested change
Expr* arraySub = BuildArraySubscript(derivedRef, {zeroIdx});
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, /*val=*/0);

Expr* arraySub = BuildArraySubscript(derivedRef, {zeroIdx});
QualType elemType = VDDerived->getType()->getPointeeType();
Expr* assignZero = BuildOp(BO_Assign, arraySub, getZeroInit(elemType));
memsetCalls.push_back(assignZero);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: no header providing "clang::BO_Assign" is directly included [misc-include-cleaner]

lib/Differentiator/ReverseModeVisitorCUDA.cpp:5:

+ #include <clang/AST/OperationKinds.h>

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

clang-tidy made some suggestions


if (!isDynamicSharedMem) {
CloneCUDASharedAttr(VD, VDDerived);
VDDerived->setStorageClass(clang::SC_Static);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: no header providing "clang::SC_Static" is directly included [misc-include-cleaner]

lib/Differentiator/ReverseModeVisitorCUDA.cpp:5:

+ #include <clang/Basic/Specifiers.h>

CloneCUDASharedAttr(VD, VDForward);
VDForward->setStorageClass(clang::SC_Static);

llvm::SmallVector<Expr*, 1> args = {BuildDeclRef(VDDerived)};

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: no header providing "llvm::SmallVector" is directly included [misc-include-cleaner]

lib/Differentiator/ReverseModeVisitorCUDA.cpp:5:

+ #include <llvm/ADT/SmallVector.h>

CloneCUDASharedAttr(VD, VDForward);
Expr* derivedRef = BuildDeclRef(VDDerived);
Expr* zeroIdx =
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: argument comment missing for literal argument 'val' [bugprone-argument-comment]

Suggested change
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);
ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, /*val=*/0);

ConstantFolder::synthesizeLiteral(m_Context.IntTy, m_Context, 0);
Expr* arraySub = BuildArraySubscript(derivedRef, {zeroIdx});
QualType elemType = VDDerived->getType()->getPointeeType();
Expr* assignZero = BuildOp(BO_Assign, arraySub, getZeroInit(elemType));

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

warning: no header providing "clang::BO_Assign" is directly included [misc-include-cleaner]

lib/Differentiator/ReverseModeVisitorCUDA.cpp:5:

+ #include <clang/AST/OperationKinds.h>

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants