Skip to content

[SYCL][CUDA][libclc] Add atomic loads and stores with various memory orders and scopes #5191

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

Merged
merged 5 commits into from
Jan 17, 2022
Merged

[SYCL][CUDA][libclc] Add atomic loads and stores with various memory orders and scopes #5191

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
db2baab
[SYCL][CUDA][libclc] Added atomic loads and stores with various memor…
t4c1 Dec 21, 2021
98aaa34
Merge branch 'sycl' into ptx_atomic_ldst
t4c1 Dec 21, 2021
be3b553
format and push/pop macros
t4c1 Dec 23, 2021
e400d17
removed redundant include and added tests for volatile intrinsics
t4c1 Jan 13, 2022
5c91b3d
removed tmp variable
t4c1 Jan 14, 2022
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
86 changes: 86 additions & 0 deletions clang/include/clang/Basic/BuiltinsNVPTX.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2037,6 +2037,92 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70)
TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70)
TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70)

#pragma push_macro("LD_VOLATILE_BUILTIN_TYPES")
#define LD_VOLATILE_BUILTIN_TYPES(ADDR_SPACE) \
BUILTIN(__nvvm_volatile_ld##ADDR_SPACE##_i, "iiD*", "n") \
BUILTIN(__nvvm_volatile_ld##ADDR_SPACE##_l, "LiLiD*", "n") \
BUILTIN(__nvvm_volatile_ld##ADDR_SPACE##_ll, "LLiLLiD*", "n") \
BUILTIN(__nvvm_volatile_ld##ADDR_SPACE##_f, "ffD*", "n") \
BUILTIN(__nvvm_volatile_ld##ADDR_SPACE##_d, "ddD*", "n")

#pragma push_macro("LD_BUILTIN_TYPES")
#define LD_BUILTIN_TYPES(ORDER, SCOPE, ADDR_SPACE) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_i, "iiD*", "n", SM_70) \
Copy link
Contributor

Choose a reason for hiding this comment

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

I think it would make things a bit more orderly if builtins would be using __nvvm_ld/__nvvm_st prefix and put modifiers as suffixes in the same order they are used in the PTX instructions the builtins translate into.

E.g. __nvvm_ld_shared_acquire_gpu_i() -> ld.shared.acquire.gpu.u32

Copy link
Contributor Author

@t4c1 t4c1 Dec 23, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

That would be nice to have, but I would argue that being consistent with other builtins is more important. I was looking especially at atomic read-modify-write operations I worked on before, where we have for example __nvvm_atom_acq_rel_sys_min_shared_i.

Copy link
Contributor

Choose a reason for hiding this comment

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

OK. We can live with the order of modifiers matching that of __nvvm_atom.

That said, __nvvm_atom does have the instruction mnemonic atom in front and so should these ld builtins.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

My understanding is that these have atom at the start because they also have atom in the PTX, which is not true for ld and st instructions.

TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_l, "LiLiD*", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_ll, "LLiLLiD*", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_f, "ffD*", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_d, "ddD*", "n", SM_70)

#pragma push_macro("LD_BUILTIN_AS_TYPES")
#define LD_BUILTIN_AS_TYPES(ORDER, SCOPE) \
LD_BUILTIN_TYPES(ORDER, SCOPE, _gen) \
LD_BUILTIN_TYPES(ORDER, SCOPE, _global) \
LD_BUILTIN_TYPES(ORDER, SCOPE, _shared)

#pragma push_macro("LD_BUILTIN_SCOPES_AS_TYPES")
#define LD_BUILTIN_SCOPES_AS_TYPES(ORDER) \
LD_BUILTIN_AS_TYPES(ORDER, ) \
LD_BUILTIN_AS_TYPES(ORDER, _cta) \
LD_BUILTIN_AS_TYPES(ORDER, _sys)

LD_BUILTIN_SCOPES_AS_TYPES()
LD_BUILTIN_SCOPES_AS_TYPES(_acquire)
LD_VOLATILE_BUILTIN_TYPES(_gen)
LD_VOLATILE_BUILTIN_TYPES(_global)
LD_VOLATILE_BUILTIN_TYPES(_shared)

#undef LD_VOLATILE_BUILTIN_TYPES
#pragma pop_macro("LD_VOLATILE_BUILTIN_TYPES")
#undef LD_BUILTIN_TYPES
#pragma pop_macro("LD_BUILTIN_TYPES")
#undef LD_BUILTIN_AS_TYPES
#pragma pop_macro("LD_BUILTIN_AS_TYPES")
#undef LD_BUILTIN_SCOPES_AS_TYPES
#pragma pop_macro("LD_BUILTIN_SCOPES_AS_TYPES")

#pragma push_macro("ST_VOLATILE_BUILTIN_TYPES")
#define ST_VOLATILE_BUILTIN_TYPES(ADDR_SPACE) \
BUILTIN(__nvvm_volatile_st##ADDR_SPACE##_i, "viD*i", "n") \
BUILTIN(__nvvm_volatile_st##ADDR_SPACE##_l, "vLiD*Li", "n") \
BUILTIN(__nvvm_volatile_st##ADDR_SPACE##_ll, "vLLiD*LLi", "n") \
BUILTIN(__nvvm_volatile_st##ADDR_SPACE##_f, "vfD*f", "n") \
BUILTIN(__nvvm_volatile_st##ADDR_SPACE##_d, "vdD*d", "n")

#pragma push_macro("ST_BUILTIN_TYPES")
#define ST_BUILTIN_TYPES(ORDER, SCOPE, ADDR_SPACE) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_i, "viD*i", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_l, "vLiD*Li", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_ll, "vLLiD*LLi", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_f, "vfD*f", "n", SM_70) \
TARGET_BUILTIN(__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_d, "vdD*d", "n", SM_70)

#pragma push_macro("ST_BUILTIN_AS_TYPES")
#define ST_BUILTIN_AS_TYPES(ORDER, SCOPE) \
ST_BUILTIN_TYPES(ORDER, SCOPE, _gen) \
ST_BUILTIN_TYPES(ORDER, SCOPE, _global) \
ST_BUILTIN_TYPES(ORDER, SCOPE, _shared)

#pragma push_macro("ST_BUILTIN_SCOPES_AS_TYPES")
#define ST_BUILTIN_SCOPES_AS_TYPES(ORDER) \
ST_BUILTIN_AS_TYPES(ORDER, ) \
ST_BUILTIN_AS_TYPES(ORDER, _cta) \
ST_BUILTIN_AS_TYPES(ORDER, _sys)

ST_BUILTIN_SCOPES_AS_TYPES()
ST_BUILTIN_SCOPES_AS_TYPES(_release)
ST_VOLATILE_BUILTIN_TYPES(_gen)
ST_VOLATILE_BUILTIN_TYPES(_global)
ST_VOLATILE_BUILTIN_TYPES(_shared)

#undef ST_VOLATILE_BUILTIN_TYPES
#pragma pop_macro("ST_VOLATILE_BUILTIN_TYPES")
#undef ST_BUILTIN_TYPES
#pragma pop_macro("ST_BUILTIN_TYPES")
#undef ST_BUILTIN_AS_TYPES
#pragma pop_macro("ST_BUILTIN_AS_TYPES")
#undef ST_BUILTIN_SCOPES_AS_TYPES
#pragma pop_macro("ST_BUILTIN_SCOPES_AS_TYPES")

// Compiler Error Warn
BUILTIN(__nvvm_compiler_error, "vcC*4", "n")
BUILTIN(__nvvm_compiler_warn, "vcC*4", "n")
Expand Down
97 changes: 97 additions & 0 deletions clang/lib/CodeGen/CGBuiltin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,6 +54,8 @@
#include "llvm/Support/X86TargetParser.h"
#include <sstream>

#include <iostream>

using namespace clang;
using namespace CodeGen;
using namespace llvm;
Expand Down Expand Up @@ -17429,6 +17431,22 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) {
Ptr->getType()}),
{Ptr, ConstantInt::get(Builder.getInt32Ty(), Align.getQuantity())});
};
auto MakeScopedLd = [&](unsigned IntrinsicID) {
Value *Ptr = EmitScalarExpr(E->getArg(0));
auto tmp = Builder.CreateCall(
CGM.getIntrinsic(IntrinsicID, {Ptr->getType()->getPointerElementType(),
Ptr->getType()}),
{Ptr});
return tmp;
};
auto MakeScopedSt = [&](unsigned IntrinsicID) {
Value *Ptr = EmitScalarExpr(E->getArg(0));
return Builder.CreateCall(
CGM.getIntrinsic(
IntrinsicID,
{Ptr->getType(), Ptr->getType()->getPointerElementType()}),
{Ptr, EmitScalarExpr(E->getArg(1))});
};
Copy link
Contributor

Choose a reason for hiding this comment

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

The pointer is in generic AS. We may want to addrspacecast it into specific AS when we generate the calls to the AS-specific builtins. That would allow compiler to eliminate unnecessary ASC to generic AS it uses when an address of a variable is taken.

auto MakeScopedAtomic = [&](unsigned IntrinsicID) {
Value *Ptr = EmitScalarExpr(E->getArg(0));
return Builder.CreateCall(
Expand All @@ -17444,6 +17462,85 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) {
{Ptr, EmitScalarExpr(E->getArg(1)), EmitScalarExpr(E->getArg(2))});
};
switch (BuiltinID) {

#define LD_VOLATILE_CASES(ADDR_SPACE) \
case NVPTX::BI__nvvm_volatile_ld##ADDR_SPACE##_i: \
case NVPTX::BI__nvvm_volatile_ld##ADDR_SPACE##_l: \
case NVPTX::BI__nvvm_volatile_ld##ADDR_SPACE##_ll: \
return MakeScopedLd(Intrinsic::nvvm_ld##ADDR_SPACE##_i_volatile); \
case NVPTX::BI__nvvm_volatile_ld##ADDR_SPACE##_f: \
case NVPTX::BI__nvvm_volatile_ld##ADDR_SPACE##_d: \
return MakeScopedLd(Intrinsic::nvvm_ld##ADDR_SPACE##_f_volatile);

#define LD_CASES(ORDER, SCOPE, ADDR_SPACE) \
case NVPTX::BI__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_i: \
case NVPTX::BI__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_l: \
case NVPTX::BI__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_ll: \
return MakeScopedLd(Intrinsic::nvvm_ld##ADDR_SPACE##_i##ORDER##SCOPE); \
case NVPTX::BI__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_f: \
case NVPTX::BI__nvvm##ORDER##SCOPE##_ld##ADDR_SPACE##_d: \
return MakeScopedLd(Intrinsic::nvvm_ld##ADDR_SPACE##_f##ORDER##SCOPE);

#define LD_CASES_AS(ORDER, SCOPE) \
LD_CASES(ORDER, SCOPE, _gen) \
LD_CASES(ORDER, SCOPE, _global) \
LD_CASES(ORDER, SCOPE, _shared)

#define LD_CASES_AS_SCOPES(ORDER) \
LD_CASES_AS(ORDER, ) \
LD_CASES_AS(ORDER, _cta) \
LD_CASES_AS(ORDER, _sys)

LD_CASES_AS_SCOPES()
LD_CASES_AS_SCOPES(_acquire)
LD_VOLATILE_CASES(_gen)
LD_VOLATILE_CASES(_global)
LD_VOLATILE_CASES(_shared)

#undef LD_VOLATILE_CASES
#undef LD_CASES
#undef LD_CASES_AS
#undef LD_CASES_AS_SCOPES

#define ST_VOLATILE_CASES(ADDR_SPACE) \
case NVPTX::BI__nvvm_volatile_st##ADDR_SPACE##_i: \
case NVPTX::BI__nvvm_volatile_st##ADDR_SPACE##_l: \
case NVPTX::BI__nvvm_volatile_st##ADDR_SPACE##_ll: \
return MakeScopedSt(Intrinsic::nvvm_st##ADDR_SPACE##_i_volatile); \
case NVPTX::BI__nvvm_volatile_st##ADDR_SPACE##_f: \
case NVPTX::BI__nvvm_volatile_st##ADDR_SPACE##_d: \
return MakeScopedSt(Intrinsic::nvvm_st##ADDR_SPACE##_f_volatile);

#define ST_CASES(ORDER, SCOPE, ADDR_SPACE) \
case NVPTX::BI__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_i: \
case NVPTX::BI__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_l: \
case NVPTX::BI__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_ll: \
return MakeScopedSt(Intrinsic::nvvm_st##ADDR_SPACE##_i##ORDER##SCOPE); \
case NVPTX::BI__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_f: \
case NVPTX::BI__nvvm##ORDER##SCOPE##_st##ADDR_SPACE##_d: \
return MakeScopedSt(Intrinsic::nvvm_st##ADDR_SPACE##_f##ORDER##SCOPE);

#define ST_CASES_AS(ORDER, SCOPE) \
ST_CASES(ORDER, SCOPE, _gen) \
ST_CASES(ORDER, SCOPE, _global) \
ST_CASES(ORDER, SCOPE, _shared)

#define ST_CASES_AS_SCOPES(ORDER) \
ST_CASES_AS(ORDER, ) \
ST_CASES_AS(ORDER, _cta) \
ST_CASES_AS(ORDER, _sys)

ST_CASES_AS_SCOPES()
ST_CASES_AS_SCOPES(_release)
ST_VOLATILE_CASES(_gen)
ST_VOLATILE_CASES(_global)
ST_VOLATILE_CASES(_shared)

#undef ST_VOLATILE_CASES
#undef ST_CASES
#undef ST_CASES_AS
#undef ST_CASES_AS_SCOPES

case NVPTX::BI__nvvm_atom_add_gen_i:
case NVPTX::BI__nvvm_atom_add_gen_l:
case NVPTX::BI__nvvm_atom_add_gen_ll:
Expand Down
Loading