Skip to content

Commit 8af88a4

Browse files
durga4githubdurga4github
authored
[MLIR][NVVM] Update PMEvent lowering to intrinsics (#171649)
The patch updates the lowering of `id` based pmevent also to intrinsics. The mask is simply (1 << event-id). Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
1 parent f8d1f53 commit 8af88a4

File tree

7 files changed

+73
-57
lines changed

7 files changed

+73
-57
lines changed

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

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -543,7 +543,7 @@ def NVVM_NanosleepOp : NVVM_Op<"nanosleep">,
543543
// NVVM Performance Monitor events
544544
//===----------------------------------------------------------------------===//
545545

546-
def NVVM_PMEventOp : NVVM_PTXBuilder_Op<"pmevent">,
546+
def NVVM_PMEventOp : NVVM_Op<"pmevent">,
547547
Arguments<(ins OptionalAttr<I16Attr>:$maskedEventId,
548548
OptionalAttr<I32Attr>:$eventId)> {
549549
let summary = "Trigger one or more Performance Monitor events.";
@@ -561,20 +561,20 @@ def NVVM_PMEventOp : NVVM_PTXBuilder_Op<"pmevent">,
561561
[For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent)
562562
}];
563563

564-
string llvmBuilder = [{
565-
llvm::Value *mId = builder.getInt16(* $maskedEventId);
566-
createIntrinsicCall(builder, llvm::Intrinsic::nvvm_pm_event_mask, {mId});
567-
}];
568-
569564
let assemblyFormat = "attr-dict (`id` `=` $eventId^)? (`mask` `=` $maskedEventId^)?";
565+
let hasVerifier = 1;
570566

571567
let extraClassDeclaration = [{
572-
bool hasIntrinsic() { return !getEventId(); }
568+
static mlir::NVVM::IDArgPair
569+
getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
570+
llvm::IRBuilderBase& builder);
573571
}];
574-
let extraClassDefinition = [{
575-
std::string $cppClass::getPtx() { return std::string("pmevent %0;"); }
572+
573+
string llvmBuilder = [{
574+
auto [id, args] = NVVM::PMEventOp::getIntrinsicIDAndArgs(
575+
*op, moduleTranslation, builder);
576+
createIntrinsicCall(builder, id, args);
576577
}];
577-
let hasVerifier = 1;
578578
}
579579

580580
//===----------------------------------------------------------------------===//

‎mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp‎

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2476,6 +2476,25 @@ mlir::NVVM::IDArgPair NVVM::BarrierOp::getIntrinsicIDAndArgs(
24762476
return {id, std::move(args)};
24772477
}
24782478

2479+
mlir::NVVM::IDArgPair
2480+
PMEventOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
2481+
llvm::IRBuilderBase &builder) {
2482+
auto thisOp = cast<NVVM::PMEventOp>(op);
2483+
llvm::Type *i16Ty = llvm::Type::getInt16Ty(mt.getLLVMContext());
2484+
2485+
// With event-id, mask is generated as (1 << event-id)
2486+
llvm::Value *maskVal;
2487+
if (auto eventAttr = thisOp.getEventIdAttr()) {
2488+
uint16_t mask = static_cast<uint16_t>(1u << eventAttr.getInt());
2489+
maskVal = llvm::ConstantInt::get(i16Ty, mask);
2490+
} else {
2491+
maskVal =
2492+
llvm::ConstantInt::get(i16Ty, thisOp.getMaskedEventIdAttr().getValue());
2493+
}
2494+
2495+
return {llvm::Intrinsic::nvvm_pm_event_mask, {maskVal}};
2496+
}
2497+
24792498
mlir::NVVM::IDArgPair MBarrierInitOp::getIntrinsicIDAndArgs(
24802499
Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
24812500
auto thisOp = cast<NVVM::MBarrierInitOp>(op);

‎mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir‎

Lines changed: 0 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -678,21 +678,6 @@ llvm.func @inline_ptx_multi_rw_r(%a : i32, %b : i32, %rw_c : f32, %rw_d : f32)
678678
llvm.return %r5 : f32
679679
}
680680

681-
682-
// -----
683-
684-
// CHECK-LABEL: @nvvm_pmevent
685-
llvm.func @nvvm_pmevent() {
686-
// CHECK: %[[S0:.+]] = llvm.mlir.constant(10 : i32) : i32
687-
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "pmevent $0;", "n" %[[S0]] : (i32) -> ()
688-
689-
nvvm.pmevent id = 10
690-
// CHECK: %[[S1:.+]] = llvm.mlir.constant(4 : i32) : i32
691-
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "pmevent $0;", "n" %[[S1]] : (i32) -> ()
692-
nvvm.pmevent id = 4
693-
llvm.return
694-
}
695-
696681
// -----
697682

698683
llvm.func @inline_ptx_pack_4i8(%src : vector<4xi8>, %mask : i32, %zero: i32) {

‎mlir/test/Target/LLVMIR/nvvm/pm_event.mlir‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
2+
3+
llvm.func @nvvm_pmevent_mask() {
4+
// CHECK-LABEL: define void @nvvm_pmevent_mask() {
5+
// CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 15000)
6+
// CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 4)
7+
// CHECK-NEXT: ret void
8+
// CHECK-NEXT: }
9+
nvvm.pmevent mask = 15000
10+
nvvm.pmevent mask = 4
11+
llvm.return
12+
}
13+
14+
llvm.func @nvvm_pmevent_id() {
15+
// CHECK-LABEL: define void @nvvm_pmevent_id() {
16+
// CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 1024)
17+
// CHECK-NEXT: call void @llvm.nvvm.pm.event.mask(i16 16)
18+
// CHECK-NEXT: ret void
19+
// CHECK-NEXT: }
20+
nvvm.pmevent id = 10
21+
nvvm.pmevent id = 4
22+
llvm.return
23+
}

‎mlir/test/Target/LLVMIR/nvvm/pm_event_invalid.mlir‎

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
2+
3+
llvm.func @pmevent_no_id() {
4+
// expected-error @below {{either `id` or `mask` must be set}}
5+
nvvm.pmevent
6+
}
7+
8+
// -----
9+
10+
llvm.func @pmevent_bigger15() {
11+
// expected-error @below {{`id` must be between 0 and 15}}
12+
nvvm.pmevent id = 16
13+
}
14+
15+
// -----
16+
17+
llvm.func @pmevent_many_ids() {
18+
// expected-error @below {{`id` and `mask` cannot be set at the same time}}
19+
nvvm.pmevent id = 1 mask = 1
20+
}
21+

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

Lines changed: 0 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1,26 +1,5 @@
11
// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s
22

3-
llvm.func @pmevent_no_id() {
4-
// expected-error @below {{either `id` or `mask` must be set}}
5-
nvvm.pmevent
6-
}
7-
8-
// -----
9-
10-
llvm.func @pmevent_bigger15() {
11-
// expected-error @below {{`id` must be between 0 and 15}}
12-
nvvm.pmevent id = 141
13-
}
14-
15-
// -----
16-
17-
llvm.func @pmevent_many_ids() {
18-
// expected-error @below {{`id` and `mask` cannot be set at the same time}}
19-
nvvm.pmevent id = 1 mask = 1
20-
}
21-
22-
// -----
23-
243
llvm.func @kernel_func(%numberOfThreads : i32) {
254
// expected-error @below {{'nvvm.barrier' op barrier id is missing, it should be set between 0 to 15}}
265
nvvm.barrier number_of_threads = %numberOfThreads

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

Lines changed: 0 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -903,17 +903,6 @@ llvm.func @nvvm_dot_accumulate_2way(%a: vector<2xi16>, %b: vector<4xi8>, %c: i32
903903

904904
// -----
905905

906-
// CHECK-LABEL: @nvvm_pmevent
907-
llvm.func @nvvm_pmevent() {
908-
// CHECK: call void @llvm.nvvm.pm.event.mask(i16 15000)
909-
nvvm.pmevent mask = 15000
910-
// CHECK: call void @llvm.nvvm.pm.event.mask(i16 4)
911-
nvvm.pmevent mask = 4
912-
llvm.return
913-
}
914-
915-
// -----
916-
917906
// CHECK-LABEL: @nanosleep
918907
llvm.func @nanosleep(%duration: i32) {
919908
// CHECK: call void @llvm.nvvm.nanosleep(i32 %{{.*}})

0 commit comments

Comments
 (0)