This is an archive of the discontinued LLVM Phabricator instance.

Differential D17056

Mark all CUDA device-side function defs and decls as convergent.
ClosedPublic

Authored by jlebar on Feb 9 2016, 4:19 PM.

Details

Reviewers
majnemer
rnk
Commits
rGddd97faeecce: [CUDA] Mark all CUDA device-side function defs, decls, and calls as convergent.
rC261779: [CUDA] Mark all CUDA device-side function defs, decls, and calls as convergent.
rL261779: [CUDA] Mark all CUDA device-side function defs, decls, and calls as convergent.
Summary

This is important for e.g. the following case:

void sync() { __syncthreads(); }
void foo() {
  do_something();
  sync();
  do_something_else():
}

Without this change, if the optimizer does not inline sync() (which it
won't because syncthreads is also marked as noduplicate, for now
anyway), it is free to perform optimizations on sync() that it would not
be able to perform on syncthreads(), because sync() is not marked as
convergent.

This chagne is conservative; the optimizer will remove these attrs where
it can, see r260318, r260319.

Diff Detail

Repository
rL LLVM

Event Timeline

jlebar updated this revision to Diff 47395.Feb 9 2016, 4:19 PM
jlebar retitled this revision from to Mark all CUDA device-side function defs and decls as convergent..
jlebar updated this object.
jlebar added a reviewer: majnemer.
jlebar added subscribers: tra, echristo, jhen, cfe-commits.
tra added inline comments.Feb 9 2016, 4:43 PM
lib/CodeGen/CodeGenModule.cpp
1880 ↗(On Diff #47395)

intrinsically?

test/CodeGenCUDA/convergent.cu
4 ↗(On Diff #47395)

You may want to add -disable-llvm-passes to make sure you're checking results of front-end actions only.

jlebar updated this revision to Diff 47409.Feb 9 2016, 5:08 PM
jlebar marked 2 inline comments as done.

Update per tra's review.

jlebar added a reviewer: rnk.Feb 16 2016, 4:44 PM
majnemer edited edge metadata.Feb 17 2016, 11:02 AM
majnemer added a subscriber: majnemer.

SetLLVMFunctionAttributesForDefinition seems like the more obvious place
for this logic.

jlebar updated this revision to Diff 48255.Feb 17 2016, 4:08 PM
jlebar edited edge metadata.

Move code into SetLLVMFunctionAttributesForDefinition.

jlebar added a comment.Feb 17 2016, 4:32 PM
In D17056#355198, @jlebar wrote:

Move code into SetLLVMFunctionAttributesForDefinition.

Actually, this doesn't work -- we don't annotate

__host__ __device__ void baz();

as convergent. (I ran the tests, but of course I didn't notice it failing...)

jlebar updated this revision to Diff 48260.Feb 17 2016, 4:35 PM

Move coded into SetLLVMFunctionAttributes (not ForDefinition).

jlebar added a comment.Feb 17 2016, 4:35 PM
In D17056#355228, @jlebar wrote:

Move coded into SetLLVMFunctionAttributes (not ForDefinition).

Much better. :)

majnemer added inline comments.Feb 17 2016, 4:42 PM
lib/CodeGen/CodeGenModule.cpp
817–823 ↗(On Diff #48260)

Sorry for the goose chase...

ConstructAttributeList seems better.

jlebar updated this revision to Diff 48261.Feb 17 2016, 4:53 PM

Move code into ConstructAttributeList. Now it applies to both functions and calls.

lib/CodeGen/CodeGenModule.cpp
817–823 ↗(On Diff #48260)

Actually, this is much better; now I don't need D17313.

jlebar added a comment.Feb 24 2016, 1:34 PM

Friendly ping -- are we happy with this?

majnemer accepted this revision.Feb 24 2016, 1:54 PM
majnemer edited edge metadata.

LGTM

This revision is now accepted and ready to land.Feb 24 2016, 1:54 PM
Closed by commit rL261779: [CUDA] Mark all CUDA device-side function defs, decls, and calls as convergent. (authored by jlebar). · Explain WhyFeb 24 2016, 1:59 PM
This revision was automatically updated to reflect the committed changes.
jlebar mentioned this in D17313: [CUDA] Annotate all calls in CUDA device mode as convergent..Feb 24 2016, 2:03 PM

Revision Contents

PathSize
cfe/
trunk/
lib/
CodeGen/
8 lines
test/
CodeGenCUDA/
39 lines
2 lines

Diff 48979

cfe/trunk/lib/CodeGen/CGCall.cpp

Show First 20 Lines • Show All 1,589 Lines • ▼ Show 20 Lines if (FD && FD->hasAttr<TargetAttr>()) {
std::sort(Features.begin(), Features.end()); std::sort(Features.begin(), Features.end());
FuncAttrs.addAttribute( FuncAttrs.addAttribute(
"target-features", "target-features",
llvm::join(Features.begin(), Features.end(), ",")); llvm::join(Features.begin(), Features.end(), ","));
} }
} }
} }
if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
// Conservatively, mark all functions and calls in CUDA as convergent
// (meaning, they may call an intrinsically convergent op, such as
// __syncthreads(), and so can't have certain optimizations applied around
// them). LLVM will remove this attribute where it safely can.
FuncAttrs.addAttribute(llvm::Attribute::Convergent);
}
ClangToLLVMArgMapping IRFunctionArgs(getContext(), FI); ClangToLLVMArgMapping IRFunctionArgs(getContext(), FI);
QualType RetTy = FI.getReturnType(); QualType RetTy = FI.getReturnType();
const ABIArgInfo &RetAI = FI.getReturnInfo(); const ABIArgInfo &RetAI = FI.getReturnInfo();
switch (RetAI.getKind()) { switch (RetAI.getKind()) {
case ABIArgInfo::Extend: case ABIArgInfo::Extend:
if (RetTy->hasSignedIntegerRepresentation()) if (RetTy->hasSignedIntegerRepresentation())
RetAttrs.addAttribute(llvm::Attribute::SExt); RetAttrs.addAttribute(llvm::Attribute::SExt);
▲ Show 20 Lines • Show All 2,109 Lines • Show Last 20 Lines

cfe/trunk/test/CodeGenCUDA/convergent.cu

// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
// RUN: -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
// RUN: -disable-llvm-passes -o - %s | \
// RUN: FileCheck -check-prefix HOST %s
#include "Inputs/cuda.h"
// DEVICE: Function Attrs:
// DEVICE-SAME: convergent
// DEVICE-NEXT: define void @_Z3foov
__device__ void foo() {}
// HOST: Function Attrs:
// HOST-NOT: convergent
// HOST-NEXT: define void @_Z3barv
// DEVICE: Function Attrs:
// DEVICE-SAME: convergent
// DEVICE-NEXT: define void @_Z3barv
__host__ __device__ void baz();
__host__ __device__ void bar() {
// DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]]
baz();
}
// DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
// DEVICE: attributes [[BAZ_ATTR]] = {
// DEVICE-SAME: convergent
// DEVICE-SAME: }
// DEVICE: attributes [[CALL_ATTR]] = { convergent }
// HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
// HOST: attributes [[BAZ_ATTR]] = {
// HOST-NOT: convergent
// NOST-SAME: }

cfe/trunk/test/CodeGenCUDA/device-var-init.cu

Show First 20 Lines • Show All 376 Lines • ▼ Show 20 Lines
// CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec) // CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec)
// CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv) // CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
// CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf) // CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
// CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs) // CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
// CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc) // CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
// CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc) // CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
// CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec) // CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
// CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1) // CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
// CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t) #3 // CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
// CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec) // CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
// CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec) // CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
// CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec) // CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
// CHECK: call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*)) // CHECK: call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*))
// CHECK: ret void // CHECK: ret void
// We should not emit global init function. // We should not emit global init function.
// CHECK-NOT: @_GLOBAL__sub_I // CHECK-NOT: @_GLOBAL__sub_I