Page MenuHomePhabricator

hliao (Michael Liao)
User

Projects

User does not belong to any projects.

User Details

User Since
Aug 7 2014, 12:01 PM (284 w, 1 d)

Recent Activity

Yesterday

hliao added a comment to D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..

yeah, here one example

Fri, Jan 17, 11:02 PM · Restricted Project
hliao added a comment to D72941: Handle ptrtoint in InferAddressSpace.

inttoptr and ptrtoint should be treated as opaque ones. If needed, they should be handled with target-specific. This pass is definitely not the place to handle them.

Fri, Jan 17, 10:53 PM · Restricted Project
hliao committed rG6d0d86a64d30: [DAG] Add helper for creating constant vector index with correct type. NFC. (authored by hliao).
[DAG] Add helper for creating constant vector index with correct type. NFC.
Fri, Jan 17, 10:24 PM

Thu, Jan 16

hliao committed rG40514a7d7a3b: [clangd] Add workaround for GCC5 host compilers. NFC. (authored by hliao).
[clangd] Add workaround for GCC5 host compilers. NFC.
Thu, Jan 16, 1:07 PM

Tue, Jan 14

hliao committed rG65c8abb14e77: [amdgpu] Fix typos in a test case. (authored by hliao).
[amdgpu] Fix typos in a test case.
Tue, Jan 14, 5:13 PM
hliao committed rG01a4b8315476: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU. (authored by hliao).
[codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU.
Tue, Jan 14, 4:33 PM
hliao closed D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..
Tue, Jan 14, 4:33 PM · Restricted Project
hliao committed rG8d07f8d98c48: [DAGCombine] Replace `getIntPtrConstant()` with `getVectorIdxTy()`. (authored by hliao).
[DAGCombine] Replace `getIntPtrConstant()` with `getVectorIdxTy()`.
Tue, Jan 14, 2:08 PM
hliao added a comment to D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..

revise after review comments

Tue, Jan 14, 2:08 PM · Restricted Project
hliao updated the diff for D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..

revise after review comments.

Tue, Jan 14, 2:08 PM · Restricted Project
hliao committed rGa3490e3e3d38: Remove trailing `;`. NFC. (authored by hliao).
Remove trailing `;`. NFC.
Tue, Jan 14, 1:58 PM
hliao added a comment to D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..

You have skipped the dead MO, but was pass reordering really necessary? It seems we have higher register pressure with this change.

Tue, Jan 14, 1:19 PM · Restricted Project
hliao added a comment to D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..

I'm not sure this logic works in the presence of an undef subregister def

Tue, Jan 14, 1:19 PM · Restricted Project
hliao added inline comments to D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..
Tue, Jan 14, 1:19 PM · Restricted Project
hliao added inline comments to D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..
Tue, Jan 14, 8:12 AM · Restricted Project
hliao created D72709: [codegen,amdgpu] Enhance MIR DIE and re-arrange it for AMDGPU..
Tue, Jan 14, 8:12 AM · Restricted Project

Thu, Jan 9

hliao added inline comments to D60386: FileCheck [6/12]: Introduce numeric variable definition.
Thu, Jan 9, 9:27 AM · Restricted Project
hliao added inline comments to D60386: FileCheck [6/12]: Introduce numeric variable definition.
Thu, Jan 9, 9:27 AM · Restricted Project
hliao added inline comments to D60386: FileCheck [6/12]: Introduce numeric variable definition.
Thu, Jan 9, 6:27 AM · Restricted Project

Wed, Jan 8

hliao committed rG07a569a0539a: [amdgpu] Remove unused header. NFC. (authored by hliao).
[amdgpu] Remove unused header. NFC.
Wed, Jan 8, 8:37 AM

Fri, Jan 3

hliao committed rG3566c75ca82f: [amdgpu] Skip non-instruction values in CF user tracing. (authored by hliao).
[amdgpu] Skip non-instruction values in CF user tracing.
Fri, Jan 3, 1:03 PM
hliao closed D72174: [amdgpu] Skip non-instruction values in CF user tracing..
Fri, Jan 3, 1:03 PM · Restricted Project
hliao added a comment to D72174: [amdgpu] Skip non-instruction values in CF user tracing..

Where is this being called from that it's so expensive?

Fri, Jan 3, 1:01 PM · Restricted Project
hliao added a comment to D72174: [amdgpu] Skip non-instruction values in CF user tracing..

one compilation is reduced from 460s to 121s after this trivial change.

Fri, Jan 3, 12:21 PM · Restricted Project
hliao created D72174: [amdgpu] Skip non-instruction values in CF user tracing..
Fri, Jan 3, 12:11 PM · Restricted Project

Thu, Jan 2

hliao added a comment to D69694: Memory leak fix for Managed Static Mutex.

Could you double-check that on other platforms? Like Linux with shared library builds. At least, there are crash in my local build.

Thu, Jan 2, 1:18 PM · Restricted Project
hliao updated the diff for D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..

code refinement after reviewers' comments.

Thu, Jan 2, 12:12 PM · Restricted Project
hliao added a comment to D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..

refinements are made after comments from reviewers.

Thu, Jan 2, 12:12 PM · Restricted Project

Tue, Dec 31

hliao committed rG79d401905fcf: [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`. (authored by hliao).
[amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`.
Tue, Dec 31, 11:27 AM
hliao closed D71866: [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`..
Tue, Dec 31, 11:27 AM · Restricted Project
hliao added a comment to D71866: [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`..

Add clarification on the test case.

Tue, Dec 31, 12:08 AM · Restricted Project
hliao retitled D71866: [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`. from [amdgpu] Fix scoreboard updating after `s_waitcnt_vscnt`. to [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`..
Tue, Dec 31, 12:07 AM · Restricted Project
hliao updated the diff for D71866: [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`..

Add test case.

Tue, Dec 31, 12:07 AM · Restricted Project

Tue, Dec 24

hliao created D71866: [amdgpu] Fix scoreboard updating on `s_waitcnt_vscnt`..
Tue, Dec 24, 9:49 PM · Restricted Project

Sat, Dec 21

hliao committed rG7cee28858674: Fix `-Wunused-variable` warning. NFC. (authored by hliao).
Fix `-Wunused-variable` warning. NFC.
Sat, Dec 21, 8:13 AM

Dec 12 2019

hliao committed rG11b2b2f4b100: [amdgpu] Fix `-Wenum-compare` warning. NFC. (authored by hliao).
[amdgpu] Fix `-Wenum-compare` warning. NFC.
Dec 12 2019, 8:51 AM

Dec 10 2019

hliao added a comment to D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..
In D71227#1778136, @tra wrote:

I wonder if this patch will help with this case:

https://godbolt.org/z/X4KdsV

__device__ float fn(int) { return threadIdx.x; };
__host__ float fn(float);

float gvar1 = []()__device__ { return fn(1);} (); // This ends up calling fn(int) on *host*

We seem to happily let host code call device function from a lambda function used as an initializer.

Dec 10 2019, 7:31 PM · Restricted Project
hliao added a comment to D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..

File PR44266 to track that bug.

Dec 10 2019, 11:32 AM · Restricted Project

Dec 9 2019

hliao updated the summary of D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..
Dec 9 2019, 1:42 PM · Restricted Project
hliao updated the diff for D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..

refine again

Dec 9 2019, 1:42 PM · Restricted Project
hliao updated the summary of D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..
Dec 9 2019, 1:40 PM · Restricted Project
hliao updated the diff for D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..

refine commit message

Dec 9 2019, 1:40 PM · Restricted Project
hliao created D71227: [cuda][hip] Fix function overload resolution in the global initiailizer..
Dec 9 2019, 1:40 PM · Restricted Project
hliao committed rG6626e5a06a99: Fix compilation warning from GCC7. NFC. (authored by hliao).
Fix compilation warning from GCC7. NFC.
Dec 9 2019, 7:19 AM

Dec 6 2019

hliao committed rGf2ace9d6005b: Add `QualType::hasAddressSpace`. NFC. (authored by hliao).
Add `QualType::hasAddressSpace`. NFC.
Dec 6 2019, 10:12 AM

Dec 4 2019

hliao committed rGfa9dd410a9a9: [opencl] Fix address space deduction on array variables. (authored by hliao).
[opencl] Fix address space deduction on array variables.
Dec 4 2019, 6:44 AM
hliao closed D70981: [opencl] Fix address space deduction on array variables..
Dec 4 2019, 6:44 AM · Restricted Project

Dec 3 2019

hliao committed rG59312cb0b81c: Fix warning on unused variable. NFC. (authored by hliao).
Fix warning on unused variable. NFC.
Dec 3 2019, 6:18 PM
hliao created D70981: [opencl] Fix address space deduction on array variables..
Dec 3 2019, 1:21 PM · Restricted Project
hliao committed rG59e69fefab88: Fix warning on extra ';'. NFC. (authored by hliao).
Fix warning on extra ';'. NFC.
Dec 3 2019, 1:12 PM
hliao committed rG3953540d8bb5: Remove unused variable. NFC. (authored by hliao).
Remove unused variable. NFC.
Dec 3 2019, 12:16 PM

Nov 26 2019

hliao accepted D70729: Workaround for EvalInfo ctor for MSVC 2017.

LGTM, anyway calling a virtual function inside the constructor is not encouraged.

Nov 26 2019, 11:08 AM · Restricted Project

Nov 21 2019

hliao committed rGc4afc6566a64: Fix compilation warning. NFC. (authored by hliao).
Fix compilation warning. NFC.
Nov 21 2019, 9:10 AM

Nov 20 2019

hliao committed rG4a308d302c33: [AMDGPU] Keep consistent check of legal addressing mode. (authored by hliao).
[AMDGPU] Keep consistent check of legal addressing mode.
Nov 20 2019, 12:11 PM
hliao closed D70473: [AMDGPU] Keep consistent check of legal addressing mode..
Nov 20 2019, 12:11 PM · Restricted Project
hliao added inline comments to D70473: [AMDGPU] Keep consistent check of legal addressing mode..
Nov 20 2019, 10:48 AM · Restricted Project

Nov 19 2019

hliao updated the diff for D70473: [AMDGPU] Keep consistent check of legal addressing mode..

revise a test

Nov 19 2019, 4:04 PM · Restricted Project
hliao created D70473: [AMDGPU] Keep consistent check of legal addressing mode..
Nov 19 2019, 4:04 PM · Restricted Project

Nov 18 2019

hliao committed rG17e37ba57a69: Fix shared lib build. (authored by hliao).
Fix shared lib build.
Nov 18 2019, 10:14 AM

Nov 13 2019

hliao committed rG2bf9b9a5a3a4: [TTI] Fix cast cost on vector types. (authored by hliao).
[TTI] Fix cast cost on vector types.
Nov 13 2019, 11:05 AM

Nov 12 2019

hliao updated the diff for D61458: [hip] Relax CUDA call restriction within `decltype` context..

This patch is revived with more changes addressing the previous concerns.

Nov 12 2019, 11:32 AM · Restricted Project
hliao committed rGceb72d07b004: Fix build with shared libraries. NFC. (authored by hliao).
Fix build with shared libraries. NFC.
Nov 12 2019, 10:46 AM

Nov 8 2019

hliao added a comment to D69322: [hip][cuda] Enable extended lambda support on Windows..

PING for review

Nov 8 2019, 9:03 AM · Restricted Project

Nov 6 2019

hliao added a comment to D69322: [hip][cuda] Enable extended lambda support on Windows..

PING for review

Nov 6 2019, 5:58 AM · Restricted Project

Nov 5 2019

hliao committed rG0a220de9e9ca: [HIP] Fix visibility for 'extern' device variables. (authored by hliao).
[HIP] Fix visibility for 'extern' device variables.
Nov 5 2019, 11:25 AM
hliao closed D63020: [HIP] Fix visibility for 'extern' device variables..
Nov 5 2019, 11:25 AM · Restricted Project
hliao added a comment to D63020: [HIP] Fix visibility for 'extern' device variables..

Sam, could you review this? Even though it has no functionality issue so far, from the code sequence, once there's an addrspacecast is inserted, we lose the chance to set target specific attributes if any.

Nov 5 2019, 10:47 AM · Restricted Project
hliao added a comment to D69826: [hip] Enable pointer argument lowering through coercing type..

I am a little bit concerned that user may have such code:

struct A { int *p; }
__global__ kernel(A a) {
  int x;
  a.p = &x;
  f(a);
}

@arsenm what happens if a private pointer is mis-used as a global pointer?

I am wondering if we should coerce byval struct kernel arg to global only if they are const, e.g.

__global__ kernel(const A a);

I understand this may lose performance. Or should we introduce an option to let user disable coerce of non-const struct kernel arg to global.

Nov 5 2019, 10:38 AM · Restricted Project
hliao committed rG15140e4bacf9: [hip] Enable pointer argument lowering through coercing type. (authored by hliao).
[hip] Enable pointer argument lowering through coercing type.
Nov 5 2019, 10:11 AM
hliao closed D69826: [hip] Enable pointer argument lowering through coercing type..
Nov 5 2019, 10:10 AM · Restricted Project
hliao abandoned D69679: [AMDGPU] Add amdgpu-promote-pointer-kernargs pass.

A different approach (D69826) is taken to address the same issue.

Nov 5 2019, 10:10 AM · Restricted Project
hliao updated the diff for D69826: [hip] Enable pointer argument lowering through coercing type..

Add host-side checks.

Nov 5 2019, 10:10 AM · Restricted Project
hliao added inline comments to D69826: [hip] Enable pointer argument lowering through coercing type..
Nov 5 2019, 7:05 AM · Restricted Project
hliao updated the diff for D69826: [hip] Enable pointer argument lowering through coercing type..

revise parameter names

Nov 5 2019, 7:05 AM · Restricted Project

Nov 4 2019

hliao added inline comments to D69826: [hip] Enable pointer argument lowering through coercing type..
Nov 4 2019, 3:57 PM · Restricted Project
hliao updated the diff for D69826: [hip] Enable pointer argument lowering through coercing type..
  • revise member function name.
  • add the test case for by-val array types.
Nov 4 2019, 3:52 PM · Restricted Project
hliao updated the diff for D69826: [hip] Enable pointer argument lowering through coercing type..

revise code following reviwers' comments.

Nov 4 2019, 2:56 PM · Restricted Project
hliao added inline comments to D69826: [hip] Enable pointer argument lowering through coercing type..
Nov 4 2019, 2:38 PM · Restricted Project
hliao updated the diff for D69826: [hip] Enable pointer argument lowering through coercing type..

add the test case for struct.

Nov 4 2019, 2:35 PM · Restricted Project
hliao created D69826: [hip] Enable pointer argument lowering through coercing type..
Nov 4 2019, 2:08 PM · Restricted Project
hliao added a comment to D69826: [hip] Enable pointer argument lowering through coercing type..

It happens that Sam has a similar patch of this one. After discussion, we agreed that this patch addresses more cases found in the workloads. Thank Sam for the test case.

Nov 4 2019, 2:08 PM · Restricted Project
hliao added a comment to D69322: [hip][cuda] Enable extended lambda support on Windows..

Looks like holidays are approaching, :). Anyway, it's really appreciated that you could review this patch to enable CUDA/HIP applications on Windows.

Nov 4 2019, 2:01 PM · Restricted Project
hliao committed rGd142ec6fef9a: Fix compilation warning. NFC. (authored by hliao).
Fix compilation warning. NFC.
Nov 4 2019, 7:04 AM

Nov 1 2019

hliao committed rG4531aee2ac16: [amdgpu] Fix known bits compuation on `MUL_I24`/`MUL_U24`. (authored by hliao).
[amdgpu] Fix known bits compuation on `MUL_I24`/`MUL_U24`.
Nov 1 2019, 2:15 PM
hliao closed D69735: [amdgpu] Fix known bits compuation on `MUL_I24`/`MUL_U24`..
Nov 1 2019, 2:15 PM · Restricted Project
hliao added a comment to D69735: [amdgpu] Fix known bits compuation on `MUL_I24`/`MUL_U24`..

What was the symptom of this?

Nov 1 2019, 2:04 PM · Restricted Project
hliao created D69735: [amdgpu] Fix known bits compuation on `MUL_I24`/`MUL_U24`..
Nov 1 2019, 1:24 PM · Restricted Project
hliao resigned from D64015: [WIP][CUDA] Use shared MangleContext for CUDA and CXX CG.
Nov 1 2019, 11:16 AM · Restricted Project
hliao added a comment to D69694: Memory leak fix for Managed Static Mutex.

could you keep the original scoped lock by adding a level of block?

Nov 1 2019, 11:16 AM · Restricted Project
hliao added a comment to D69322: [hip][cuda] Enable extended lambda support on Windows..

kindly PING for review

Nov 1 2019, 10:00 AM · Restricted Project

Oct 31 2019

hliao created D69679: [AMDGPU] Add amdgpu-promote-pointer-kernargs pass.
Oct 31 2019, 1:25 PM · Restricted Project

Oct 30 2019

hliao added inline comments to D69322: [hip][cuda] Enable extended lambda support on Windows..
Oct 30 2019, 7:55 PM · Restricted Project
hliao updated the diff for D69322: [hip][cuda] Enable extended lambda support on Windows..

simplify again following suggestion.

Oct 30 2019, 7:53 PM · Restricted Project
hliao added inline comments to D69322: [hip][cuda] Enable extended lambda support on Windows..
Oct 30 2019, 7:19 AM · Restricted Project
hliao updated the diff for D69322: [hip][cuda] Enable extended lambda support on Windows..

revise MSHIPNumberingContext

Oct 30 2019, 7:19 AM · Restricted Project

Oct 29 2019

hliao added a comment to D69322: [hip][cuda] Enable extended lambda support on Windows..

PING for review

Oct 29 2019, 5:48 AM · Restricted Project

Oct 24 2019

hliao committed rG45787e56829f: Fix compilation warning. NFC. (authored by hliao).
Fix compilation warning. NFC.
Oct 24 2019, 10:09 PM
hliao committed rG5a48678a6a16: [hip] Allow the declaration of functions with variadic arguments in HIP. (authored by hliao).
[hip] Allow the declaration of functions with variadic arguments in HIP.
Oct 24 2019, 9:40 PM
hliao closed D69389: [hip] Allow the declaration of functions with variadic arguments in HIP..
Oct 24 2019, 9:40 PM · Restricted Project
hliao added inline comments to D69389: [hip] Allow the declaration of functions with variadic arguments in HIP..
Oct 24 2019, 12:48 PM · Restricted Project