Home
last modified time | relevance | path

Searched refs:mmaShape (Results 1 – 6 of 6) sorted by relevance

/llvm-project-15.0.7/mlir/test/Dialect/NVGPU/
H A Dinvalid.mlir53 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x4xf16>, vector<2x2x…
60 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x4x…
67 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2x…
74 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<2x4xf16>, vector<2x2x…
81 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<4x2xf32>, vector<2x1xf…
88 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<1x4xf32>, vector<2x1xf…
95 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [8, 8, 4]} : (vector<1x2xf64>, vector<1x1xf6…
102 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [8, 8, 4]} : (vector<1x1xf64>, vector<1x1xf6…
109 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<4x4xi…
116 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xi32>, vector<2x4x…
H A Droundtrip.mlir16 // CHECK: nvgpu.mma.sync(%{{.*}}, %{{.*}}, %{{.*}}) {mmaShape = [16, 8, 16]} : (vector<4x2xf1…
17 %d = nvgpu.mma.sync(%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} :
/llvm-project-15.0.7/mlir/lib/Dialect/LLVMIR/IR/
H A DNVVMDialect.cpp357 std::array<int64_t, 3> mmaShape{getShapeAttr().getM(), getShapeAttr().getN(), in verify() local
372 if (mmaShape[0] == 16) { in verify()
414 int64_t unitA = (mmaShape[0] / 8) * (mmaShape[2] / kFactor); in verify()
415 int64_t unitB = (mmaShape[1] / 8) * (mmaShape[2] / kFactor); in verify()
423 if (mmaShape[0] == 8) { in verify()
463 [&](const auto &allowed) { return allowed == mmaShape; })) { in verify()
465 llvm::interleaveComma(mmaShape, errorStream); in verify()
/llvm-project-15.0.7/mlir/test/Conversion/NVGPUToNVVM/
H A Dnvgpu-to-nvvm.mlir16 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2x…
35 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2x…
65 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<2x2xf16>, vector<1x2xf…
99 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<2x4xi…
120 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<2x8xi4>, vector<1x8xi…
147 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 64]} : (vector<4x8xi4>, vector<2x8xi…
160 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [8, 8, 4]} : (vector<1x1xf64>, vector<1x1xf6…
222 …%d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 4]} : (vector<2x1xf32>, vector<1x1xf…
/llvm-project-15.0.7/mlir/test/Conversion/VectorToGPU/
H A Dvector-to-mma-ops-mma-sync.mlir90 …// CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<2x…
155 …// CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [8, 8, 4]} : (vector<1x1xf64>, vector<1x1…
333 // CHECK-SAME: mmaShape = [16, 8, 4]
393 // CHECK-SAME: mmaShape = [16, 8, 8]
459 …// CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 64]} : (vector<4x8xi4>, vector<2x…
527 …// CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<2x…
594 // CHECK-SAME: mmaShape = [16, 8, 8]
/llvm-project-15.0.7/mlir/include/mlir/Dialect/NVGPU/IR/
H A DNVGPU.td97 `mmaShape` attribute in the form of a list of dimensions `[m, n, k]`.
114 AnyVector:$matrixC, I64ArrayAttr:$mmaShape);