yaxunl (Yaxun Liu)
User

Projects

User does not belong to any projects.

User Details

User Since
May 13 2015, 10:16 AM (175 w, 1 d)

Recent Activity

Yesterday

yaxunl created D52320: AMDGPU: add __builtin_amdgcn_update_dpp.
Thu, Sep 20, 12:50 PM
yaxunl added a comment to D51809: [CUDA][HIP] Fix assertion in LookupSpecialMember.

ping

Thu, Sep 20, 9:05 AM

Fri, Sep 14

yaxunl added a comment to D43783: [OpenCL] Remove block invoke function from emitted block literal struct.

Ping! Do you still plan to do this? :)

Fri, Sep 14, 10:31 AM

Fri, Sep 7

yaxunl created D51809: [CUDA][HIP] Fix assertion in LookupSpecialMember.
Fri, Sep 7, 12:36 PM

Wed, Sep 5

yaxunl added a comment to D51411: [OpenCL] Improve diagnostic of argument in address space conversion builtins .

Is this a feature requested by users?

I can understand that this may be useful to pinpoint some conversions which are potentially harmful to performance. However such conversions can often be eliminated by optimization, which makes this warning less useful.

On the other hand, too many warnings is a nuance to users, therefore I am wondering whether this warning should be off by default.

Do you have any chance to have any feedback from users about how useful or intrusive this warning is.

Hi Sam, no feedback from the user. If you think this is not that useful we won't commit or make is off by default as you suggested. Let us know what you think makes more sense.

Wed, Sep 5, 7:08 AM

Thu, Aug 30

yaxunl committed rC341077: [HIP] Add -fvisibility hidden option to clang.
[HIP] Add -fvisibility hidden option to clang
Thu, Aug 30, 8:11 AM
yaxunl committed rL341077: [HIP] Add -fvisibility hidden option to clang.
[HIP] Add -fvisibility hidden option to clang
Thu, Aug 30, 8:11 AM
yaxunl closed D51434: [HIP] Add -fvisibility hidden option to clang.
Thu, Aug 30, 8:11 AM
yaxunl accepted D51480: AMDGPU: Remove remnants of old address space mapping.

LGTM. except the typos due to script. Thanks a lot for this effort!

Thu, Aug 30, 8:01 AM
yaxunl added inline comments to D51434: [HIP] Add -fvisibility hidden option to clang.
Thu, Aug 30, 7:42 AM
yaxunl added inline comments to D51434: [HIP] Add -fvisibility hidden option to clang.
Thu, Aug 30, 7:37 AM

Wed, Aug 29

yaxunl updated the diff for D51434: [HIP] Add -fvisibility hidden option to clang.

Revised by Artem's comments.

Wed, Aug 29, 2:42 PM
yaxunl updated the diff for D51434: [HIP] Add -fvisibility hidden option to clang.

Use -fvisibility hidden.

Wed, Aug 29, 1:59 PM
yaxunl accepted D51209: AMDGPU: Default to hidden visibility.

LGTM. Thanks!

Wed, Aug 29, 1:55 PM
yaxunl committed rL340967: Add predefined macro __gnu_linux__ for proper aux-triple.
Add predefined macro __gnu_linux__ for proper aux-triple
Wed, Aug 29, 1:40 PM
yaxunl committed rC340967: Add predefined macro __gnu_linux__ for proper aux-triple.
Add predefined macro __gnu_linux__ for proper aux-triple
Wed, Aug 29, 1:40 PM
yaxunl closed D51441: Add predefined macro __gnu_linux__ for proper aux-triple.
Wed, Aug 29, 1:40 PM
yaxunl added a comment to D51441: Add predefined macro __gnu_linux__ for proper aux-triple.
In D51441#1218010, @tra wrote:

While we're here, perhaps Builder.defineMacro("__linux__") should be changed to DefineStd("linux") which defines linux/__linux/__linux__?

Will do when committing. Thanks.

Wed, Aug 29, 1:17 PM
yaxunl added a comment to D51441: Add predefined macro __gnu_linux__ for proper aux-triple.
In D51441#1218010, @tra wrote:

While we're here, perhaps Builder.defineMacro("__linux__") should be changed to DefineStd("linux") which defines linux/__linux/__linux__?

Wed, Aug 29, 11:59 AM
yaxunl added a comment to D51434: [HIP] Add -fvisibility hidden option to clang.

D51209 is the patch. I think HIP will need an additional patch, since I think it isn’t subclassing the amdgpu toolchain

Wed, Aug 29, 11:42 AM
yaxunl updated the diff for D51441: Add predefined macro __gnu_linux__ for proper aux-triple.

Revised by Artem's comments.

Wed, Aug 29, 11:11 AM
yaxunl added a reviewer for D51434: [HIP] Add -fvisibility hidden option to clang: scott.linder.
Wed, Aug 29, 10:31 AM
yaxunl added a comment to D51434: [HIP] Add -fvisibility hidden option to clang.
In D51434#1217772, @tra wrote:

Could you elaborate on what exactly is the problem this patch fixes?
I don't see how internalizing the symbols connects to PLTs. My understanding is that PLTs are used to provide stubs for symbols to be resolved by dynamic linker at runtime. AFAICT AMD does not use shared libs on device side. What do I miss?

Wed, Aug 29, 10:12 AM
yaxunl created D51441: Add predefined macro __gnu_linux__ for proper aux-triple.
Wed, Aug 29, 10:06 AM
yaxunl created D51434: [HIP] Add -fvisibility hidden option to clang.
Wed, Aug 29, 8:52 AM
yaxunl added a comment to D43783: [OpenCL] Remove block invoke function from emitted block literal struct.

Sorry for digging up an old commit...

Apparently this broke block arguments, e.g. the following test case:

int foo(int (^ bl)(void)) {
  return bl();
}

int get21() {
  return foo(^{return 21;});
}

int get42() {
  return foo(^{return 42;});
}

In particular, the VarDecl that getBlockExpr() sees doesn't have an initializer when the called block comes from an argument (causing clang to crash).

Sorry for the delay. I think block should not be allowed as function argument since this generally leads indirect function calls therefore requires support of function pointer. It will rely on optimizations to get rid of indirect function calls.

The idea was to allow blocks as function parameters because they are statically known at each function call. This can be resolved later at IR level instead of frontend. I am also not sure there can be other corner cases where it wouldn't work in Clang since we can't leverage analysis passes here. I feel this might be a risky thing to do in Clang. Currently it doesn't work for the examples provided and therefore breaking the compliance with the spec.

Wed, Aug 29, 7:09 AM
yaxunl added inline comments to D51411: [OpenCL] Improve diagnostic of argument in address space conversion builtins .
Wed, Aug 29, 6:58 AM
yaxunl added a comment to D51411: [OpenCL] Improve diagnostic of argument in address space conversion builtins .

Is this a feature requested by users?

Wed, Aug 29, 6:52 AM

Tue, Aug 28

yaxunl committed rL340873: [HIP] Fix output file extension.
[HIP] Fix output file extension
Tue, Aug 28, 2:10 PM
yaxunl committed rC340873: [HIP] Fix output file extension.
[HIP] Fix output file extension
Tue, Aug 28, 2:10 PM
yaxunl closed D51336: [HIP] Fix output file extension.
Tue, Aug 28, 2:10 PM

Mon, Aug 27

yaxunl created D51336: [HIP] Fix output file extension.
Mon, Aug 27, 9:16 PM

Fri, Aug 24

yaxunl added a comment to D43783: [OpenCL] Remove block invoke function from emitted block literal struct.

Sorry for digging up an old commit...

Apparently this broke block arguments, e.g. the following test case:

int foo(int (^ bl)(void)) {
  return bl();
}

int get21() {
  return foo(^{return 21;});
}

int get42() {
  return foo(^{return 42;});
}

In particular, the VarDecl that getBlockExpr() sees doesn't have an initializer when the called block comes from an argument (causing clang to crash).

Fri, Aug 24, 9:17 AM
yaxunl accepted D51212: [OpenCL][Docs] Release notes for OpenCL in Clang .

LGTM. Thanks.

Fri, Aug 24, 9:01 AM

Aug 17 2018

yaxunl accepted D50259: [OpenCL] Disallow negative attribute arguments.

LGTM. Thanks.

Aug 17 2018, 2:58 PM
yaxunl committed rL340056: [HIP] Make __hip_gpubin_handle hidden to avoid being merged across different….
[HIP] Make __hip_gpubin_handle hidden to avoid being merged across different…
Aug 17 2018, 10:48 AM
yaxunl committed rC340056: [HIP] Make __hip_gpubin_handle hidden to avoid being merged across different….
[HIP] Make __hip_gpubin_handle hidden to avoid being merged across different…
Aug 17 2018, 10:48 AM
yaxunl closed D50596: [HIP] Make __hip_gpubin_handle hidden to avoid being merged across different shared libraries.
Aug 17 2018, 10:48 AM

Aug 10 2018

yaxunl created D50596: [HIP] Make __hip_gpubin_handle hidden to avoid being merged across different shared libraries.
Aug 10 2018, 6:55 PM

Aug 2 2018

yaxunl committed rC338805: Sema: Fix explicit address space cast involving void pointers.
Sema: Fix explicit address space cast involving void pointers
Aug 2 2018, 8:19 PM
yaxunl committed rL338805: Sema: Fix explicit address space cast involving void pointers.
Sema: Fix explicit address space cast involving void pointers
Aug 2 2018, 8:19 PM
yaxunl closed D50003: Sema: Fix explicit address space cast involving void pointers.
Aug 2 2018, 8:19 PM
yaxunl accepted D50104: [OpenCL] Always emit alloca in entry block for enqueue_kernel builtin.

LGTM. Thanks!

Aug 2 2018, 6:59 PM
yaxunl added a comment to D50104: [OpenCL] Always emit alloca in entry block for enqueue_kernel builtin.

I still don't quite see what you describe; with that change all of the lifetime.end calls pile up just before the enclosing function returns, not after each call to enqueue_kernel. Looking at https://clang.llvm.org/doxygen/EHScopeStack_8h_source.html#l00078 I don't see any option which isn't based on scope. The lifetime.start calls do occur where I would expect, though, so I will update the patch.

Aug 2 2018, 10:00 AM
yaxunl added a comment to D50104: [OpenCL] Always emit alloca in entry block for enqueue_kernel builtin.

Address feedback; I hope I understood correctly what debug info to check for.

I don't see where in CreateMemTemp and friends EmitLifetimeStart gets called, and I don't see any lifetime intrinsics in the IR even at -O1.

Aug 2 2018, 6:50 AM
yaxunl added inline comments to D50003: Sema: Fix explicit address space cast involving void pointers.
Aug 2 2018, 5:44 AM

Aug 1 2018

yaxunl added inline comments to D50003: Sema: Fix explicit address space cast involving void pointers.
Aug 1 2018, 9:52 PM

Jul 31 2018

yaxunl added inline comments to D50104: [OpenCL] Always emit alloca in entry block for enqueue_kernel builtin.
Jul 31 2018, 2:51 PM

Jul 30 2018

yaxunl accepted D49930: [DebugInfo][OpenCL] Generate correct block literal debug info for OpenCL.

Thank you for taking a look @yaxunl. Should I wait for another reviewer or can I commit this?

Jul 30 2018, 12:39 PM
yaxunl added a reviewer for D49930: [DebugInfo][OpenCL] Generate correct block literal debug info for OpenCL: yaxunl.
Jul 30 2018, 12:37 PM
yaxunl created D50003: Sema: Fix explicit address space cast involving void pointers.
Jul 30 2018, 12:30 PM
yaxunl accepted D47154: Try to make builtin address space declarations not useless.

I missed the addr space casts you added to CodeGenFunction::EmitBuiltinExpr. With those casts it should work.

Jul 30 2018, 4:36 AM

Jul 27 2018

yaxunl added a comment to D49930: [DebugInfo][OpenCL] Generate correct block literal debug info for OpenCL.

LGTM. Thanks.

Jul 27 2018, 8:31 PM
yaxunl accepted D49723: [OpenCL] Check for invalid kernel arguments in array types.

LGTM. Thanks.

Jul 27 2018, 8:28 PM
yaxunl committed rL338188: [CUDA][HIP] Allow function-scope static const variable.
[CUDA][HIP] Allow function-scope static const variable
Jul 27 2018, 8:06 PM
yaxunl committed rC338188: [CUDA][HIP] Allow function-scope static const variable.
[CUDA][HIP] Allow function-scope static const variable
Jul 27 2018, 8:05 PM
yaxunl closed D49931: [CUDA][HIP] Allow function-scope static const variable.
Jul 27 2018, 8:05 PM
yaxunl updated the diff for D49931: [CUDA][HIP] Allow function-scope static const variable.

Revised by Artem's comments.

Jul 27 2018, 2:58 PM
yaxunl added inline comments to D49931: [CUDA][HIP] Allow function-scope static const variable.
Jul 27 2018, 2:38 PM
yaxunl added a comment to D49931: [CUDA][HIP] Allow function-scope static const variable.
In D49931#1178720, @tra wrote:

This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space.

What does NVCC do with local static const variables?

Jul 27 2018, 1:44 PM
yaxunl updated the diff for D49931: [CUDA][HIP] Allow function-scope static const variable.

update diagnostic message.

Jul 27 2018, 12:40 PM
yaxunl created D49931: [CUDA][HIP] Allow function-scope static const variable.
Jul 27 2018, 12:39 PM
yaxunl added a comment to D47154: Try to make builtin address space declarations not useless.
In D47154#1108813, @tra wrote:

CUDA does not expose explicit AS on clang size. All pointers are treated as generic and we infer specific address space only in LLVM.
__nvvm_atom_*_[sg]_* builtins should probably be removed as they are indeed useless without pointers with explicit AS and NVCC itself does not have such builtins either. Instead, we should convert the generic AS builtin to address-space specific instruction somewhere in LLVM.

Using attribute((address_space()) should probably produce an error during CUDA compilation.

Jul 27 2018, 10:26 AM
yaxunl added inline comments to D47154: Try to make builtin address space declarations not useless.
Jul 27 2018, 9:59 AM
yaxunl added inline comments to D47154: Try to make builtin address space declarations not useless.
Jul 27 2018, 9:43 AM
yaxunl requested changes to D47154: Try to make builtin address space declarations not useless.
Jul 27 2018, 9:18 AM

Jul 26 2018

yaxunl accepted D49725: [OpenCL] Forbid size dependent types used as kernel arguments.

LGTM. Thanks!

Jul 26 2018, 8:17 AM

Jul 24 2018

yaxunl added a comment to D49725: [OpenCL] Forbid size dependent types used as kernel arguments.

This patch also adds check for array of structs. Can you include this in title or split to a separate patch?

Jul 24 2018, 6:02 AM
yaxunl accepted D49723: [OpenCL] Check for invalid kernel arguments in array types.

LGTM

Jul 24 2018, 5:41 AM

Jul 23 2018

yaxunl committed rL337797: Attempt to fix regression due to r337791.
Attempt to fix regression due to r337791
Jul 23 2018, 7:12 PM
yaxunl committed rC337797: Attempt to fix regression due to r337791.
Attempt to fix regression due to r337791
Jul 23 2018, 7:12 PM
yaxunl committed rL337793: [HIP] pass -target-cpu when running the device-mode compiler.
[HIP] pass -target-cpu when running the device-mode compiler
Jul 23 2018, 6:41 PM
yaxunl committed rC337793: [HIP] pass -target-cpu when running the device-mode compiler.
[HIP] pass -target-cpu when running the device-mode compiler
Jul 23 2018, 6:41 PM
yaxunl closed D49643: [HIP] pass `-target-cpu` when running the device-mode compiler.
Jul 23 2018, 6:41 PM
yaxunl committed rC337791: Enable .hip files for test/Driver.
Enable .hip files for test/Driver
Jul 23 2018, 6:04 PM
yaxunl committed rL337791: Enable .hip files for test/Driver.
Enable .hip files for test/Driver
Jul 23 2018, 6:04 PM
yaxunl added inline comments to D47618: __c11_atomic_load's _Atomic can be const.
Jul 23 2018, 8:27 AM

Jul 22 2018

yaxunl retitled D49643: [HIP] pass `-target-cpu` when running the device-mode compiler from [HIP] Add -target-cpu option for clang -cc1 to [HIP] pass `-target-cpu` when running the device-mode compiler.
Jul 22 2018, 7:28 PM
yaxunl added a comment to D49643: [HIP] pass `-target-cpu` when running the device-mode compiler.

The commit message here could be better. You're passing -target-cpu when running the device-mode compiler.

Jul 22 2018, 7:27 PM

Jul 21 2018

yaxunl created D49643: [HIP] pass `-target-cpu` when running the device-mode compiler.
Jul 21 2018, 9:39 PM

Jul 20 2018

yaxunl committed rL337639: [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn.
[HIP] Support -fcuda-flush-denormals-to-zero for amdgcn
Jul 20 2018, 7:07 PM
yaxunl committed rC337639: [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn.
[HIP] Support -fcuda-flush-denormals-to-zero for amdgcn
Jul 20 2018, 7:07 PM
yaxunl closed D48287: [HIP] Support -fcuda-flush-denormals-to-zero for amdgcn.
Jul 20 2018, 7:07 PM
yaxunl committed rL337631: [HIP] Register/unregister device fat binary only once.
[HIP] Register/unregister device fat binary only once
Jul 20 2018, 3:51 PM
yaxunl committed rC337631: [HIP] Register/unregister device fat binary only once.
[HIP] Register/unregister device fat binary only once
Jul 20 2018, 3:50 PM
yaxunl closed D49083: [HIP] Register/unregister device fat binary only once.
Jul 20 2018, 3:50 PM
yaxunl committed rC337540: Sema: Fix explicit address space cast in C++.
Sema: Fix explicit address space cast in C++
Jul 20 2018, 4:38 AM
yaxunl committed rL337540: Sema: Fix explicit address space cast in C++.
Sema: Fix explicit address space cast in C++
Jul 20 2018, 4:38 AM
yaxunl closed D49294: Sema: Fix explicit address space cast in C++.
Jul 20 2018, 4:38 AM

Jul 19 2018

yaxunl updated the summary of D49083: [HIP] Register/unregister device fat binary only once.
Jul 19 2018, 8:38 PM
yaxunl updated the diff for D49083: [HIP] Register/unregister device fat binary only once.

Revised by John's comments.

Jul 19 2018, 8:37 PM
yaxunl added inline comments to D49083: [HIP] Register/unregister device fat binary only once.
Jul 19 2018, 7:55 PM
yaxunl updated the diff for D49294: Sema: Fix explicit address space cast in C++.

Revised by John's comments.

Jul 19 2018, 5:12 PM
yaxunl updated the diff for D49083: [HIP] Register/unregister device fat binary only once.

Revised by John's comments.

Jul 19 2018, 4:28 PM
yaxunl added inline comments to D49083: [HIP] Register/unregister device fat binary only once.
Jul 19 2018, 4:16 PM
yaxunl added inline comments to D49294: Sema: Fix explicit address space cast in C++.
Jul 19 2018, 11:43 AM
yaxunl added inline comments to D49083: [HIP] Register/unregister device fat binary only once.
Jul 19 2018, 7:28 AM

Jul 18 2018

yaxunl updated the diff for D49083: [HIP] Register/unregister device fat binary only once.

Added comments about thread safety of ctor functions.

Jul 18 2018, 4:57 PM
yaxunl added inline comments to D49083: [HIP] Register/unregister device fat binary only once.
Jul 18 2018, 11:01 AM
yaxunl added inline comments to D49294: Sema: Fix explicit address space cast in C++.
Jul 18 2018, 9:07 AM