Skip to content

Fix build break #72

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
manbearian merged 6 commits into from
Dec 1, 2023
Merged

Fix build break #72

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
a274484
replace memremf.tensor_store with bufferization.materialize_in_destin…
manbearian Dec 1, 2023
e4d40c6
mark output memref as 'writeable'
manbearian Dec 1, 2023
ec51be0
update LKG Triton as well
manbearian Dec 1, 2023
5da65ee
try recommended hash
manbearian Dec 1, 2023
9c84a93
Merge branch 'main' into dev/ianb/excise-memref-tensor-store-op
nhat-nguyen Dec 1, 2023
755fff3
empty commit
manbearian Dec 1, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions .github/workflows/integration-tests.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,6 @@ jobs:
call-workflow:
uses: ./.github/workflows/test-plugin.yml
with:
# [BACKEND] Fix layouts in nvgpu RewriteTensorPointer. (#2634)
triton-ref: '0e3bf3f58061bd62725e7c502e7e0af19690df1c'
# [BACKEND] Refactor wgmma descriptor creation
triton-ref: '56c284cf7e39f249cdf1d8d5dba7892deb0286d6'
triton-shared-ref: ${{ github.ref }}
2 changes: 1 addition & 1 deletion README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -102,7 +102,7 @@ Important details to note:
%reinterpret_cast = memref.reinterpret_cast %arg2 to offset: [...] memref<*xf32> to memref<1024xf32>
%extracted_slice = tensor.extract_slice %15[0] [%21] [1] : tensor<1024xf32> to tensor<?xf32>
%subview = memref.subview %reinterpret_cast[0] [%21] [1] : memref<1024xf32> to memref<?xf32>
memref.tensor_store %extracted_slice, %subview : memref<?xf32>
bufferization.materialize_in_destination %extracted_slice in writable %subview
```

+ element-wise `arith` and `math` operators are converted to their corresponding `linalg.generic` version.
Expand Down
8 changes: 6 additions & 2 deletions lib/Conversion/TritonToLinalg/TritonToLinalg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -744,7 +744,9 @@ struct StoreConverter : public OpConversionPattern<triton::StoreOp> {

// 1. Simple case where no mask is used.
if (!mask) {
rewriter.create<memref::TensorStoreOp>(loc, val, ptr);
auto storeOp = rewriter.create<bufferization::MaterializeInDestinationOp>(
loc, val, ptr);
storeOp.setWritable(true);
rewriter.eraseOp(op);
return success();
}
Expand All @@ -761,7 +763,9 @@ struct StoreConverter : public OpConversionPattern<triton::StoreOp> {
auto srcSlice = mstate.getExtractSlice(val, loc, rewriter);
auto dstSubview = mstate.getSubview(ptr, loc, rewriter);

rewriter.create<memref::TensorStoreOp>(loc, srcSlice, dstSubview);
auto storeOp = rewriter.create<bufferization::MaterializeInDestinationOp>(
loc, srcSlice, dstSubview);
storeOp.setWritable(true);
rewriter.eraseOp(op);

return success();
Expand Down
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_2d_example.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,6 +64,6 @@ module {
// CHECK: } -> tensor<4x256xbf16>
// CHECK: %[[VAL_21:.*]] = arith.index_cast %[[VAL_3]] : i32 to index
// CHECK: %[[VAL_22:.*]] = memref.reinterpret_cast %[[VAL_2]] to offset: {{\[}}%[[VAL_21]]], sizes: [4, 256], strides: [1, %[[VAL_7]]] : memref<*xbf16> to memref<4x256xbf16, strided<[1, ?], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_23:.*]], %[[VAL_22]] : memref<4x256xbf16, strided<[1, ?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_23:.*]] in writable %[[VAL_22]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_add_value.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,6 @@ module {
// CHECK: %[[VAL_19:.*]] = memref.alloc() : memref<4x256xbf16>
// CHECK: memref.copy %[[VAL_13]], %[[VAL_19]] : memref<4x256xbf16, strided<[1, ?], offset: ?>> to memref<4x256xbf16>
// CHECK: %[[VAL_20:.*]] = bufferization.to_tensor %[[VAL_19]] restrict writable : memref<4x256xbf16>
// CHECK: memref.tensor_store %[[VAL_20]], %[[VAL_18]] : memref<4x256xbf16, strided<[1, ?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_20]] in writable %[[VAL_18]]
// CHECK: return
// CHECK: }
4 changes: 2 additions & 2 deletions test/Conversion/TritonToLinalg/addptr_dim1.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -81,7 +81,7 @@ module {
// CHECK-DAG: [[VAR_2_:%.+]] = bufferization.to_tensor [[RES_]] restrict writable : memref<1x256xbf16>
// CHECK-DAG: [[VAR_3_:%.+]] = arith.index_cast [[PARAM_1_]] : i32 to index
// CHECK: [[VAR_reinterpret_cast_0_:%.+]] = memref.reinterpret_cast [[PARAM_0_]] to offset: {{.}}[[VAR_3_]]{{.}}, sizes: [1, 256], strides: [256, 1] : memref<*xbf16> to memref<1x256xbf16, strided<[256, 1], offset: ?>>
// CHECK: memref.tensor_store [[VAR_2_]], [[VAR_reinterpret_cast_0_]] : memref<1x256xbf16, strided<[256, 1], offset: ?>>
// CHECK: bufferization.materialize_in_destination [[VAR_2_]] in writable [[VAR_reinterpret_cast_0_]]
// CHECK-DAG: [[VAR_4_:%.+]]:3 = scf.for [[VAR_arg5_:%.+]] = [[CST_0_]] to [[CST_12_]] step [[CST_3_]] iter_args([[VAR_arg6_:%.+]] = [[VAR_1_]], [[VAR_arg7_:%.+]] = [[CST_0_]], [[VAR_arg8_:%.+]] = [[CST_0_]]) -> (tensor<4x256xbf16>, index, index) {
// CHECK-DAG: [[VAR_5_:%.+]] = arith.index_cast [[VAR_arg5_]] : index to i32
// CHECK: [[VAR_6_:%.+]] = arith.muli [[VAR_5_]], [[CST_256_1_]] : i32
Expand All @@ -102,6 +102,6 @@ module {
// CHECK: scf.yield [[VAR_10_]], [[VAR_12_]], [[CST_0_]] : tensor<4x256xbf16>, index, index
// CHECK: }
// CHECK: [[VAR_reinterpret_cast_1_:%.+]] = memref.reinterpret_cast [[PARAM_0_]] to offset: [0], sizes: [4, 256], strides: {{.}}[[CST_256_]], 1] : memref<*xbf16> to memref<4x256xbf16, strided<[?, 1]>>
// CHECK: memref.tensor_store [[VAR_4_]]#0, [[VAR_reinterpret_cast_1_]] : memref<4x256xbf16, strided<[?, 1]>>
// CHECK: bufferization.materialize_in_destination [[VAR_4_]]#0 in writable [[VAR_reinterpret_cast_1_]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_for_accumulation.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -88,6 +88,6 @@ module {
// CHECK: }
// CHECK: %[[VAL_36:.*]] = arith.index_cast %[[VAL_3]] : i32 to index
// CHECK: %[[VAL_37:.*]] = memref.reinterpret_cast %[[VAL_2]] to offset: {{\[}}%[[VAL_36]]], sizes: [4, 256], strides: [1, %[[VAL_8]]] : memref<*xbf16> to memref<4x256xbf16, strided<[1, ?], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_38:.*]]#0, %[[VAL_37]] : memref<4x256xbf16, strided<[1, ?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_38:.*]]#0 in writable %[[VAL_37]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_for_expand_ptr.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,7 @@ module {
// CHECK: %[[VAL_15:.*]] = memref.alloc() : memref<256x256xbf16>
// CHECK: memref.copy %[[VAL_14]], %[[VAL_15]] : memref<256x256xbf16, strided<[?, 1], offset: ?>> to memref<256x256xbf16>
// CHECK: %[[VAL_16:.*]] = bufferization.to_tensor %[[VAL_15]] restrict writable : memref<256x256xbf16>
// CHECK: memref.tensor_store %[[VAL_16]], %[[VAL_14]] : memref<256x256xbf16, strided<[?, 1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_16]] in writable %[[VAL_14]]
// CHECK: %[[VAL_17:.*]] = arith.addi %[[VAL_12]], %[[VAL_9]] : index
// CHECK: scf.yield %[[VAL_17]] : index
// CHECK: }
Expand Down
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_for_more_init_args.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -55,7 +55,7 @@ module {
// CHECK: %[[VAL_23:.*]] = memref.alloc() : memref<256xbf16>
// CHECK: memref.copy %[[VAL_17]], %[[VAL_23]] : memref<256xbf16, strided<[?], offset: ?>> to memref<256xbf16>
// CHECK: %[[VAL_24:.*]] = bufferization.to_tensor %[[VAL_23]] restrict writable : memref<256xbf16>
// CHECK: memref.tensor_store %[[VAL_24]], %[[VAL_19]] : memref<256xbf16, strided<[?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_24]] in writable %[[VAL_19]]
// CHECK: %[[VAL_25:.*]] = arith.addi %[[VAL_21]], %[[VAL_10]] : index
// CHECK: %[[VAL_26:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_25]]], sizes: [256], strides: {{\[}}%[[VAL_8]]] : memref<*xbf16> to memref<256xbf16, strided<[?], offset: ?>>
// CHECK: %[[VAL_27:.*]] = arith.addi %[[VAL_16]], %[[VAL_10]] : index
Expand Down
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_for_used_after_update.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,7 +91,7 @@ module {
// CHECK: %[[VAL_14:.*]] = memref.alloc() : memref<256xbf16>
// CHECK: memref.copy %[[VAL_13]], %[[VAL_14]] : memref<256xbf16, strided<[?], offset: ?>> to memref<256xbf16>
// CHECK: %[[VAL_15:.*]] = bufferization.to_tensor %[[VAL_14]] restrict writable : memref<256xbf16>
// CHECK: memref.tensor_store %[[VAL_15]], %[[VAL_13]] : memref<256xbf16, strided<[?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_15]] in writable %[[VAL_13]]
// CHECK: scf.yield %[[VAL_12]] : index
// CHECK: }
// CHECK: return
Expand Down
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_for_used_before_update.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,7 @@ module {
// CHECK: %[[VAL_14:.*]] = memref.alloc() : memref<256xbf16>
// CHECK: memref.copy %[[VAL_12]], %[[VAL_14]] : memref<256xbf16, strided<[?], offset: ?>> to memref<256xbf16>
// CHECK: %[[VAL_15:.*]] = bufferization.to_tensor %[[VAL_14]] restrict writable : memref<256xbf16>
// CHECK: memref.tensor_store %[[VAL_15]], %[[VAL_12]] : memref<256xbf16, strided<[?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_15]] in writable %[[VAL_12]]
// CHECK: %[[VAL_16:.*]] = arith.addi %[[VAL_13]], %[[VAL_8]] : index
// CHECK: %[[VAL_17:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_16]]], sizes: [256], strides: {{\[}}%[[VAL_4]]] : memref<*xbf16> to memref<256xbf16, strided<[?], offset: ?>>
// CHECK: scf.yield %[[VAL_17]], %[[VAL_16]] : memref<256xbf16, strided<[?], offset: ?>>, index
Expand Down
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_loopback.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,6 @@ module {
// CHECK: %[[VAL_11:.*]] = memref.alloc() : memref<4x256xbf16>
// CHECK: memref.copy %[[VAL_8]], %[[VAL_11]] : memref<4x256xbf16, strided<[1, ?], offset: ?>> to memref<4x256xbf16>
// CHECK: %[[VAL_12:.*]] = bufferization.to_tensor %[[VAL_11]] restrict writable : memref<4x256xbf16>
// CHECK: memref.tensor_store %[[VAL_12]], %[[VAL_10]] : memref<4x256xbf16, strided<[1, ?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_12]] in writable %[[VAL_10]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_mul_const_const.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,6 @@ module {
// CHECK: %[[VAL_13:.*]] = memref.alloc() : memref<1024xbf16>
// CHECK: memref.copy %[[VAL_10]], %[[VAL_13]] : memref<1024xbf16, strided<[?], offset: ?>> to memref<1024xbf16>
// CHECK: %[[VAL_14:.*]] = bufferization.to_tensor %[[VAL_13]] restrict writable : memref<1024xbf16>
// CHECK: memref.tensor_store %[[VAL_14]], %[[VAL_12]] : memref<1024xbf16, strided<[1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_14]] in writable %[[VAL_12]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_mul_value_const.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,6 @@ module {
// CHECK: %[[VAL_18:.*]] = memref.alloc() : memref<1024xbf16>
// CHECK: memref.copy %[[VAL_15]], %[[VAL_18]] : memref<1024xbf16, strided<[?], offset: ?>> to memref<1024xbf16>
// CHECK: %[[VAL_19:.*]] = bufferization.to_tensor %[[VAL_18]] restrict writable : memref<1024xbf16>
// CHECK: memref.tensor_store %[[VAL_19]], %[[VAL_17]] : memref<1024xbf16, strided<[1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_19]] in writable %[[VAL_17]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_nested.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,6 +68,6 @@ module {
// CHECK: %[[VAL_26:.*]] = arith.index_cast %[[VAL_1]] : i32 to index
// CHECK: %[[VAL_27:.*]] = arith.addi %[[VAL_25]], %[[VAL_26]] : index
// CHECK: %[[VAL_28:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_27]]], sizes: [4, 256], strides: [3, %[[VAL_5]]] : memref<*xbf16> to memref<4x256xbf16, strided<[3, ?], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_29:.*]], %[[VAL_28]] : memref<4x256xbf16, strided<[3, ?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_29:.*]] in writable %[[VAL_28]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_reshape_broadcast.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,6 @@ module {
// CHECK: %[[VAL_8:.*]] = memref.alloc() : memref<256x128xbf16>
// CHECK: memref.copy %[[VAL_7]], %[[VAL_8]] : memref<256x128xbf16, strided<[1, ?], offset: 6656>> to memref<256x128xbf16>
// CHECK: %[[VAL_9:.*]] = bufferization.to_tensor %[[VAL_8]] restrict writable : memref<256x128xbf16>
// CHECK: memref.tensor_store %[[VAL_9]], %[[VAL_7]] : memref<256x128xbf16, strided<[1, ?], offset: 6656>>
// CHECK: bufferization.materialize_in_destination %[[VAL_9]] in writable %[[VAL_7]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_scalar_broadcast.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,6 +60,6 @@ module {
// CHECK: %[[VAL_17:.*]] = arith.muli %[[ARG_8]], %[[VAL_3]] : i32
// CHECK: %[[VAL_18:.*]] = arith.index_cast %[[VAL_17]] : i32 to index
// CHECK: %[[VAL_19:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_18]]], sizes: [1024, 1024], strides: [1, 1] : memref<*xf32> to memref<1024x1024xf32, strided<[1, 1], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_20:.*]], %[[VAL_19]] : memref<1024x1024xf32, strided<[1, 1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_20:.*]] in writable %[[VAL_19]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_scalar_for.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,6 +65,6 @@ module {
// CHECK: %[[VAL_35:.*]] = arith.muli %[[ARG_8]], %[[VAL_3]] : i32
// CHECK: %[[VAL_36:.*]] = arith.index_cast %[[VAL_35]] : i32 to index
// CHECK: %[[VAL_37:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_36]]], sizes: [1024], strides: [1] : memref<*xf32> to memref<1024xf32, strided<[1], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_38:.*]]#0, %[[VAL_37]] : memref<1024xf32, strided<[1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_38:.*]]#0 in writable %[[VAL_37]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_scalar_for_2d.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,6 +87,6 @@ module {
// CHECK: %[[VAL_38:.*]] = arith.index_cast %[[VAL_37]] : i32 to index
// CHECK: %[[VAL_39:.*]] = arith.addi %[[VAL_38]], %[[VAL_8]] : index
// CHECK: %[[VAL_40:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_39]]], sizes: [128, 128], strides: [1, 1] : memref<*xf32> to memref<128x128xf32, strided<[1, 1], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_41:.*]]#0, %[[VAL_40]] : memref<128x128xf32, strided<[1, 1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_41:.*]]#0 in writable %[[VAL_40]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_scalar_nested.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,6 +52,6 @@ module {
// CHECK: %[[VAL_23:.*]] = arith.muli %[[ARG_8]], %[[VAL_3]] : i32
// CHECK: %[[VAL_24:.*]] = arith.index_cast %[[VAL_23]] : i32 to index
// CHECK: %[[VAL_25:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_24]]], sizes: [1024], strides: [1] : memref<*xf32> to memref<1024xf32, strided<[1], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_26:.*]], %[[VAL_25]] : memref<1024xf32, strided<[1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_26:.*]] in writable %[[VAL_25]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_scalar_splat.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,6 @@ module {
// CHECK: %[[VAL_17:.*]] = arith.muli %[[ARG_8]], %[[VAL_3]] : i32
// CHECK: %[[VAL_18:.*]] = arith.index_cast %[[VAL_17]] : i32 to index
// CHECK: %[[VAL_19:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_18]]], sizes: [1024], strides: [1] : memref<*xf32> to memref<1024xf32, strided<[1], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_20:.*]], %[[VAL_19]] : memref<1024xf32, strided<[1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_20:.*]] in writable %[[VAL_19]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/addptr_scalar_splat_2d.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,6 +51,6 @@ module {
// CHECK: %[[VAL_20:.*]] = arith.index_cast %[[VAL_19]] : i32 to index
// CHECK: %[[VAL_21:.*]] = arith.addi %[[VAL_20]], %[[VAL_8]] : index
// CHECK: %[[VAL_22:.*]] = memref.reinterpret_cast %[[VAL_0]] to offset: {{\[}}%[[VAL_21]]], sizes: [128, 128], strides: [1, 1] : memref<*xf32> to memref<128x128xf32, strided<[1, 1], offset: ?>>
// CHECK: memref.tensor_store %[[VAL_23:.*]], %[[VAL_22]] : memref<128x128xf32, strided<[1, 1], offset: ?>>
// CHECK: bufferization.materialize_in_destination %[[VAL_23:.*]] in writable %[[VAL_22]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/arith_not_ptr_arith.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,6 @@ module {
// CHECK: %[[VAL_15:.*]] = arith.addi %[[VAL_12]], %[[VAL_13]] : i32
// CHECK: linalg.yield %[[VAL_15]] : i32
// CHECK: } -> tensor<1024xi32>
// CHECK: memref.tensor_store %[[VAL_16:.*]], %[[VAL_6]] : memref<1024xi32, strided<[1]>>
// CHECK: bufferization.materialize_in_destination %[[VAL_16:.*]] in writable %[[VAL_6]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/bitcast.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@ module {
// CHECK: [[VAR_5_:%.+]] = arith.bitcast %in : i32 to f32
// CHECK: linalg.yield [[VAR_5_]] : f32
// CHECK: } -> tensor<1024xf32>
// CHECK: memref.tensor_store [[VAR_2_]], [[RC_0_]] : memref<1024xf32, strided<[1]>>
// CHECK: bufferization.materialize_in_destination [[VAR_2_]] in writable [[RC_0_]]
// CHECK: return
// CHECK: }
// CHECK: }
Expand Down
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/block_ptr_advance.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,7 +87,7 @@ module {
// CHECK: %15 = arith.muli %11, %13 : index
// CHECK: %16 = arith.addi %14, %15 : index
// CHECK: %reinterpret_cast_1 = memref.reinterpret_cast %arg2 to offset: [%16], sizes: [128, 64], strides: [%12, %13] : memref<*xbf16> to memref<128x64xbf16, strided<[?, ?], offset: ?>>
// CHECK: memref.tensor_store %8#0, %reinterpret_cast_1 : memref<128x64xbf16, strided<[?, ?], offset: ?>>
// CHECK: bufferization.materialize_in_destination %8#0 in writable %reinterpret_cast_1
// CHECK: return
// CHECK: }
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/convert_1d_elemwise_arith_binary.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,6 +67,6 @@ module {
// CHECK: %[[VAL_48:.*]] = arith.select %[[VAL_44]], %[[VAL_45]], %[[VAL_46]] : f32
// CHECK: linalg.yield %[[VAL_48]] : f32
// CHECK: } -> tensor<1024xf32>
// CHECK: memref.tensor_store %[[VAL_49:.*]], %[[VAL_2]] : memref<1024xf32>
// CHECK: bufferization.materialize_in_destination %[[VAL_49:.*]] in writable %[[VAL_2]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/convert_1d_elemwise_arith_ternary.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,6 @@ module {
// CHECK: %[[VAL_21:.*]] = arith.select %[[VAL_17]], %[[VAL_18]], %[[VAL_19]] : f32
// CHECK: linalg.yield %[[VAL_21]] : f32
// CHECK: } -> tensor<1024xf32>
// CHECK: memref.tensor_store %[[VAL_22:.*]], %[[VAL_3]] : memref<1024xf32>
// CHECK: bufferization.materialize_in_destination %[[VAL_22:.*]] in writable %[[VAL_3]]
// CHECK: return
// CHECK: }
10 changes: 5 additions & 5 deletions test/Conversion/TritonToLinalg/convert_1d_elemwise_arith_unary.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,10 +79,10 @@ module {
// CHECK: %[[VAL_42:.*]] = math.sqrt %[[VAL_40]] : f32
// CHECK: linalg.yield %[[VAL_42]] : f32
// CHECK: } -> tensor<1024xf32>
// CHECK: memref.tensor_store %[[VAL_43:.*]], %[[VAL_3]] : memref<1024xbf16>
// CHECK: memref.tensor_store %[[VAL_44:.*]], %[[VAL_4]] : memref<1024xf32>
// CHECK: memref.tensor_store %[[VAL_45:.*]], %[[VAL_5]] : memref<1024xf32>
// CHECK: memref.tensor_store %[[VAL_46:.*]], %[[VAL_6]] : memref<1024xf32>
// CHECK: memref.tensor_store %[[VAL_47:.*]], %[[VAL_7]] : memref<1024xf32>
// CHECK: bufferization.materialize_in_destination %[[VAL_43:.*]] in writable %[[VAL_3]]
// CHECK: bufferization.materialize_in_destination %[[VAL_44:.*]] in writable %[[VAL_4]]
// CHECK: bufferization.materialize_in_destination %[[VAL_45:.*]] in writable %[[VAL_5]]
// CHECK: bufferization.materialize_in_destination %[[VAL_46:.*]] in writable %[[VAL_6]]
// CHECK: bufferization.materialize_in_destination %[[VAL_47:.*]] in writable %[[VAL_7]]
// CHECK: return
// CHECK: }
4 changes: 2 additions & 2 deletions test/Conversion/TritonToLinalg/convert_2d_elemwise_arith_binary.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,7 @@ module {
// CHECK: %[[VAL_22:.*]] = arith.subf %[[VAL_19]], %[[VAL_20]] : f32
// CHECK: linalg.yield %[[VAL_22]] : f32
// CHECK: } -> tensor<128x128xf32>
// CHECK: memref.tensor_store %[[VAL_23:.*]], %[[VAL_2]] : memref<128x128xf32>
// CHECK: memref.tensor_store %[[VAL_24:.*]], %[[VAL_3]] : memref<128x128xf32>
// CHECK: bufferization.materialize_in_destination %[[VAL_23:.*]] in writable %[[VAL_2]]
// CHECK: bufferization.materialize_in_destination %[[VAL_24:.*]] in writable %[[VAL_3]]
// CHECK: return
// CHECK: }
2 changes: 1 addition & 1 deletion test/Conversion/TritonToLinalg/convert_2d_elemwise_arith_ternary.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,6 @@ module {
// CHECK: %[[VAL_21:.*]] = arith.select %[[VAL_17]], %[[VAL_18]], %[[VAL_19]] : f32
// CHECK: linalg.yield %[[VAL_21]] : f32
// CHECK: } -> tensor<128x128xf32>
// CHECK: memref.tensor_store %[[VAL_22:.*]], %[[VAL_3]] : memref<128x128xf32>
// CHECK: bufferization.materialize_in_destination %[[VAL_22:.*]] in writable %[[VAL_3]]
// CHECK: return
// CHECK: }
Loading