Page MenuHomePhabricator

herhut (Stephan Herhut)
Animal

Projects

User does not belong to any projects.

User Details

User Since
Jan 7 2020, 7:38 AM (31 w, 3 d)

Recent Activity

Tue, Aug 11

herhut added a comment to D85634: [MLIR][Standard] Fix cast materialization for index types.

Is there anything about this specific to index type? Can this issue happen with any user-defined type?

The solution in this revision looks like a hack. Returning the input when a materialization is requested defeats the whole purpose of materializations, and only works because the assert isn't turned on yet. Realistically, the underlying issue can happen with any type right now. I have a proper fix in the pipeline that should solve this case, where the conversion infra removes casts that become redundant.

Tue, Aug 11, 3:22 AM · Restricted Project
herhut added a comment to D85634: [MLIR][Standard] Fix cast materialization for index types.

The simultaneous application of patterns from SCF to standard and from standard to LLVM rewrites the following to invalid IR.

Do we actually need a simultaneous application?

Tue, Aug 11, 3:08 AM · Restricted Project

Mon, Aug 10

herhut accepted D85631: [MLIR] Adding gpu.host_register op and lower it to a runtime call..

Thanks!

Mon, Aug 10, 10:59 AM · Restricted Project
herhut accepted D85073: [MLIR] Make gpu.launch_func rewrite pattern part of the LLVM lowering pass..
Mon, Aug 10, 7:50 AM · Restricted Project
herhut added inline comments to D85513: [mlir] WIP: Added support for loops in BufferPlacement transformation..
Mon, Aug 10, 3:50 AM · Restricted Project
herhut added inline comments to D85634: [MLIR][Standard] Fix cast materialization for index types.
Mon, Aug 10, 3:00 AM · Restricted Project

Fri, Aug 7

herhut added inline comments to D85518: BEGIN_PUBLIC [mlir] Add support for unranked case for `tensor_store` and `tensor_load` ops. END_PUBLIC.
Fri, Aug 7, 6:13 AM · Restricted Project
herhut accepted D85518: BEGIN_PUBLIC [mlir] Add support for unranked case for `tensor_store` and `tensor_load` ops. END_PUBLIC.
Fri, Aug 7, 5:15 AM · Restricted Project
herhut accepted D85306: [MLIR][Shape] Add custom assembly format for `shape.any`.

Thanks for cleaning this up!

Fri, Aug 7, 2:43 AM · Restricted Project

Thu, Aug 6

herhut accepted D85361: [mlir] Lower DimOp to LLVM for unranked memrefs..

Thanks for clarifying!

Thu, Aug 6, 2:41 AM · Restricted Project
herhut added inline comments to D85361: [mlir] Lower DimOp to LLVM for unranked memrefs..
Thu, Aug 6, 2:25 AM · Restricted Project
herhut added inline comments to D85361: [mlir] Lower DimOp to LLVM for unranked memrefs..
Thu, Aug 6, 12:33 AM · Restricted Project

Wed, Aug 5

herhut added a comment to D85073: [MLIR] Make gpu.launch_func rewrite pattern part of the LLVM lowering pass..

Nice, some nits.

Wed, Aug 5, 12:27 AM · Restricted Project
herhut added a comment to D85133: [mlir] Extend BufferAssignmentTypeConverter with result conversion callbacks.

Just some quick comments.

Wed, Aug 5, 12:14 AM · Restricted Project

Tue, Aug 4

herhut accepted D85273: [mlir] Lower RankOp to LLVM for unranked memrefs..

Thanks!

Tue, Aug 4, 11:56 PM · Restricted Project
herhut accepted D84946: [MLIR] Change GpuLaunchFuncToGpuRuntimeCallsPass to wrap a RewritePattern with the same functionality..
Tue, Aug 4, 6:47 AM · Restricted Project

Fri, Jul 31

herhut accepted D85019: [mlir] translate types between MLIR LLVM dialect and LLVM IR.

It is a bit sad the we have two dynamic dispatch implementations but I guess there is no way around it. A LLVMTypeNew::translate overloaded method would be great :)

Fri, Jul 31, 9:02 AM · Restricted Project
herhut added inline comments to D84946: [MLIR] Change GpuLaunchFuncToGpuRuntimeCallsPass to wrap a RewritePattern with the same functionality..
Fri, Jul 31, 9:02 AM · Restricted Project
herhut accepted D85027: [MLIR][Shape] Lower `shape.broadcast` to `scf`.

Looks good. memref<index> has landed upstream so please fix before landing.

Fri, Jul 31, 9:02 AM · Restricted Project
herhut added a comment to D84933: [MLIR][Shape] Allow unsafe `shape.broadcast`.

Thank you!

Fri, Jul 31, 7:18 AM · Restricted Project
herhut accepted D85017: [mlir] Extended Buffer Assignment to support AllocaOps..
Fri, Jul 31, 5:45 AM · Restricted Project
herhut added inline comments to D84946: [MLIR] Change GpuLaunchFuncToGpuRuntimeCallsPass to wrap a RewritePattern with the same functionality..
Fri, Jul 31, 3:29 AM · Restricted Project
herhut accepted D84933: [MLIR][Shape] Allow unsafe `shape.broadcast`.

Thanks. With one more test this is good to land!

Fri, Jul 31, 1:51 AM · Restricted Project

Thu, Jul 30

herhut added inline comments to D84933: [MLIR][Shape] Allow unsafe `shape.broadcast`.
Thu, Jul 30, 9:09 AM · Restricted Project
herhut committed rG85defd23aa09: [mlir][shape] Use memref of index in shape lowering (authored by herhut).
[mlir][shape] Use memref of index in shape lowering
Thu, Jul 30, 6:20 AM
herhut closed D84938: [mlir][shape] Use memref of index in shape lowering.
Thu, Jul 30, 6:20 AM · Restricted Project
herhut added a comment to D84918: [mlir] Add TFFramework dialect to DialectSymbolRegistry..

Looks reasonable to me. However, we are registering this for a dialect that is in very local use. Maybe @rriddle can advise.

Thu, Jul 30, 5:46 AM · Restricted Project
herhut requested review of D84938: [mlir][shape] Use memref of index in shape lowering.
Thu, Jul 30, 5:38 AM · Restricted Project
herhut committed rGe12db3ed997d: [mlir] Allow index as element type of memref (authored by herhut).
[mlir] Allow index as element type of memref
Thu, Jul 30, 5:37 AM
herhut closed D84934: [mlir] Allow index as element type of memref.
Thu, Jul 30, 5:36 AM · Restricted Project
herhut accepted D84917: [mlir] NFC: Expose `getElementPtrType` and `getSizes` methods of AllocOpLowering..

Please fix linter error and wait for @ftynse to comment.

Thu, Jul 30, 5:23 AM · Restricted Project
herhut added inline comments to D84933: [MLIR][Shape] Allow unsafe `shape.broadcast`.
Thu, Jul 30, 5:21 AM · Restricted Project
herhut requested review of D84934: [mlir] Allow index as element type of memref.
Thu, Jul 30, 5:03 AM · Restricted Project
herhut added inline comments to D84917: [mlir] NFC: Expose `getElementPtrType` and `getSizes` methods of AllocOpLowering..
Thu, Jul 30, 4:18 AM · Restricted Project
herhut added inline comments to D84917: [mlir] NFC: Expose `getElementPtrType` and `getSizes` methods of AllocOpLowering..
Thu, Jul 30, 2:03 AM · Restricted Project

Wed, Jul 29

herhut accepted D84861: [mlir] fix error handling in rocm runtime wrapper.

Thanks for fixing this!

Wed, Jul 29, 11:04 AM · Restricted Project
herhut accepted D84852: [MLIR][Shape] Limit `shape.rank` lowering to its extent tensor variant.
Wed, Jul 29, 8:25 AM · Restricted Project
herhut added inline comments to D84832: [MLIR] Forward extent tensors through shape.broadcast..
Wed, Jul 29, 6:12 AM · Restricted Project
herhut committed rG823ffef00915: [mlir][Standard] Allow unranked memrefs as operands to dim and rank (authored by herhut).
[mlir][Standard] Allow unranked memrefs as operands to dim and rank
Wed, Jul 29, 5:43 AM
herhut closed D84790: [mlir][Standard] Allow unranked memrefs as operands to dim and rank.
Wed, Jul 29, 5:43 AM · Restricted Project
herhut committed rG5d9f33aaa00c: [MLIR][Shape] Add conversion for missing ops to standard (authored by herhut).
[MLIR][Shape] Add conversion for missing ops to standard
Wed, Jul 29, 4:05 AM
herhut closed D84745: [MLIR][Shape] Add conversion for missing ops to standard.
Wed, Jul 29, 4:05 AM · Restricted Project
herhut updated the diff for D84790: [mlir][Standard] Allow unranked memrefs as operands to dim and rank.

Fix test

Wed, Jul 29, 4:04 AM · Restricted Project
herhut accepted D84443: [MLIR][Shape] Limit shape to standard lowerings to their supported types.
Wed, Jul 29, 4:04 AM · Restricted Project
herhut added inline comments to D84790: [mlir][Standard] Allow unranked memrefs as operands to dim and rank.
Wed, Jul 29, 4:04 AM · Restricted Project
herhut updated the diff for D84745: [MLIR][Shape] Add conversion for missing ops to standard.

Rebase

Wed, Jul 29, 4:04 AM · Restricted Project

Tue, Jul 28

herhut added reviewers for D84790: [mlir][Standard] Allow unranked memrefs as operands to dim and rank: ftynse, mehdi_amini.
Tue, Jul 28, 12:36 PM · Restricted Project
herhut requested review of D84790: [mlir][Standard] Allow unranked memrefs as operands to dim and rank.
Tue, Jul 28, 12:35 PM · Restricted Project
herhut updated the diff for D84745: [MLIR][Shape] Add conversion for missing ops to standard.

Comments

Tue, Jul 28, 10:09 AM · Restricted Project
herhut added inline comments to D84745: [MLIR][Shape] Add conversion for missing ops to standard.
Tue, Jul 28, 10:00 AM · Restricted Project
herhut added inline comments to D84745: [MLIR][Shape] Add conversion for missing ops to standard.
Tue, Jul 28, 5:49 AM · Restricted Project
herhut published D84745: [MLIR][Shape] Add conversion for missing ops to standard for review.
Tue, Jul 28, 5:44 AM · Restricted Project
herhut committed rG6d10d317d8b0: [MLIR][Shape] Support transforming shape.num_elements on tensors (authored by herhut).
[MLIR][Shape] Support transforming shape.num_elements on tensors
Tue, Jul 28, 5:13 AM
herhut closed D84744: [MLIR][Shape] Support transforming shape.num_elements on tensors.
Tue, Jul 28, 5:13 AM · Restricted Project
herhut added a comment to D84744: [MLIR][Shape] Support transforming shape.num_elements on tensors.

Thanks!

Tue, Jul 28, 5:11 AM · Restricted Project
herhut updated the diff for D84744: [MLIR][Shape] Support transforming shape.num_elements on tensors.

Comments on formatting

Tue, Jul 28, 5:10 AM · Restricted Project
herhut published D84744: [MLIR][Shape] Support transforming shape.num_elements on tensors for review.
Tue, Jul 28, 4:27 AM · Restricted Project
herhut accepted D84619: [mlir][GPUToSPIRV] Add a test pass to set workgroup size for kernel functions..
Tue, Jul 28, 1:51 AM · Restricted Project
herhut accepted D84660: Clean up cuda-runtime-wrappers API..

Thanks. I only have a minor nit.

Tue, Jul 28, 1:39 AM · Restricted Project
herhut added inline comments to D84632: [MLIR] Add shape.extend op..
Tue, Jul 28, 1:34 AM · Restricted Project

Mon, Jul 27

herhut accepted D84529: [mlir][NFC] Polish copy removal transform.

Maybe mark it as NFC. Looks good with @bondhugula's comment addressed.

Mon, Jul 27, 2:12 AM · Restricted Project
herhut accepted D84441: [MLIR][Shape] Allow `shape.add` to operate on indices.
Mon, Jul 27, 2:10 AM · Restricted Project
herhut accepted D84439: [MLIR][Shape] Refactor verification.
Mon, Jul 27, 2:08 AM · Restricted Project
herhut accepted D84583: [MLIR] NFC: Rename mcuMemHostRegister* to mgpuMemHostRegister* to make it consistent with the other cuda-runner functions and ROCm..

Thanks.

Mon, Jul 27, 2:05 AM · Restricted Project

Wed, Jul 22

herhut accepted D84161: [MLIR][Shape] Simplify shape lowering.
Wed, Jul 22, 12:35 AM · Restricted Project
herhut accepted D84157: [MLIR][Shape] Generalize `shape.const_shape` to extent tensors.
Wed, Jul 22, 12:32 AM · Restricted Project

Tue, Jul 21

herhut added a comment to D84217: [mlir] Add shape.with_shape op.

Adding the op SGTM.

Tue, Jul 21, 1:43 AM · Restricted Project

Mon, Jul 20

herhut added a comment to D83799: [MLIR] Support function signature conversions with tuples of results types.

What you say is completely doable and meaningful as a rewrite in HLO, but I think you are assuming that eliminating such tuples completely is always desirable or possible before applying buffer conversion. On a side note, the real source of such return value tuples in TF/MLIR context is -tf-promote-resources-to-args, which creates a single tuple value for all TF outputs and resource var output aliases. The handling of the tuples itself is about 6-7 lines of code here, but there is a mix of concerns (as you suggest) if one assumes the tuples could always be eliminated and it'd be desirable to do so. But that's a big assumption because such tuples may only ultimately go away while we do the tensor to buffer conversion (where return value to output arg is happening). I have a conversion pattern on the HLO side that deals with xla_hlo.tuple (now mhlo.tuple) pretty much doing what the test dialect tensor tuple op is doing here. A tuple tensor being returned could potentially be an operand for another intermediate op for example. How do we remove tuples on the tensor op form without breaking op semantics?

Mon, Jul 20, 7:39 AM · Restricted Project
herhut accepted D84160: [MLIR][Shape] Allow for `shape_of` to return extent tensors.

We should discuss whether we in the long run want to allow known-size tensors for known-rank inputs.

Mon, Jul 20, 6:02 AM · Restricted Project
herhut accepted D84158: [MLIR][Shape] Allow `shape.get_extent` to operate on extent tensors.
Mon, Jul 20, 5:54 AM · Restricted Project
herhut added a comment to D84157: [MLIR][Shape] Generalize `shape.const_shape` to extent tensors.

Shouldn't shape.const_shape always return a tensor? It cannot return an error currently, right?

Mon, Jul 20, 5:52 AM · Restricted Project
herhut accepted D84156: [MLIR][Shape] Allow `shape.rank` to accept extent tensors `tensor?xindex>`.
Mon, Jul 20, 5:49 AM · Restricted Project
herhut accepted D84155: [MLIR][Shape] Allow `cstr_broadcastable` to accept extent tensors.

I wonder whether we need to duplicate all the tests to also cover tensor types. Under the hood it would be the same code-path anyway, though.

Mon, Jul 20, 5:48 AM · Restricted Project
herhut accepted D84149: [mlir] Support translating function linkage between MLIR and LLVM IR.

Thanks!

Mon, Jul 20, 5:01 AM · Restricted Project

Thu, Jul 16

herhut accepted D83943: [MLIR][Shape] Allow `shape.reduce` to operate on extent tensors.

Thanks!

Thu, Jul 16, 6:14 AM · Restricted Project
herhut accepted D83944: [MLIR] Lower `shape.reduce` to `scf.for` only when argument is `tensor<?xindex>`.

Thanks for cleaning this up!

Thu, Jul 16, 6:13 AM · Restricted Project
herhut added a comment to D83943: [MLIR][Shape] Allow `shape.reduce` to operate on extent tensors.

Nice, thank you!

Thu, Jul 16, 5:28 AM · Restricted Project
herhut accepted D83932: [MLIR][Shape] Use callback builder again.

Thanks!

Thu, Jul 16, 2:57 AM · Restricted Project
herhut added a comment to D83799: [MLIR] Support function signature conversions with tuples of results types.

Could we just get rid of tuples completely and instead use multiple return values from the start? Do we actually need tuples as a data type?

Are you referring to TF? We are in MLIR here, and this is client agnostic. The tuples were there to model HLO ops many of which pack into and extract from tuples.

Thu, Jul 16, 2:12 AM · Restricted Project
herhut added a reviewer for D83799: [MLIR] Support function signature conversions with tuples of results types: herhut.

I don't think there is a document; I can update it if one exists. Waiting for @dfki-ehna to comment where.

Thu, Jul 16, 12:32 AM · Restricted Project

Jul 15 2020

herhut committed rG8ef47244b95f: [mlir][shape] Fold shape.broadcast with one scalar operand (authored by herhut).
[mlir][shape] Fold shape.broadcast with one scalar operand
Jul 15 2020, 9:49 AM
herhut closed D83854: [mlir][shape] Fold shape.broadcast with scalar operands.
Jul 15 2020, 9:49 AM · Restricted Project
herhut updated the diff for D83854: [mlir][shape] Fold shape.broadcast with scalar operands.

Add back test and unbreak.

Jul 15 2020, 9:49 AM · Restricted Project
herhut committed rG412b60531edd: [mlir][shape] Mark some operations as commutative (authored by herhut).
[mlir][shape] Mark some operations as commutative
Jul 15 2020, 9:33 AM
herhut closed D83856: [mlir][shape] Mark some operations as commutative.
Jul 15 2020, 9:33 AM · Restricted Project
herhut added inline comments to D83854: [mlir][shape] Fold shape.broadcast with scalar operands.
Jul 15 2020, 8:58 AM · Restricted Project
herhut updated the diff for D83854: [mlir][shape] Fold shape.broadcast with scalar operands.

Comments

Jul 15 2020, 8:56 AM · Restricted Project
herhut accepted D83873: [MLIR][Shape] Fix `shape_of` lowering to `scf`.
Jul 15 2020, 8:33 AM · Restricted Project
herhut updated the diff for D83856: [mlir][shape] Mark some operations as commutative.

Comments.

Jul 15 2020, 8:31 AM · Restricted Project
herhut accepted D82528: [MLIR][Shape] Add `shape.shape_eq` operation.

I think this is ready to land after yesterday's discussion.

Jul 15 2020, 3:13 AM · Restricted Project
herhut added a reviewer for D83856: [mlir][shape] Mark some operations as commutative: frgossen.
Jul 15 2020, 12:34 AM · Restricted Project
Herald added a project to D83856: [mlir][shape] Mark some operations as commutative: Restricted Project.
Jul 15 2020, 12:33 AM · Restricted Project
herhut added a reviewer for D83854: [mlir][shape] Fold shape.broadcast with scalar operands: tpopp.
Jul 15 2020, 12:28 AM · Restricted Project
Herald added a project to D83854: [mlir][shape] Fold shape.broadcast with scalar operands: Restricted Project.
Jul 15 2020, 12:27 AM · Restricted Project
herhut committed rG1919c8bfe837: Make linalg::ReshapeOp implement ViewLikeOpInterface (authored by herhut).
Make linalg::ReshapeOp implement ViewLikeOpInterface
Jul 15 2020, 12:24 AM
herhut closed D83773: Make linalg::ReshapeOp implement ViewLikeOpInterface.
Jul 15 2020, 12:24 AM · Restricted Project

Jul 14 2020

herhut updated the diff for D83773: Make linalg::ReshapeOp implement ViewLikeOpInterface.

Format tablegen.

Jul 14 2020, 7:27 AM · Restricted Project
Herald added a project to D83773: Make linalg::ReshapeOp implement ViewLikeOpInterface: Restricted Project.
Jul 14 2020, 7:25 AM · Restricted Project

Jul 10 2020

herhut added a comment to D82574: Merge TableGen files used for clang options.

Could you add the normalization back? This is in line with the comment to make sure the old and new files align.

Jul 10 2020, 1:51 AM · Restricted Project

Jul 8 2020

herhut accepted D83385: [mlir] Add ViewLikeOpInterface to std.memref_cast..

Can you update the commit message so that it has a summary (first sentence) and description?

Jul 8 2020, 5:11 AM · Restricted Project