yaxunl (Yaxun Liu)
User

Projects

User does not belong to any projects.

User Details

User Since
May 13 2015, 10:16 AM (161 w, 6 d)

Recent Activity

Yesterday

yaxunl updated the diff for D48287: [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn.

Correct comments in test.

Mon, Jun 18, 10:45 AM
yaxunl created D48287: [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn.
Mon, Jun 18, 10:07 AM
yaxunl added a reviewer for D48287: [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn: scchan.
Mon, Jun 18, 10:07 AM

Sat, Jun 16

yaxunl committed rC334886: Add RUN line for amdgcn to lit test conditional-temporaries.cpp.
Add RUN line for amdgcn to lit test conditional-temporaries.cpp
Sat, Jun 16, 5:34 AM
yaxunl committed rL334886: Add RUN line for amdgcn to lit test conditional-temporaries.cpp.
Add RUN line for amdgcn to lit test conditional-temporaries.cpp
Sat, Jun 16, 5:33 AM

Fri, Jun 15

yaxunl committed rC334879: Call CreateTempAllocaWithoutCast for ActiveFlag.
Call CreateTempAllocaWithoutCast for ActiveFlag
Fri, Jun 15, 6:26 PM
yaxunl committed rL334879: Call CreateTempAllocaWithoutCast for ActiveFlag.
Call CreateTempAllocaWithoutCast for ActiveFlag
Fri, Jun 15, 6:26 PM
yaxunl committed rC334837: [NFC] Add CreateMemTempWithoutCast and CreateTempAllocaWithoutCast.
[NFC] Add CreateMemTempWithoutCast and CreateTempAllocaWithoutCast
Fri, Jun 15, 8:38 AM
yaxunl committed rL334837: [NFC] Add CreateMemTempWithoutCast and CreateTempAllocaWithoutCast.
[NFC] Add CreateMemTempWithoutCast and CreateTempAllocaWithoutCast
Fri, Jun 15, 8:37 AM

Wed, Jun 13

yaxunl committed rL334625: [AMDGPU] Change enqueue kernel handle type.
[AMDGPU] Change enqueue kernel handle type
Wed, Jun 13, 10:36 AM
yaxunl closed D48094: [AMDGPU] Change enqueue kernel handle type.
Wed, Jun 13, 10:36 AM

Tue, Jun 12

yaxunl committed rC334561: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes.
[CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes
Tue, Jun 12, 5:05 PM
yaxunl committed rL334561: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes.
[CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes
Tue, Jun 12, 5:05 PM
yaxunl closed D47958: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes.
Tue, Jun 12, 5:04 PM
yaxunl edited reviewers for D48094: [AMDGPU] Change enqueue kernel handle type, added: b-sumner; removed: brian.
Tue, Jun 12, 2:05 PM
yaxunl created D48094: [AMDGPU] Change enqueue kernel handle type.
Tue, Jun 12, 2:04 PM

Mon, Jun 11

yaxunl updated the diff for D47958: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes.

Improve error msg.

Mon, Jun 11, 5:53 PM
yaxunl committed rC334457: [CUDA][HIP] Set kernel calling convention before arrange function.
[CUDA][HIP] Set kernel calling convention before arrange function
Mon, Jun 11, 5:21 PM
yaxunl committed rL334457: [CUDA][HIP] Set kernel calling convention before arrange function.
[CUDA][HIP] Set kernel calling convention before arrange function
Mon, Jun 11, 5:21 PM
yaxunl closed D47733: [CUDA][HIP] Set kernel calling convention before arrange function.
Mon, Jun 11, 5:20 PM
yaxunl accepted D47740: [AMDGPU] Do not consider indirect acces through phi for wave limiter.

LGTM. Thanks.

Mon, Jun 11, 8:30 AM

Fri, Jun 8

yaxunl added a comment to D47958: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes.
In D47958#1126875, @tra wrote:

Drive-by review:

The patch could use a better description.
Something that describes *what* the patch does (E.g. enforce that attributes X/Y/Z are only applied to global functions.)
*why* the change is needed is relevant, too, but it's not very useful without the *what* part.

Fri, Jun 8, 2:07 PM
yaxunl retitled D47958: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes from [CUDA][HIP] Allow CUDA kernel to have amdgpu kernel attributes to [CUDA][HIP] Allow CUDA `__global__` functions to have amdgpu kernel attributes.
Fri, Jun 8, 2:06 PM
yaxunl updated the diff for D47733: [CUDA][HIP] Set kernel calling convention before arrange function.

Wrap long RUN lines in test.

Fri, Jun 8, 1:28 PM
yaxunl created D47958: [CUDA][HIP] Allow CUDA __global__ functions to have amdgpu kernel attributes.
Fri, Jun 8, 1:20 PM

Wed, Jun 6

yaxunl committed rL334128: [HIP] Fix unbundling.
[HIP] Fix unbundling
Wed, Jun 6, 12:48 PM
yaxunl committed rC334128: [HIP] Fix unbundling.
[HIP] Fix unbundling
Wed, Jun 6, 12:48 PM
yaxunl closed D47555: [HIP] Fix unbundling.
Wed, Jun 6, 12:48 PM
yaxunl added inline comments to D47555: [HIP] Fix unbundling.
Wed, Jun 6, 12:11 PM
yaxunl updated the diff for D47733: [CUDA][HIP] Set kernel calling convention before arrange function.

Revised by Artem's comments.

Wed, Jun 6, 4:22 AM

Tue, Jun 5

yaxunl added a reviewer for D47555: [HIP] Fix unbundling: jlebar.
Tue, Jun 5, 10:31 AM
yaxunl added reviewers for D47376: [CUDA][HIP] Do not offload for -M: jlebar, rjmccall.
Tue, Jun 5, 10:31 AM
yaxunl committed rL334021: [CUDA][HIP] Do not emit type info when compiling for device.
[CUDA][HIP] Do not emit type info when compiling for device
Tue, Jun 5, 8:15 AM
yaxunl committed rC334021: [CUDA][HIP] Do not emit type info when compiling for device.
[CUDA][HIP] Do not emit type info when compiling for device
Tue, Jun 5, 8:15 AM
yaxunl closed D47694: [CUDA][HIP] Do not emit type info when compiling for device.
Tue, Jun 5, 8:15 AM

Mon, Jun 4

yaxunl updated the diff for D47694: [CUDA][HIP] Do not emit type info when compiling for device.

Revised by John's comments.

Mon, Jun 4, 6:53 PM
yaxunl created D47733: [CUDA][HIP] Set kernel calling convention before arrange function.
Mon, Jun 4, 11:24 AM
yaxunl added a comment to D47694: [CUDA][HIP] Do not emit type info when compiling for device.

Why not just have the driver disable RTTI in the frontend invocation?

CUDA/HIP uses single source for device and host. The host code may depend on RTTI,
e.g., an application may include some boost headers which will fail if RTTI is disabled,
therefore RTTI cannot be disabled when compiling device code.

It's a single source file, yes, but behind the scenes you do a separate invocation of the compiler to re-parse that file for device code-generation, right?

Mon, Jun 4, 9:19 AM

Sun, Jun 3

yaxunl added a comment to D47694: [CUDA][HIP] Do not emit type info when compiling for device.

Why not just have the driver disable RTTI in the frontend invocation?

Sun, Jun 3, 8:15 PM
yaxunl added a comment to D47376: [CUDA][HIP] Do not offload for -M.

ping

Sun, Jun 3, 4:45 PM
yaxunl added a reviewer for D47555: [HIP] Fix unbundling: tra.
Sun, Jun 3, 4:44 PM
yaxunl created D47694: [CUDA][HIP] Do not emit type info when compiling for device.
Sun, Jun 3, 4:41 PM

Wed, May 30

yaxunl created D47555: [HIP] Fix unbundling.
Wed, May 30, 1:51 PM

Tue, May 29

yaxunl committed rL333484: Add HIP toolchain.
Add HIP toolchain
Tue, May 29, 5:58 PM
yaxunl committed rC333484: Add HIP toolchain.
Add HIP toolchain
Tue, May 29, 5:58 PM
yaxunl closed D45212: Add HIP toolchain.
Tue, May 29, 5:57 PM
yaxunl committed rL333483: Add action builder for HIP.
Add action builder for HIP
Tue, May 29, 5:53 PM
yaxunl committed rC333483: Add action builder for HIP.
Add action builder for HIP
Tue, May 29, 5:53 PM
yaxunl closed D46476: [HIP] Add action builder for HIP.
Tue, May 29, 5:53 PM
yaxunl added inline comments to D46476: [HIP] Add action builder for HIP.
Tue, May 29, 12:08 PM

Mon, May 28

yaxunl added a comment to D46476: [HIP] Add action builder for HIP.

Any further changes are needed? Thanks.

Mon, May 28, 6:05 AM

Fri, May 25

yaxunl updated the summary of D47376: [CUDA][HIP] Do not offload for -M.
Fri, May 25, 8:24 AM
yaxunl created D47376: [CUDA][HIP] Do not offload for -M.
Fri, May 25, 8:23 AM

Wed, May 23

yaxunl added inline comments to D45212: Add HIP toolchain.
Wed, May 23, 1:47 PM
yaxunl updated the diff for D45212: Add HIP toolchain.

Revised by Artem's comments.

Wed, May 23, 1:34 PM
yaxunl added inline comments to D45212: Add HIP toolchain.
Wed, May 23, 1:31 PM
yaxunl updated the diff for D45212: Add HIP toolchain.

Revised by Artem's comments.

Wed, May 23, 8:05 AM
yaxunl added a comment to D45212: Add HIP toolchain.
In D45212#1105177, @tra wrote:

Hi,

Sorry about the long silence. I'm back to continue the reviews. I'll handle what I can today and will continue with the rest on Tuesday.

It looks like patch description needs to be updated:

Use clang-offload-bindler to create binary for device ISA.

I don't see anything related to offload-bundler in this patch any more.

Wed, May 23, 7:59 AM

Tue, May 22

yaxunl updated the diff for D46476: [HIP] Add action builder for HIP.

Revised by Artem's comments.

Tue, May 22, 10:38 AM
yaxunl added inline comments to D46476: [HIP] Add action builder for HIP.
Tue, May 22, 10:33 AM
yaxunl added a comment to D47099: Call CreateTempAllocaWithoutCast for ActiveFlag.

I revert it since it caused regression on arm and some other arch's.

Tue, May 22, 9:23 AM
yaxunl committed rL332991: Revert r332982 Call CreateTempMemWithoutCast for ActiveFlag.
Revert r332982 Call CreateTempMemWithoutCast for ActiveFlag
Tue, May 22, 9:17 AM
yaxunl committed rC332991: Revert r332982 Call CreateTempMemWithoutCast for ActiveFlag.
Revert r332982 Call CreateTempMemWithoutCast for ActiveFlag
Tue, May 22, 9:17 AM
yaxunl committed rL332982: Call CreateTempMemWithoutCast for ActiveFlag.
Call CreateTempMemWithoutCast for ActiveFlag
Tue, May 22, 7:40 AM
yaxunl committed rC332982: Call CreateTempMemWithoutCast for ActiveFlag.
Call CreateTempMemWithoutCast for ActiveFlag
Tue, May 22, 7:40 AM
yaxunl closed D47099: Call CreateTempAllocaWithoutCast for ActiveFlag.
Tue, May 22, 7:40 AM

Mon, May 21

yaxunl retitled D47099: Call CreateTempAllocaWithoutCast for ActiveFlag from Disable casting of alloca for ActiveFlag to Call CreateTempAllocaWithoutCast for ActiveFlag.
Mon, May 21, 5:15 PM
yaxunl updated the diff for D47099: Call CreateTempAllocaWithoutCast for ActiveFlag.

Revised by John's comments.

Mon, May 21, 4:54 PM
yaxunl updated the diff for D47099: Call CreateTempAllocaWithoutCast for ActiveFlag.

Add CreateMemTempWithoutCast and CreateTempAllocaWithoutCast by John's comments.

Mon, May 21, 1:44 PM
yaxunl added inline comments to D46472: [HIP] Support offloading by linker script.
Mon, May 21, 7:00 AM

May 19 2018

yaxunl added a comment to D47099: Call CreateTempAllocaWithoutCast for ActiveFlag.

Maybe there should just be a method that makes a primitive alloca without the casting, and then you can call that in CreateTempAlloca.

May 19 2018, 5:20 AM

May 18 2018

yaxunl committed rL332794: Fix evaluator for non-zero alloca addr space.
Fix evaluator for non-zero alloca addr space
May 18 2018, 8:02 PM
yaxunl closed D47081: Fix evaluator for non-zero alloca addr space.
May 18 2018, 8:02 PM
yaxunl created D47099: Call CreateTempAllocaWithoutCast for ActiveFlag.
May 18 2018, 7:30 PM
yaxunl updated the diff for D47081: Fix evaluator for non-zero alloca addr space.

Add positive check.

May 18 2018, 2:20 PM
yaxunl added a reviewer for D45212: Add HIP toolchain: t-tye.
May 18 2018, 1:42 PM
yaxunl created D47081: Fix evaluator for non-zero alloca addr space.
May 18 2018, 11:24 AM
yaxunl updated the summary of D46476: [HIP] Add action builder for HIP.
May 18 2018, 8:35 AM
yaxunl committed rL332724: [HIP] Support offloading by linker script.
[HIP] Support offloading by linker script
May 18 2018, 8:12 AM
yaxunl committed rC332724: [HIP] Support offloading by linker script.
[HIP] Support offloading by linker script
May 18 2018, 8:12 AM
yaxunl closed D46472: [HIP] Support offloading by linker script.
May 18 2018, 8:12 AM
yaxunl added a comment to D45212: Add HIP toolchain.

Hi Artem,

May 18 2018, 7:27 AM
yaxunl added a comment to D46472: [HIP] Support offloading by linker script.

LGTM except for minor suggestions.

May 18 2018, 7:03 AM
yaxunl added inline comments to D46992: [AMDGPU] Add perf hints to functions.
May 18 2018, 4:08 AM

May 17 2018

yaxunl added a comment to D45212: Add HIP toolchain.

ping

May 17 2018, 11:39 AM
yaxunl added a comment to D46476: [HIP] Add action builder for HIP.

ping

May 17 2018, 11:39 AM
yaxunl added a comment to D46472: [HIP] Support offloading by linker script.

ping

May 17 2018, 11:39 AM
yaxunl committed rC332593: CodeGen: Fix invalid bitcast for lifetime.start/end.
CodeGen: Fix invalid bitcast for lifetime.start/end
May 17 2018, 4:20 AM
yaxunl committed rL332593: CodeGen: Fix invalid bitcast for lifetime.start/end.
CodeGen: Fix invalid bitcast for lifetime.start/end
May 17 2018, 4:20 AM
yaxunl closed D45900: CodeGen: Fix invalid bitcast for lifetime.start/end.
May 17 2018, 4:20 AM

May 16 2018

yaxunl removed a dependent revision for D45212: Add HIP toolchain: D46489: [HIP] Let assembler output bitcode for amdgcn.
May 16 2018, 12:40 PM
yaxunl removed a dependency for D46489: [HIP] Let assembler output bitcode for amdgcn: D45212: Add HIP toolchain.
May 16 2018, 12:40 PM
yaxunl abandoned D46489: [HIP] Let assembler output bitcode for amdgcn.

I have updated D46476 to skip backend and assembler phases for amdgcn, therefore this patch is no longer needed.

May 16 2018, 12:39 PM
yaxunl updated the diff for D46476: [HIP] Add action builder for HIP.

Skip backend and assemmbler phases for amdgcn since it does not support linking of object files.

May 16 2018, 12:37 PM
yaxunl added reviewers for D46472: [HIP] Support offloading by linker script: rsmith, t-tye.
May 16 2018, 6:19 AM

May 15 2018

yaxunl updated the diff for D45900: CodeGen: Fix invalid bitcast for lifetime.start/end.

Add optional argument to CreateMemTemp and CreateTempAlloca to get the original alloca and use it for lifetime intrinsic.

May 15 2018, 8:02 PM
yaxunl added a comment to D46489: [HIP] Let assembler output bitcode for amdgcn.

I think the right solution here is to make a CompileJobAction with type TY_LLVM_BC in the first place. You should get the advice of a driver expert, though.

May 15 2018, 12:13 PM
yaxunl added reviewers for D46489: [HIP] Let assembler output bitcode for amdgcn: hfinkel, Hahnfeld.
May 15 2018, 10:01 AM
yaxunl added reviewers for D45212: Add HIP toolchain: hfinkel, Hahnfeld.
May 15 2018, 10:00 AM
yaxunl retitled D46476: [HIP] Add action builder for HIP from [HIP] Add action builder to [HIP] Add action builder for HIP.
May 15 2018, 9:59 AM