This is an archive of the discontinued LLVM Phabricator instance.

Differential D121666

[mlir][nvvm] Fix bug in ldmatrix intrinsic conversion
ClosedPublic

Authored by ThomasRaoux on Mar 14 2022, 9:57 PM.

Details

Reviewers
ftynse
mravishankar
christopherbate
Commits
rG6d007e0278ad: [mlir][nvvm] Fix bug in ldmatrix intrinsic conversion
Summary

The ldmatrix intrinsic trans option was inverted. Bug reported by @ctb on discord.

Diff Detail

Event Timeline

ThomasRaoux created this revision.Mar 14 2022, 9:57 PM
Herald added a reviewer: ftynse. · View Herald TranscriptMar 14 2022, 9:57 PM
Herald added a reviewer: mravishankar. · View Herald Transcript
Herald added a project: Restricted Project. · View Herald Transcript
Herald added subscribers: awarzynski, sdasgup3, wenzhicui and 21 others. · View Herald Transcript
ThomasRaoux requested review of this revision.Mar 14 2022, 9:57 PM
Herald added a project: Restricted Project. · View Herald TranscriptMar 14 2022, 9:57 PM
Herald added subscribers: stephenneuendorffer, nicolasvasilache. · View Herald Transcript
ThomasRaoux edited the summary of this revision. (Show Details)Mar 14 2022, 10:00 PM
christopherbate accepted this revision.Mar 14 2022, 10:02 PM
This revision is now accepted and ready to land.Mar 14 2022, 10:02 PM
This revision was landed with ongoing or failed builds.Mar 14 2022, 10:04 PM
Closed by commit rG6d007e0278ad: [mlir][nvvm] Fix bug in ldmatrix intrinsic conversion (authored by ThomasRaoux). · Explain Why
This revision was automatically updated to reflect the committed changes.
ThomasRaoux added a commit: rG6d007e0278ad: [mlir][nvvm] Fix bug in ldmatrix intrinsic conversion.
Harbormaster completed remote builds in B154249: Diff 415317.Mar 14 2022, 10:26 PM

Revision Contents

PathSize
mlir/
lib/
Target/
LLVMIR/
Dialect/
NVVM/
2 lines
test/
Target/
LLVMIR/
12 lines

Diff 415317

mlir/lib/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.cpp

Show First 20 Lines • Show All 61 Lines • ▼ Show 20 Lines if (withPredicate) {
} }
} }
llvm_unreachable("unknown shuffle kind"); llvm_unreachable("unknown shuffle kind");
} }
/// Return the intrinsic ID associated with ldmatrix for the given paramters. /// Return the intrinsic ID associated with ldmatrix for the given paramters.
static llvm::Intrinsic::ID getLdMatrixIntrinsicId(NVVM::MMALayout layout, static llvm::Intrinsic::ID getLdMatrixIntrinsicId(NVVM::MMALayout layout,
int32_t num) { int32_t num) {
if (layout == NVVM::MMALayout::col) { if (layout == NVVM::MMALayout::row) {
switch (num) { switch (num) {
case 1: case 1:
return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16; return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16;
case 2: case 2:
return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16; return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16;
case 4: case 4:
return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16; return llvm::Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16;
default: default:
▲ Show 20 Lines • Show All 74 Lines • Show Last 20 Lines

mlir/test/Target/LLVMIR/nvvmir.mlir

Show First 20 Lines • Show All 172 Lines • ▼ Show 20 Lines// CHECK: call void @llvm.nvvm.cp.async.commit.group()
nvvm.cp.async.commit.group nvvm.cp.async.commit.group
// CHECK: call void @llvm.nvvm.cp.async.wait.group(i32 0) // CHECK: call void @llvm.nvvm.cp.async.wait.group(i32 0)
nvvm.cp.async.wait.group 0 nvvm.cp.async.wait.group 0
llvm.return llvm.return
} }
// CHECK-LABEL: @ld_matrix( // CHECK-LABEL: @ld_matrix(
llvm.func @ld_matrix(%arg0: !llvm.ptr<i32, 3>) { llvm.func @ld_matrix(%arg0: !llvm.ptr<i32, 3>) {
// CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16.p3i32(i32 addrspace(3)* %{{.*}}) // CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.b16.p3i32(i32 addrspace(3)* %{{.*}})
%l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32 %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
// CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16.p3i32(i32 addrspace(3)* %{{.*}}) // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.b16.p3i32(i32 addrspace(3)* %{{.*}})
%l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)> %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
// CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3i32(i32 addrspace(3)* %{{.*}}) // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3i32(i32 addrspace(3)* %{{.*}})
%l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)> %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
// CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16.p3i32(i32 addrspace(3)* %{{.*}})
%l1t = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<col>} : (!llvm.ptr<i32, 3>) -> i32
// CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16.p3i32(i32 addrspace(3)* %{{.*}})
%l2t = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<col>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
// CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3i32(i32 addrspace(3)* %{{.*}})
%l4t = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<col>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
llvm.return llvm.return
} }
// This function has the "kernel" attribute attached and should appear in the // This function has the "kernel" attribute attached and should appear in the
// NVVM annotations after conversion. // NVVM annotations after conversion.
llvm.func @kernel_func() attributes {nvvm.kernel} { llvm.func @kernel_func() attributes {nvvm.kernel} {
llvm.return llvm.return
} }
// CHECK: !nvvm.annotations = // CHECK: !nvvm.annotations =
// CHECK-NOT: {i32 ()* @nvvm_special_regs, !"kernel", i32 1} // CHECK-NOT: {i32 ()* @nvvm_special_regs, !"kernel", i32 1}
// CHECK: {void ()* @kernel_func, !"kernel", i32 1} // CHECK: {void ()* @kernel_func, !"kernel", i32 1}