Page MenuHomePhabricator

christopherbate (Christopher Bate)
User

Projects

User does not belong to any projects.

User Details

User Since
Mar 14 2022, 9:59 PM (9 w, 6 d)

Recent Activity

Today

christopherbate committed rG7085cb6011d4: [mlir][NvGpuToNVVM] Fix byte size calculation in async copy lowering (authored by christopherbate).
[mlir][NvGpuToNVVM] Fix byte size calculation in async copy lowering
Mon, May 23, 9:56 AM · Restricted Project, Restricted Project
christopherbate closed D125838: [mlir][NvGpuToNVVM] Fix byte size calculation in async copy lowering.
Mon, May 23, 9:56 AM · Restricted Project, Restricted Project
christopherbate committed rG334f63e7c39f: [mlir][NvGpuToNVVM] Fix missing i4 support for nvgpu.mma.sync (authored by christopherbate).
[mlir][NvGpuToNVVM] Fix missing i4 support for nvgpu.mma.sync
Mon, May 23, 9:53 AM · Restricted Project, Restricted Project
christopherbate closed D126092: [mlir][NvGpuToNVVM] Fix missing i4 support for nvgpu.mma.sync.
Mon, May 23, 9:53 AM · Restricted Project, Restricted Project

Fri, May 20

christopherbate requested review of D126092: [mlir][NvGpuToNVVM] Fix missing i4 support for nvgpu.mma.sync.
Fri, May 20, 2:11 PM · Restricted Project, Restricted Project
christopherbate updated the diff for D125838: [mlir][NvGpuToNVVM] Fix byte size calculation in async copy lowering.

Added a test for the i4 case.

Fri, May 20, 1:26 PM · Restricted Project, Restricted Project
christopherbate committed rG873a3e2c1d04: [mlir] Add missing NVGPU link dependency to VectorToGPU (authored by christopherbate).
[mlir] Add missing NVGPU link dependency to VectorToGPU
Fri, May 20, 9:47 AM · Restricted Project, Restricted Project
christopherbate committed rG1ca772ed951e: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass (authored by christopherbate).
[MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass
Fri, May 20, 8:44 AM · Restricted Project, Restricted Project
christopherbate closed D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.
Fri, May 20, 8:44 AM · Restricted Project, Restricted Project

Tue, May 17

christopherbate added a comment to D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Addressed Thomas' comments.

Tue, May 17, 4:59 PM · Restricted Project, Restricted Project
christopherbate updated the diff for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Address reviewer comments:
Remove use of GPU Dialect's WarpMatrixFragmentType in NvGpu conversion path.
Fix variable naming in FileCheck tests to allow for CHECK-LABEL.
Cleanup misc typos/small fixes.

Tue, May 17, 4:58 PM · Restricted Project, Restricted Project
christopherbate requested review of D125838: [mlir][NvGpuToNVVM] Fix byte size calculation in async copy lowering.
Tue, May 17, 2:53 PM · Restricted Project, Restricted Project

Fri, May 13

christopherbate added a comment to D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

I think this should complete any functional updates. I still need to address Thomas' comments.

Fri, May 13, 9:14 PM · Restricted Project, Restricted Project
christopherbate retitled D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass from [MLIR][GPU] Add VectorToGPU, GPU to nvvm.mma.sync lowering path to [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.
Fri, May 13, 9:11 PM · Restricted Project, Restricted Project
christopherbate updated the diff for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Small functional changes and additional tests based on TF32 and Int8 testing.

Fri, May 13, 9:09 PM · Restricted Project, Restricted Project

Tue, May 10

christopherbate accepted D125244: [mlir][gpu] Move async copy ops to NVGPU and add caching hints.

LGTM, thanks!

Tue, May 10, 2:18 PM · Restricted Project, Restricted Project

Mon, May 9

christopherbate added a comment to D125244: [mlir][gpu] Move async copy ops to NVGPU and add caching hints.

A couple of minor things, otherwise LGTM.

Mon, May 9, 11:40 AM · Restricted Project, Restricted Project
christopherbate accepted D125241: [mlir][nvvm] Add attribute to nvvm.cpAsyncOp to control l1 bypass.

LGTM, thanks for adding this.

Mon, May 9, 10:40 AM · Restricted Project, Restricted Project

Sun, May 8

christopherbate committed rG9879807393d3: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types (authored by christopherbate).
[mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types
Sun, May 8, 8:59 PM · Restricted Project, Restricted Project
christopherbate closed D124533: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types.
Sun, May 8, 8:58 PM · Restricted Project, Restricted Project

Thu, May 5

christopherbate added inline comments to D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.
Thu, May 5, 12:36 PM · Restricted Project, Restricted Project
christopherbate retitled D124533: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types from [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32 types to [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types.
Thu, May 5, 10:42 AM · Restricted Project, Restricted Project
christopherbate updated the diff for D124533: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types.

Updated the diff to include changes that support TF32 operands.

Thu, May 5, 10:40 AM · Restricted Project, Restricted Project
christopherbate committed rG22c6e7b277fb: [mlir][nvvm] Fix support for tf32 data type in mma.sync (authored by christopherbate).
[mlir][nvvm] Fix support for tf32 data type in mma.sync
Thu, May 5, 10:04 AM · Restricted Project, Restricted Project
christopherbate closed D124975: [mlir][nvvm] Fix support for tf32 data type in mma.sync.
Thu, May 5, 10:03 AM · Restricted Project, Restricted Project
christopherbate updated the diff for D124975: [mlir][nvvm] Fix support for tf32 data type in mma.sync.

Fix type in test name: 'nvvm_mma_m16n8k4_f64_f64' should have been 'nvvm_mma_m16n8k4_tf32_f32'

Thu, May 5, 9:13 AM · Restricted Project, Restricted Project

Wed, May 4

christopherbate requested review of D124975: [mlir][nvvm] Fix support for tf32 data type in mma.sync.
Wed, May 4, 7:19 PM · Restricted Project, Restricted Project
christopherbate added a comment to D124533: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types.

I have a patch for NVVM mma.sync TF32 support. I should land t hat before this one, then update this one to include additional fixes for TF32

Wed, May 4, 2:03 PM · Restricted Project, Restricted Project

Wed, Apr 27

christopherbate requested review of D124533: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types.
Wed, Apr 27, 9:15 AM · Restricted Project, Restricted Project

Tue, Apr 26

christopherbate updated the summary of D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.
Tue, Apr 26, 8:14 PM · Restricted Project, Restricted Project
christopherbate added inline comments to D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.
Tue, Apr 26, 8:11 PM · Restricted Project, Restricted Project
christopherbate updated the diff for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Add documentation / address comments.

Tue, Apr 26, 8:10 PM · Restricted Project, Restricted Project
christopherbate added a comment to D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

I introduced an option use-nvgpu to the VectorToGPU pass rather than introducing a VectorToNvGPU pass. The reason is because a fair amount of code can be shared. I think we should still break this conversion into its own pass (VectorToNvGPU) and factor out the common utilities that can be shared by both. Pending discussion/input from others, I'm happy to do that in this patch or land this patch and follow up with the creation of the new pass (or keep it as is/something else)

Tue, Apr 26, 8:02 PM · Restricted Project, Restricted Project
christopherbate updated the diff for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Rebased on top-of-tree. The first part of this patch was already merged
(the ldmatrix and mma.sync ops in nvgpu dialect). What remains is the lowerings
from vector to nvgpu.

Tue, Apr 26, 8:00 PM · Restricted Project, Restricted Project

Apr 19 2022

christopherbate added inline comments to D123266: [mlir][nvgpu] Add NVGPU dialect (architectural specific gpu dialect).
Apr 19 2022, 12:23 PM · Restricted Project, Restricted Project

Apr 18 2022

christopherbate added a comment to D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Need to rebase this on the most recent movement of mma.sync and ldmatrix ops from the gpu dialect to nvgpu dialect. should have this updated tomorrow morning.

Apr 18 2022, 4:41 PM · Restricted Project, Restricted Project

Apr 14 2022

christopherbate accepted D123824: [mlir][nvgpu] Move mma.sync and ldmatrix in nvgpu dialect.

It looks good to me! thanks

Apr 14 2022, 3:55 PM · Restricted Project, Restricted Project

Apr 13 2022

christopherbate added a comment to D123647: [MLIR][GPU] Add GPU ops nvvm.mma.sync, nvvm.mma.ldmatrix, lane_id.

I updated the revision to address the comments so far, thanks!

Apr 13 2022, 9:08 AM · Restricted Project, Restricted Project
christopherbate updated the diff for D123647: [MLIR][GPU] Add GPU ops nvvm.mma.sync, nvvm.mma.ldmatrix, lane_id.

Address reviewer comments.

Apr 13 2022, 9:07 AM · Restricted Project, Restricted Project

Apr 12 2022

christopherbate requested review of D123647: [MLIR][GPU] Add GPU ops nvvm.mma.sync, nvvm.mma.ldmatrix, lane_id.
Apr 12 2022, 5:05 PM · Restricted Project, Restricted Project

Apr 7 2022

Herald added a reviewer for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass: ThomasRaoux.

Thanks for contributing these. Can you separate out the VectorToGPU changes into another revision?

Apr 7 2022, 12:03 PM · Restricted Project, Restricted Project

Apr 4 2022

christopherbate updated the diff for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Fix ldmatrix use condition to check address space

Apr 4 2022, 10:43 PM · Restricted Project, Restricted Project

Apr 1 2022

christopherbate updated the diff for D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.

Removed dead code

Apr 1 2022, 2:18 PM · Restricted Project, Restricted Project
christopherbate requested review of D122940: [MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass.
Apr 1 2022, 2:11 PM · Restricted Project, Restricted Project

Mar 24 2022

christopherbate added inline comments to D122410: [mlir][NVVM] Add support for nvvm mma.sync ops.
Mar 24 2022, 9:18 PM · Restricted Project, Restricted Project
christopherbate updated the diff for D122410: [mlir][NVVM] Add support for nvvm mma.sync ops.

Address comments and add missing f64 mma cases.

Mar 24 2022, 9:17 PM · Restricted Project, Restricted Project
christopherbate published D122410: [mlir][NVVM] Add support for nvvm mma.sync ops for review.
Mar 24 2022, 10:49 AM · Restricted Project, Restricted Project
christopherbate updated christopherbate.
Mar 24 2022, 9:59 AM

Mar 14 2022

christopherbate accepted D121666: [mlir][nvvm] Fix bug in ldmatrix intrinsic conversion.
Mar 14 2022, 10:02 PM · Restricted Project, Restricted Project