Page MenuHomePhabricator

guraypp (guray ozen)
User

Projects

User does not belong to any projects.

User Details

User Since
May 30 2018, 9:30 AM (234 w, 5 d)

Recent Activity

Yesterday

guraypp committed rG135977c92a17: [mlir] Export `collapseGenericOpIterationDims` (NFC) (authored by guraypp).
[mlir] Export `collapseGenericOpIterationDims` (NFC)
Mon, Nov 28, 4:54 AM · Restricted Project, Restricted Project
guraypp closed D138697: [mlir] Export `collapseGenericOpIterationDims` (NFC).
Mon, Nov 28, 4:54 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138697: [mlir] Export `collapseGenericOpIterationDims` (NFC).

rebase

Mon, Nov 28, 3:49 AM · Restricted Project, Restricted Project

Fri, Nov 25

guraypp requested review of D138697: [mlir] Export `collapseGenericOpIterationDims` (NFC).
Fri, Nov 25, 1:53 AM · Restricted Project, Restricted Project

Thu, Nov 17

guraypp committed rGc5798fae05af: [mlir] [transform] Error for duplicated processor mapping (authored by guraypp).
[mlir] [transform] Error for duplicated processor mapping
Thu, Nov 17, 11:39 PM · Restricted Project, Restricted Project
guraypp closed D138032: [mlir] [transform] Error for duplicated processor mapping.
Thu, Nov 17, 11:39 PM · Restricted Project, Restricted Project
guraypp updated the diff for D138032: [mlir] [transform] Error for duplicated processor mapping.

remove unnecessary include

Thu, Nov 17, 7:16 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138032: [mlir] [transform] Error for duplicated processor mapping.

address @ftynse comments

Thu, Nov 17, 7:13 AM · Restricted Project, Restricted Project

Wed, Nov 16

guraypp committed rG5ce68f4284c6: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase` (authored by guraypp).
[mlir] Introduce `replaceUsesOfWith` to `RewriterBase`
Wed, Nov 16, 8:53 AM · Restricted Project, Restricted Project
guraypp closed D138110: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase`.
Wed, Nov 16, 8:53 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138110: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase`.

Address @mehdi_amini comments

Wed, Nov 16, 8:24 AM · Restricted Project, Restricted Project
guraypp added a reviewer for D138110: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase`: silvas.
Wed, Nov 16, 3:35 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138110: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase`.

get rid of vector and use make_early_inc_range

Wed, Nov 16, 3:28 AM · Restricted Project, Restricted Project
guraypp added inline comments to D138110: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase`.
Wed, Nov 16, 3:20 AM · Restricted Project, Restricted Project
guraypp requested review of D138110: [mlir] Introduce `replaceUsesOfWith` to `RewriterBase`.
Wed, Nov 16, 3:03 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138032: [mlir] [transform] Error for duplicated processor mapping.

rebase
add newline

Wed, Nov 16, 1:02 AM · Restricted Project, Restricted Project
guraypp committed rG63ca939783eb: [mlir] [transform] Fix for RAUW error in transform gpu dialect (authored by guraypp).
[mlir] [transform] Fix for RAUW error in transform gpu dialect
Wed, Nov 16, 12:55 AM · Restricted Project, Restricted Project
guraypp closed D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect.
Wed, Nov 16, 12:55 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect.

address @springerm comments

Wed, Nov 16, 12:51 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect.

Iterate over the uses instead of users
Use SmallVector instead of SetVector

Wed, Nov 16, 12:28 AM · Restricted Project, Restricted Project

Tue, Nov 15

guraypp updated the diff for D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect.

Use vector constructor to build users.
Fix the same potential bug in map_foreach_to_blocks

Tue, Nov 15, 9:50 AM · Restricted Project, Restricted Project
guraypp committed rGbeaffb041c68: [mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom… (authored by guraypp).
[mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom…
Tue, Nov 15, 9:16 AM · Restricted Project, Restricted Project
guraypp closed D138020: [mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom dialect code generator.
Tue, Nov 15, 9:16 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect.

Move invariant bvm.lookup(threadIdx) outside of the loop

Tue, Nov 15, 8:39 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138020: [mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom dialect code generator.

rebase

Tue, Nov 15, 8:28 AM · Restricted Project, Restricted Project
guraypp requested review of D138032: [mlir] [transform] Error for duplicated processor mapping.
Tue, Nov 15, 6:52 AM · Restricted Project, Restricted Project
guraypp added a reviewer for D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect: springerm.
Tue, Nov 15, 6:28 AM · Restricted Project, Restricted Project
guraypp requested review of D138029: [mlir] [transform] Fix for RAUW error in transform gpu dialect.
Tue, Nov 15, 6:28 AM · Restricted Project, Restricted Project
guraypp added a reviewer for D138020: [mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom dialect code generator: ftynse.
Tue, Nov 15, 3:03 AM · Restricted Project, Restricted Project
guraypp updated the diff for D138020: [mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom dialect code generator.

remove accidentally edited lines

Tue, Nov 15, 3:03 AM · Restricted Project, Restricted Project
guraypp requested review of D138020: [mlir][transform] Decouple GPUDeviceMapping attribute from the GPU transfrom dialect code generator.
Tue, Nov 15, 2:59 AM · Restricted Project, Restricted Project

Mon, Nov 14

guraypp accepted D137906: [mlir][Transform]Significantly cleanup scf.foreach_thread and GPU transform permutation handling.

Looks great! Thanks for fixing that.

Mon, Nov 14, 12:39 AM · Restricted Project, Restricted Project

Sat, Nov 12

guraypp committed rGd93be483eaf5: [mlir][transform] Make `tile_to_foreach_thread_op` builder to use ArrayAttr (authored by guraypp).
[mlir][transform] Make `tile_to_foreach_thread_op` builder to use ArrayAttr
Sat, Nov 12, 10:27 AM · Restricted Project, Restricted Project
guraypp closed D137891: [mlir][transform] Make `tile_to_foreach_thread_op` builder to use ArrayAttr.
Sat, Nov 12, 10:27 AM · Restricted Project, Restricted Project
guraypp added a reviewer for D137891: [mlir][transform] Make `tile_to_foreach_thread_op` builder to use ArrayAttr: springerm.
Sat, Nov 12, 4:02 AM · Restricted Project, Restricted Project
guraypp requested review of D137891: [mlir][transform] Make `tile_to_foreach_thread_op` builder to use ArrayAttr.
Sat, Nov 12, 4:02 AM · Restricted Project, Restricted Project

Fri, Nov 11

guraypp accepted D137830: [mlir][bufferize][NFC] Consolidate transform header files.

The code looks good.
Out of curiosity why do we want to combine all the headers files into one?

These transforms are usually used together (I'm cleaning up some code in IREE and noticed a long list of includes). Also to make it consistent with the other dialects, e.g., SCF/Transforms/Transforms.h.

Fri, Nov 11, 5:12 AM · Restricted Project, Restricted Project
guraypp added a comment to D137830: [mlir][bufferize][NFC] Consolidate transform header files.

The code looks good.
Out of curiosity why do we want to combine all the headers files into one?

Fri, Nov 11, 4:58 AM · Restricted Project, Restricted Project
guraypp committed rG0d845660f4e2: [mlir] Fix asan errors in gpu transform dialect (authored by guraypp).
[mlir] Fix asan errors in gpu transform dialect
Fri, Nov 11, 2:57 AM · Restricted Project, Restricted Project

Thu, Nov 10

guraypp committed rG6663f3470417: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to… (authored by guraypp).
[mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to…
Thu, Nov 10, 11:45 PM · Restricted Project, Restricted Project
guraypp closed D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Thu, Nov 10, 11:45 PM · Restricted Project, Restricted Project
guraypp updated the summary of D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Thu, Nov 10, 8:55 AM · Restricted Project, Restricted Project
guraypp updated the summary of D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Thu, Nov 10, 8:52 AM · Restricted Project, Restricted Project
guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

address @ftynse comments

Thu, Nov 10, 8:51 AM · Restricted Project, Restricted Project
guraypp abandoned D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

We implemented a better solution in https://reviews.llvm.org/D137413

Thu, Nov 10, 8:41 AM · Restricted Project, Restricted Project

Wed, Nov 9

guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

rebase

Wed, Nov 9, 9:05 AM · Restricted Project, Restricted Project
guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

minor: fix comments

Wed, Nov 9, 4:52 AM · Restricted Project, Restricted Project
guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

bazel fix

Wed, Nov 9, 3:16 AM · Restricted Project, Restricted Project
guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

Address @rriddle comments, rebase, add description to gpu device mapping attributes

Wed, Nov 9, 2:07 AM · Restricted Project, Restricted Project

Tue, Nov 8

guraypp updated the summary of D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Tue, Nov 8, 10:49 AM · Restricted Project, Restricted Project
guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

Moved DeviceMappingInterface to SCF.

Tue, Nov 8, 10:48 AM · Restricted Project, Restricted Project

Mon, Nov 7

guraypp updated the diff for D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

address ftynse comments

Mon, Nov 7, 8:52 AM · Restricted Project, Restricted Project
guraypp added a comment to D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.

I agree that this goes in the right direction, however, I am unsure what this direction is? Is there a writeup of where this should be going? That would make reviewing these changes easier.

Also, it would be nice to unify this with the mapping attributes on the scf.parallel operation, so that we have one way of doing this. Having a different encoding was fine for the initial experimentation but now that this turns into an interface, we should ensure it works with other scf operations, as well.

Mon, Nov 7, 8:47 AM · Restricted Project, Restricted Project
guraypp updated the summary of D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Mon, Nov 7, 8:42 AM · Restricted Project, Restricted Project
guraypp added inline comments to D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Mon, Nov 7, 8:40 AM · Restricted Project, Restricted Project

Fri, Nov 4

guraypp requested review of D137438: [mlir][transform] remove the barrier before terminator.
Fri, Nov 4, 9:06 AM · Restricted Project, Restricted Project
guraypp updated the diff for D135252: [mlir][transform] Assing blockDim automatically.

update the test

Fri, Nov 4, 7:49 AM · Restricted Project, Restricted Project
guraypp updated the diff for D135252: [mlir][transform] Assing blockDim automatically.

rebase and ping

Fri, Nov 4, 7:43 AM · Restricted Project, Restricted Project
guraypp requested review of D137424: [mlir][transform] Introduce `gpu.map_foreach`.
Fri, Nov 4, 7:27 AM · Restricted Project, Restricted Project
guraypp updated the summary of D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Fri, Nov 4, 5:32 AM · Restricted Project, Restricted Project
guraypp added a comment to D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

@ftynse thanks for clear explanation. As it is different that this current PR, I implement in another in D137413. Let me know what do you think.

Fri, Nov 4, 5:27 AM · Restricted Project, Restricted Project
guraypp added a comment to D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Fri, Nov 4, 5:27 AM · Restricted Project, Restricted Project
guraypp requested review of D137413: [mlir] Introduce device mapper attribute for `thread_dim_map` and `mapped to dims`.
Fri, Nov 4, 5:25 AM · Restricted Project, Restricted Project

Mon, Oct 31

guraypp added a comment to D136931: [mlir][nvvm] Introduce performance tuning directives.
Mon, Oct 31, 5:46 AM · Restricted Project, Restricted Project

Oct 28 2022

guraypp committed rG3ac17449cf98: [mlir][nvvm] Introduce performance tuning directives (authored by guraypp).
[mlir][nvvm] Introduce performance tuning directives
Oct 28 2022, 5:03 AM · Restricted Project, Restricted Project
guraypp closed D136931: [mlir][nvvm] Introduce performance tuning directives.
Oct 28 2022, 5:02 AM · Restricted Project, Restricted Project
guraypp updated the diff for D136931: [mlir][nvvm] Introduce performance tuning directives.

Add attribute verifiers

Oct 28 2022, 4:46 AM · Restricted Project, Restricted Project
guraypp added a comment to D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

Then I have a proposal. As Alex said, let's relex verifier of foreach_thread, for example as follows (here I improved the syntax a bit).

scf.foreach_thread (%bi, %bj)  {
  scf.foreach_thread (%ti, %tj) {
  }  {map = "unit", dimensions = ["x", "y"]}
}  {map = "group", dimensions = ["x", "y"]}

Next, let's add a new generic op to the GPU dialect. One can map named parallelism unit to block and thread.

// transform dialect for GPU
transform.gpu.map_foreach %op "group" to blocks
transform.gpu.map_foreach %op "unit" to threads

Another accelerator transform dialect uses only the outer level parallelism for example.

// transform dialect for other accelerator
transform.other_accelerator.map_foreach %op "group" to threads
Oct 28 2022, 3:52 AM · Restricted Project, Restricted Project
guraypp added a comment to D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

Have you considered making it an actual attribute in the GPU dialect, e.g., an enum attribute? Passing strings around feels a lot like unsafe JSON even if it is a readability improvement.

One aspect here is that the op semantics and transformations should be retargetable.
Is there a way to have something like a BaseForeachThreadTargetEnumAttr for which the GPU dialect could define Thread/Block; the IREE xxx dialect could define Workgroupyy and future ACME dialect could define its own?

A (far from) ideal solution would be to have users define their own names "foo", "bar", "baz" and explicitly have the map_to_target specify the map ["foo" -> threadIdx.x ...].
This is so ugly that I don't think it should be considered but that should give a feeling on what a better preferred way to interact with the system could look like.

If you don't constrain the kind of attributes in the op verifier, any downstream client can use whatever they want there. Lowerings can then check if they support the specific attribute kind and fail early if they don't. If you want more guarantees, introduce an attribute interface. There is no need to invent new infrastructure layers here.

Cool thnks for the guidnce, I had not followed that part of the stack.

Oct 28 2022, 3:36 AM · Restricted Project, Restricted Project
guraypp updated the summary of D136931: [mlir][nvvm] Introduce performance tuning directives.
Oct 28 2022, 2:40 AM · Restricted Project, Restricted Project
guraypp requested review of D136931: [mlir][nvvm] Introduce performance tuning directives.
Oct 28 2022, 2:40 AM · Restricted Project, Restricted Project
guraypp added a comment to D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

Have you considered making it an actual attribute in the GPU dialect, e.g., an enum attribute? Passing strings around feels a lot like unsafe JSON even if it is a readability improvement.

Oct 28 2022, 1:28 AM · Restricted Project, Restricted Project
guraypp updated the diff for D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

remove size of the vector

Oct 28 2022, 1:28 AM · Restricted Project, Restricted Project
guraypp updated the diff for D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

remove comments

Oct 28 2022, 12:14 AM · Restricted Project, Restricted Project

Oct 27 2022

guraypp added a reviewer for D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`: ftynse.
Oct 27 2022, 12:41 PM · Restricted Project, Restricted Project
guraypp updated the diff for D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.

update the description

Oct 27 2022, 10:42 AM · Restricted Project, Restricted Project
guraypp updated the summary of D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.
Oct 27 2022, 9:00 AM · Restricted Project, Restricted Project
guraypp requested review of D136851: [mlir] Use processing unit names for `thread_dim_map` and `mapped to dims`.
Oct 27 2022, 8:59 AM · Restricted Project, Restricted Project

Oct 11 2022

guraypp added inline comments to D135598: [mlir][transform] Add debug messages to gpu transform.
Oct 11 2022, 4:50 AM · Restricted Project, Restricted Project
guraypp added a reviewer for D135598: [mlir][transform] Add debug messages to gpu transform: ftynse.
Oct 11 2022, 4:49 AM · Restricted Project, Restricted Project
guraypp accepted D135661: [mlir] drop unnecssary transform.with_pdl_patterns from tests, NFC.

looks good, thanks for doing that

Oct 11 2022, 4:49 AM · Restricted Project, Restricted Project

Oct 10 2022

guraypp updated the summary of D135598: [mlir][transform] Add debug messages to gpu transform.
Oct 10 2022, 9:01 AM · Restricted Project, Restricted Project
guraypp updated the summary of D135598: [mlir][transform] Add debug messages to gpu transform.
Oct 10 2022, 9:00 AM · Restricted Project, Restricted Project
guraypp requested review of D135598: [mlir][transform] Add debug messages to gpu transform.
Oct 10 2022, 9:00 AM · Restricted Project, Restricted Project
guraypp requested review of D135566: [mlir][transform] Fail if thre is no `thread_dim_mapping`.
Oct 10 2022, 2:11 AM · Restricted Project, Restricted Project

Oct 6 2022

guraypp committed rG040805dc4748: [mlir] Add bar.warp.sync to NVVM (authored by guraypp).
[mlir] Add bar.warp.sync to NVVM
Oct 6 2022, 3:28 AM · Restricted Project, Restricted Project
guraypp closed D135253: [mlir] Add bar.warp.sync to NVVM.
Oct 6 2022, 3:28 AM · Restricted Project, Restricted Project

Oct 5 2022

guraypp added inline comments to D135252: [mlir][transform] Assing blockDim automatically.
Oct 5 2022, 6:10 AM · Restricted Project, Restricted Project
guraypp updated the diff for D135252: [mlir][transform] Assing blockDim automatically.

address nicolasvasilache comments

Oct 5 2022, 6:06 AM · Restricted Project, Restricted Project
guraypp updated the summary of D135252: [mlir][transform] Assing blockDim automatically.
Oct 5 2022, 4:35 AM · Restricted Project, Restricted Project
guraypp committed rGe68a7bed5997: [mlir][transform] Add failing test for GPU transform dialect (authored by guraypp).
[mlir][transform] Add failing test for GPU transform dialect
Oct 5 2022, 4:10 AM · Restricted Project, Restricted Project
guraypp closed D135063: [mlir][transform] Add failing test for GPU transform dialect.
Oct 5 2022, 4:10 AM · Restricted Project, Restricted Project
guraypp updated the diff for D135063: [mlir][transform] Add failing test for GPU transform dialect.

rebase

Oct 5 2022, 4:09 AM · Restricted Project, Restricted Project
guraypp committed rG78305720f387: [mlir][transform][nfc] typo fix (authored by guraypp).
[mlir][transform][nfc] typo fix
Oct 5 2022, 4:06 AM · Restricted Project, Restricted Project
guraypp closed D135242: [mlir][transform][nfc] typo fix.
Oct 5 2022, 4:05 AM · Restricted Project, Restricted Project
guraypp requested review of D135253: [mlir] Add bar.warp.sync to NVVM.
Oct 5 2022, 3:24 AM · Restricted Project, Restricted Project
guraypp requested review of D135252: [mlir][transform] Assing blockDim automatically.
Oct 5 2022, 3:14 AM · Restricted Project, Restricted Project
guraypp updated the diff for D135063: [mlir][transform] Add failing test for GPU transform dialect.

rebase and change error message

Oct 5 2022, 1:35 AM · Restricted Project, Restricted Project
guraypp updated the diff for D135242: [mlir][transform][nfc] typo fix.

rebase

Oct 5 2022, 1:25 AM · Restricted Project, Restricted Project
guraypp added a reviewer for D135242: [mlir][transform][nfc] typo fix: ftynse.
Oct 5 2022, 12:25 AM · Restricted Project, Restricted Project