Skip to content

Commit ed3fe39

Browse files
AlexMacleanAlexMaclean
committed
updates from previous version
1 parent a46075f commit ed3fe39

File tree

11 files changed

+43
-41
lines changed

11 files changed

+43
-41
lines changed

‎clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1174,7 +1174,7 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
11741174
EmitScalarExpr(E->getArg(0)));
11751175
case NVPTX::BI__nvvm_barrier_sync_cnt:
11761176
return Builder.CreateCall(
1177-
CGM.getIntrinsic(Intrinsic::nvvm_barrier_cta_sync),
1177+
CGM.getIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count),
11781178
{EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1))});
11791179
default:
11801180
return nullptr;

‎clang/test/CodeGen/builtins-nvptx-ptx60.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,7 @@ __device__ void nvvm_sync(unsigned mask, int i, float f, int a, int b,
3535
// CHECK: call void @llvm.nvvm.barrier.cta.sync.all(i32
3636
// expected-error@+1 {{'__nvvm_barrier_sync' needs target feature ptx60}}
3737
__nvvm_barrier_sync(mask);
38-
// CHECK: call void @llvm.nvvm.barrier.cta.sync(i32
38+
// CHECK: call void @llvm.nvvm.barrier.cta.sync.count(i32
3939
// expected-error@+1 {{'__nvvm_barrier_sync_cnt' needs target feature ptx60}}
4040
__nvvm_barrier_sync_cnt(mask, i);
4141

‎llvm/docs/NVPTXUsage.rst

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -207,13 +207,13 @@ Syntax:
207207

208208
.. code-block:: llvm
209209
210-
declare void @llvm.nvvm.barrier.cta.sync(i32 %id, i32 %n)
210+
declare void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 %n)
211211
declare void @llvm.nvvm.barrier.cta.sync.all(i32 %id)
212-
declare void @llvm.nvvm.barrier.cta.arrive(i32 %id, i32 %n)
212+
declare void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 %n)
213213
214-
declare void @llvm.nvvm.barrier.cta.sync.aligned(i32 %id, i32 %n)
214+
declare void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 %n)
215215
declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %id)
216-
declare void @llvm.nvvm.barrier.cta.arrive.aligned(i32 %id, i32 %n)
216+
declare void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 %n)
217217
218218
Overview:
219219
"""""""""
@@ -230,7 +230,8 @@ Operand %id specifies a logical barrier resource and must fall within the range
230230
participating in the barrier. When specifying a thread count, the value must be
231231
a multiple of the warp size. With the '``@llvm.nvvm.barrier.cta.sync.*``'
232232
variants, the '``.all``' suffix indicates that all threads in the CTA should
233-
participate in the barrier and the %n operand is not present.
233+
participate in the barrier while the '``.count``' suffix indicates that only
234+
the threads specified by the %n operand should participate in the barrier.
234235

235236
All forms of the '``@llvm.nvvm.barrier.cta.*``' intrinsic cause the executing
236237
thread to wait for all non-exited threads from its warp and then marks the

‎llvm/include/llvm/IR/IntrinsicsNVVM.td

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1300,9 +1300,9 @@ let TargetPrefix = "nvvm" in {
13001300
foreach align = ["", "_aligned"] in {
13011301
def int_nvvm_barrier_cta_sync # align # _all :
13021302
Intrinsic<[], [llvm_i32_ty]>;
1303-
def int_nvvm_barrier_cta_sync # align :
1303+
def int_nvvm_barrier_cta_sync # align # _count :
13041304
Intrinsic<[], [llvm_i32_ty, llvm_i32_ty]>;
1305-
def int_nvvm_barrier_cta_arrive # align :
1305+
def int_nvvm_barrier_cta_arrive # align # _count :
13061306
Intrinsic<[], [llvm_i32_ty, llvm_i32_ty]>;
13071307
}
13081308
}

‎llvm/lib/IR/AutoUpgrade.cpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2492,13 +2492,14 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
24922492
Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned_all,
24932493
{}, {Arg});
24942494
} else if (Name == "barrier") {
2495-
Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_aligned, {},
2496-
{CI->getArgOperand(0), CI->getArgOperand(1)});
2495+
Rep = Builder.CreateIntrinsic(
2496+
Intrinsic::nvvm_barrier_cta_sync_aligned_count, {},
2497+
{CI->getArgOperand(0), CI->getArgOperand(1)});
24972498
} else if (Name == "barrier.sync") {
24982499
Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_all, {},
24992500
{CI->getArgOperand(0)});
25002501
} else if (Name == "barrier.sync.cnt") {
2501-
Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync, {},
2502+
Rep = Builder.CreateIntrinsic(Intrinsic::nvvm_barrier_cta_sync_count, {},
25022503
{CI->getArgOperand(0), CI->getArgOperand(1)});
25032504
} else {
25042505
Intrinsic::ID IID = shouldUpgradeNVPTXBF16Intrinsic(Name);

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -132,12 +132,12 @@ multiclass BARRIER2<string asmstr, Intrinsic intrinsic, list<Predicate> requires
132132
// "barrier.*.aligned" variants. We use the older syntax for compatibility with
133133
// older versions of the PTX ISA.
134134
defm BARRIER_CTA_SYNC_ALIGNED_ALL : BARRIER1<"bar.sync", int_nvvm_barrier_cta_sync_aligned_all>;
135-
defm BARRIER_CTA_SYNC_ALIGNED : BARRIER2<"bar.sync", int_nvvm_barrier_cta_sync_aligned>;
136-
defm BARRIER_CTA_ARRIVE_ALIGNED : BARRIER2<"bar.arrive", int_nvvm_barrier_cta_arrive_aligned>;
135+
defm BARRIER_CTA_SYNC_ALIGNED : BARRIER2<"bar.sync", int_nvvm_barrier_cta_sync_aligned_count>;
136+
defm BARRIER_CTA_ARRIVE_ALIGNED : BARRIER2<"bar.arrive", int_nvvm_barrier_cta_arrive_aligned_count>;
137137

138138
defm BARRIER_CTA_SYNC_ALL : BARRIER1<"barrier.sync", int_nvvm_barrier_cta_sync_all, [hasPTX<60>]>;
139-
defm BARRIER_CTA_SYNC : BARRIER2<"barrier.sync", int_nvvm_barrier_cta_sync, [hasPTX<60>]>;
140-
defm BARRIER_CTA_ARRIVE : BARRIER2<"barrier.arrive", int_nvvm_barrier_cta_arrive, [hasPTX<60>]>;
139+
defm BARRIER_CTA_SYNC : BARRIER2<"barrier.sync", int_nvvm_barrier_cta_sync_count, [hasPTX<60>]>;
140+
defm BARRIER_CTA_ARRIVE : BARRIER2<"barrier.arrive", int_nvvm_barrier_cta_arrive_count, [hasPTX<60>]>;
141141

142142
class INT_BARRIER_CLUSTER<string variant, Intrinsic Intr,
143143
list<Predicate> Preds = [hasPTX<78>, hasSM<90>]>:

‎llvm/lib/Transforms/IPO/AttributorAttributes.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2151,7 +2151,7 @@ struct AANoUnwindCallSite final
21512151
bool AANoSync::isAlignedBarrier(const CallBase &CB, bool ExecutedAligned) {
21522152
switch (CB.getIntrinsicID()) {
21532153
case Intrinsic::nvvm_barrier_cta_sync_aligned_all:
2154-
case Intrinsic::nvvm_barrier_cta_sync_aligned:
2154+
case Intrinsic::nvvm_barrier_cta_sync_aligned_count:
21552155
case Intrinsic::nvvm_barrier0_and:
21562156
case Intrinsic::nvvm_barrier0_or:
21572157
case Intrinsic::nvvm_barrier0_popc:

‎llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -335,9 +335,9 @@ define void @cta_barriers(i32 %x, i32 %y) {
335335
; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)
336336
; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %x)
337337
; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %x)
338-
; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned(i32 %x, i32 %y)
338+
; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %x, i32 %y)
339339
; CHECK: call void @llvm.nvvm.barrier.cta.sync.all(i32 %x)
340-
; CHECK: call void @llvm.nvvm.barrier.cta.sync(i32 %x, i32 %y)
340+
; CHECK: call void @llvm.nvvm.barrier.cta.sync.count(i32 %x, i32 %y)
341341
call void @llvm.nvvm.barrier0()
342342
call void @llvm.nvvm.barrier.n(i32 %x)
343343
call void @llvm.nvvm.bar.sync(i32 %x)

‎llvm/test/CodeGen/NVPTX/barrier.ll

Lines changed: 20 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -4,11 +4,11 @@
44

55
declare void @llvm.nvvm.bar.warp.sync(i32)
66
declare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32)
7-
declare void @llvm.nvvm.barrier.cta.sync.aligned(i32, i32)
7+
declare void @llvm.nvvm.barrier.cta.sync.aligned.count(i32, i32)
88
declare void @llvm.nvvm.barrier.cta.sync.all(i32)
9-
declare void @llvm.nvvm.barrier.cta.sync(i32, i32)
10-
declare void @llvm.nvvm.barrier.cta.arrive(i32, i32)
11-
declare void @llvm.nvvm.barrier.cta.arrive.aligned(i32, i32)
9+
declare void @llvm.nvvm.barrier.cta.sync.count(i32, i32)
10+
declare void @llvm.nvvm.barrier.cta.arrive.count(i32, i32)
11+
declare void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32, i32)
1212

1313
define void @barrier_warp_sync(i32 %id) {
1414
; CHECK-LABEL: barrier_warp_sync(
@@ -53,10 +53,10 @@ define void @barrier_cta_sync_aligned(i32 %id, i32 %cnt) {
5353
; CHECK-NEXT: bar.sync %r1, 64;
5454
; CHECK-NEXT: bar.sync 4, 64;
5555
; CHECK-NEXT: ret;
56-
call void @llvm.nvvm.barrier.cta.sync.aligned(i32 %id, i32 %cnt)
57-
call void @llvm.nvvm.barrier.cta.sync.aligned(i32 3, i32 %cnt)
58-
call void @llvm.nvvm.barrier.cta.sync.aligned(i32 %id, i32 64)
59-
call void @llvm.nvvm.barrier.cta.sync.aligned(i32 4, i32 64)
56+
call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 %cnt)
57+
call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 3, i32 %cnt)
58+
call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %id, i32 64)
59+
call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 4, i32 64)
6060
ret void
6161
}
6262

@@ -73,10 +73,10 @@ define void @barrier_cta_arrive_aligned(i32 %id, i32 %cnt) {
7373
; CHECK-NEXT: bar.arrive %r1, 64;
7474
; CHECK-NEXT: bar.arrive 4, 64;
7575
; CHECK-NEXT: ret;
76-
call void @llvm.nvvm.barrier.cta.arrive.aligned(i32 %id, i32 %cnt)
77-
call void @llvm.nvvm.barrier.cta.arrive.aligned(i32 3, i32 %cnt)
78-
call void @llvm.nvvm.barrier.cta.arrive.aligned(i32 %id, i32 64)
79-
call void @llvm.nvvm.barrier.cta.arrive.aligned(i32 4, i32 64)
76+
call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 %cnt)
77+
call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 3, i32 %cnt)
78+
call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 %id, i32 64)
79+
call void @llvm.nvvm.barrier.cta.arrive.aligned.count(i32 4, i32 64)
8080
ret void
8181
}
8282

@@ -108,10 +108,10 @@ define void @barrier_cta_sync(i32 %id, i32 %cnt) {
108108
; CHECK-NEXT: barrier.sync %r1, 64;
109109
; CHECK-NEXT: barrier.sync 4, 64;
110110
; CHECK-NEXT: ret;
111-
call void @llvm.nvvm.barrier.cta.sync(i32 %id, i32 %cnt)
112-
call void @llvm.nvvm.barrier.cta.sync(i32 3, i32 %cnt)
113-
call void @llvm.nvvm.barrier.cta.sync(i32 %id, i32 64)
114-
call void @llvm.nvvm.barrier.cta.sync(i32 4, i32 64)
111+
call void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 %cnt)
112+
call void @llvm.nvvm.barrier.cta.sync.count(i32 3, i32 %cnt)
113+
call void @llvm.nvvm.barrier.cta.sync.count(i32 %id, i32 64)
114+
call void @llvm.nvvm.barrier.cta.sync.count(i32 4, i32 64)
115115
ret void
116116
}
117117

@@ -128,9 +128,9 @@ define void @barrier_cta_arrive(i32 %id, i32 %cnt) {
128128
; CHECK-NEXT: barrier.arrive %r1, 64;
129129
; CHECK-NEXT: barrier.arrive 4, 64;
130130
; CHECK-NEXT: ret;
131-
call void @llvm.nvvm.barrier.cta.arrive(i32 %id, i32 %cnt)
132-
call void @llvm.nvvm.barrier.cta.arrive(i32 3, i32 %cnt)
133-
call void @llvm.nvvm.barrier.cta.arrive(i32 %id, i32 64)
134-
call void @llvm.nvvm.barrier.cta.arrive(i32 4, i32 64)
131+
call void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 %cnt)
132+
call void @llvm.nvvm.barrier.cta.arrive.count(i32 3, i32 %cnt)
133+
call void @llvm.nvvm.barrier.cta.arrive.count(i32 %id, i32 64)
134+
call void @llvm.nvvm.barrier.cta.arrive.count(i32 4, i32 64)
135135
ret void
136136
}

‎mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -552,7 +552,7 @@ def NVVM_BarrierOp : NVVM_Op<"barrier", [AttrSizedOperandSegments]> {
552552
llvm::Value *id = $barrierId ? $barrierId : builder.getInt32(0);
553553
if ($numberOfThreads)
554554
createIntrinsicCall(
555-
builder, llvm::Intrinsic::nvvm_barrier_cta_sync_aligned,
555+
builder, llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_count,
556556
{id, $numberOfThreads});
557557
else
558558
createIntrinsicCall(

‎mlir/test/Target/LLVMIR/nvvmir.mlir

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -174,7 +174,7 @@ llvm.func @llvm_nvvm_barrier(%barID : i32, %numberOfThreads : i32) {
174174
nvvm.barrier
175175
// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %[[barId]])
176176
nvvm.barrier id = %barID
177-
// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned(i32 %[[barId]], i32 %[[numThreads]])
177+
// CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %[[barId]], i32 %[[numThreads]])
178178
nvvm.barrier id = %barID number_of_threads = %numberOfThreads
179179
llvm.return
180180
}

0 commit comments

Comments
 (0)