Page MenuHomePhabricator

yaxunl (Yaxun Liu)
User

Projects

User does not belong to any projects.

User Details

User Since
May 13 2015, 10:16 AM (244 w, 2 d)

Recent Activity

Thu, Jan 16

yaxunl abandoned D72868: Workaround build failure with gcc5.4.

it is already there

Thu, Jan 16, 1:36 PM
yaxunl created D72868: Workaround build failure with gcc5.4.
Thu, Jan 16, 12:56 PM
yaxunl added a comment to D72806: [HIP] fix paths for executables not in clang bin directory.

What's the use case of this change?

Thu, Jan 16, 10:22 AM · Restricted Project
yaxunl added a reviewer for D72806: [HIP] fix paths for executables not in clang bin directory: tra.
Thu, Jan 16, 10:22 AM · Restricted Project

Wed, Jan 15

yaxunl accepted D72723: Built-in functions for AMDGPU MFMA instructions..

LGTM. Thanks.

Wed, Jan 15, 7:54 AM · Restricted Project, Restricted Project, Restricted Project

Thu, Jan 9

yaxunl added a comment to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

Most uses of the destructor do not use the delete operator, though, and therefore should not trigger the diagnostics in f to be emitted. And this really doesn't require a fully-realized use graph; you could very easily track the current use stack when making a later pass over the entities used.

Thu, Jan 9, 11:20 AM
yaxunl added a comment to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

This doesn't look quite right to me. I don't think we should treat the delete this; for a destructor as being emitted-for-device in any translation unit in which the vtable is marked used. (For example, if in your testcase MSEmitDeletingDtor::CFileStream::CFileStream() were a __host__ function, I think you'd still diagnose, but presumably shouldn't do so, because the vtable -- and therefore CFileStream::operator delete -- is never referenced / emitted for the device.) Instead, I think we should treat the delete this; as being emitted in any translation unit in which the vtable itself is emitted-for-device. Presumably, this means you will need to model / track usage of the vtable itself in your "call graph".

Thu, Jan 9, 10:52 AM
yaxunl updated the diff for D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

Add tests for device compilation.

Thu, Jan 9, 10:23 AM
yaxunl added a comment to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

I thought you were saying that the destructor decl hadn't been created yet, but I see now that you're saying something more subtle.

CurContext is set to the destructor because the standard says in [class.dtor]p13:

At the point of definition of a virtual destructor (including an implicit definition), the non-array deallocation function is determined as if for the expression `delete this` appearing in a non-virtual destructor of the destructor’s class.

Which is to say that, semantically, the context is as if it were within the destructor, to the extent that this affects access control and so on.

I can see why this causes problems for your call graph (really a use graph), since it's a use in the apparent context of the destructor at a point where the destructor is not being defined. A similar thing happens with default arguments, but because we don't consider uses from default arguments to be true ODR-uses until the default argument is used, that probably doesn't cause problems for you.

I don't think the destructor -> deallocation function edge is actually interesting for your use graph. It'd be more appropriate to treat the deallocation function as used by the v-table than by the destructor; I don't know whether you make any attempt to model v-tables as nodes in your use graph. You might consider finding a simple way to suppress adding this edge, like just not adding edges from a destructor that's not currently being defined (D->willHaveBody()).

With all that said, maintaining a use graph for all the functions you might emit in the entire translation unit seems very expensive and brittle. Have you considered doing this walk in a final pass? You could just build up a set of all the functions you know you're going to emit and then walk their bodies looking for uses of lazy-emitted entities. If we don't already have a function that calls a callback for every declaration ODR-used by a function body, we should.

Thu, Jan 9, 10:23 AM

Tue, Jan 7

yaxunl added a comment to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

Richard is definitely our main expert in the implicit synthesis of special members. It seems to me that if we need the destructor declaration at some point, we should be forcing it to exist at that point.

Tue, Jan 7, 1:56 PM
yaxunl added a comment to D68578: [HIP] Fix device stub name.

ping

Tue, Jan 7, 8:49 AM
yaxunl committed rG9f2d8b5c0cdb: [HIP] Add option --gpu-max-threads-per-block=n (authored by yaxunl).
[HIP] Add option --gpu-max-threads-per-block=n
Tue, Jan 7, 8:21 AM
yaxunl closed D71221: [HIP] Add option --gpu-max-threads-per-block=n.
Tue, Jan 7, 8:21 AM · Restricted Project

Sun, Jan 5

yaxunl added a comment to D71221: [HIP] Add option --gpu-max-threads-per-block=n.

ping

Sun, Jan 5, 9:53 AM · Restricted Project

Fri, Dec 27

yaxunl committed rG134ef0fb4b92: [OpenCL] Fix inconsistency between opencl and c11 atomic fetch max/min (authored by yaxunl).
[OpenCL] Fix inconsistency between opencl and c11 atomic fetch max/min
Fri, Dec 27, 8:51 AM
yaxunl closed D71725: [OpenCL] Fix inconsistency between opencl and c11 atomic fetch max/min.
Fri, Dec 27, 8:51 AM · Restricted Project

Mon, Dec 23

yaxunl committed rGeca40066ebb5: [NFC] Move OptionUtils from Basic to Driver (authored by yaxunl).
[NFC] Move OptionUtils from Basic to Driver
Mon, Dec 23, 5:53 AM
yaxunl closed D71802: [NFC] Move OptionUtils from Basic to Driver.
Mon, Dec 23, 5:53 AM · Restricted Project

Sat, Dec 21

yaxunl created D71802: [NFC] Move OptionUtils from Basic to Driver.
Sat, Dec 21, 9:37 PM · Restricted Project
yaxunl added a comment to D71080: [NFC] Separate getLastArgIntValue to Basic.

Why move this to Basic instead of to Driver if you want to use it in Driver? (Frontend depends on Driver so it could then still use it.)

Sat, Dec 21, 9:07 PM · Restricted Project
yaxunl committed rG7376d9eb3891: [NFC] Separate getLastArgIntValue to Basic (authored by yaxunl).
[NFC] Separate getLastArgIntValue to Basic
Sat, Dec 21, 5:40 PM
yaxunl closed D71080: [NFC] Separate getLastArgIntValue to Basic.
Sat, Dec 21, 5:40 PM · Restricted Project
yaxunl updated the diff for D71221: [HIP] Add option --gpu-max-threads-per-block=n.

revised by Artem's comments.

Sat, Dec 21, 12:27 PM · Restricted Project

Fri, Dec 20

yaxunl added a comment to D71726: Let clang atomic builtins fetch add/sub support floating point types.
In D71726#1791904, @jfb wrote:

This generally seems fine. Does it work on most backends? I want to make sure it doesn't fail in backends :)

Fri, Dec 20, 7:38 AM
yaxunl added a comment to D71221: [HIP] Add option --gpu-max-threads-per-block=n.
In D71221#1791802, @tra wrote:

What's the use case for this flag?

Fri, Dec 20, 5:58 AM · Restricted Project

Thu, Dec 19

yaxunl added a comment to D71080: [NFC] Separate getLastArgIntValue to Basic.

ping

Thu, Dec 19, 1:44 PM · Restricted Project
yaxunl added a comment to D71221: [HIP] Add option --gpu-max-threads-per-block=n.

ping

Thu, Dec 19, 1:44 PM · Restricted Project
yaxunl created D71726: Let clang atomic builtins fetch add/sub support floating point types.
Thu, Dec 19, 1:14 PM
yaxunl created D71725: [OpenCL] Fix inconsistency between opencl and c11 atomic fetch max/min.
Thu, Dec 19, 1:13 PM · Restricted Project

Dec 10 2019

yaxunl committed rG21b43885b81a: Fix bug 44190 - wrong code with #pragma pack(1) (authored by yaxunl).
Fix bug 44190 - wrong code with #pragma pack(1)
Dec 10 2019, 11:05 AM
yaxunl closed D71282: Fix bug 44190 - wrong code with #pragma pack(1).
Dec 10 2019, 11:05 AM · Restricted Project
yaxunl updated the diff for D71282: Fix bug 44190 - wrong code with #pragma pack(1).

remove parenthesis

Dec 10 2019, 10:09 AM · Restricted Project
yaxunl added inline comments to D71282: Fix bug 44190 - wrong code with #pragma pack(1).
Dec 10 2019, 10:09 AM · Restricted Project
yaxunl created D71282: Fix bug 44190 - wrong code with #pragma pack(1).
Dec 10 2019, 9:13 AM · Restricted Project

Dec 9 2019

yaxunl created D71221: [HIP] Add option --gpu-max-threads-per-block=n.
Dec 9 2019, 12:07 PM · Restricted Project

Dec 6 2019

yaxunl updated the diff for D71080: [NFC] Separate getLastArgIntValue to Basic.

revised by Artem's comments

Dec 6 2019, 1:29 PM · Restricted Project
yaxunl added inline comments to D71080: [NFC] Separate getLastArgIntValue to Basic.
Dec 6 2019, 1:20 PM · Restricted Project
yaxunl updated the diff for D71080: [NFC] Separate getLastArgIntValue to Basic.

revised by Artem's comments.

Dec 6 2019, 9:44 AM · Restricted Project
yaxunl added inline comments to D71080: [NFC] Separate getLastArgIntValue to Basic.
Dec 6 2019, 9:44 AM · Restricted Project

Dec 5 2019

yaxunl added a parent revision for D69582: Let clang driver support parallel jobs: D71080: [NFC] Separate getLastArgIntValue to Basic.
Dec 5 2019, 11:11 AM
yaxunl added a child revision for D71080: [NFC] Separate getLastArgIntValue to Basic: D69582: Let clang driver support parallel jobs.
Dec 5 2019, 11:11 AM · Restricted Project
yaxunl updated the diff for D69582: Let clang driver support parallel jobs.

split the patch

Dec 5 2019, 11:11 AM
yaxunl created D71080: [NFC] Separate getLastArgIntValue to Basic.
Dec 5 2019, 11:01 AM · Restricted Project
yaxunl accepted D70987: [AMDGPU][HIP] Improve opt-level handling.
Dec 5 2019, 7:43 AM · Restricted Project

Dec 4 2019

yaxunl added inline comments to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.
Dec 4 2019, 1:41 PM
yaxunl updated the diff for D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

remove unnecessary states added to Sema.

Dec 4 2019, 1:34 PM
yaxunl committed rG7d0e1117c929: [HIP] Remove opencl.amdgcn.lib (authored by yaxunl).
[HIP] Remove opencl.amdgcn.lib
Dec 4 2019, 9:42 AM
yaxunl closed D70980: [HIP] Remove opencl.amdgcn.lib.
Dec 4 2019, 9:42 AM · Restricted Project
yaxunl accepted D70987: [AMDGPU][HIP] Improve opt-level handling.

LGTM

Dec 4 2019, 8:17 AM · Restricted Project

Dec 3 2019

yaxunl updated the diff for D70980: [HIP] Remove opencl.amdgcn.lib.

fix test

Dec 3 2019, 1:40 PM · Restricted Project
yaxunl created D70980: [HIP] Remove opencl.amdgcn.lib.
Dec 3 2019, 1:12 PM · Restricted Project
yaxunl updated the diff for D68578: [HIP] Fix device stub name.

clean up and fix assertions.

Dec 3 2019, 12:53 PM
yaxunl updated the diff for D68578: [HIP] Fix device stub name.

use calling convention to mangle the stub differently.

Dec 3 2019, 9:37 AM

Nov 26 2019

yaxunl committed rGded249049429: Workaround for EvalInfo ctor for MSVC 2017 (authored by yaxunl).
Workaround for EvalInfo ctor for MSVC 2017
Nov 26 2019, 6:45 PM
yaxunl closed D70729: Workaround for EvalInfo ctor for MSVC 2017.
Nov 26 2019, 6:45 PM · Restricted Project
yaxunl updated the diff for D70729: Workaround for EvalInfo ctor for MSVC 2017.

add a test

Nov 26 2019, 10:40 AM · Restricted Project
yaxunl created D70729: Workaround for EvalInfo ctor for MSVC 2017.
Nov 26 2019, 10:21 AM · Restricted Project

Nov 25 2019

yaxunl added a comment to D70424: clang/AMDGPU: Fix default for frame-pointer attribute.

LGTM. But I am wondering how it affects -g. Do we need to keep frame pointer when -g is specified? Should we add a test for -O3 -g?

Nov 25 2019, 8:24 AM

Nov 20 2019

yaxunl added inline comments to D59321: WIP: AMDGPU: Teach toolchain to link rocm device libs.
Nov 20 2019, 1:56 PM
yaxunl added a comment to D59321: WIP: AMDGPU: Teach toolchain to link rocm device libs.

Add Scott since this may affect comgr. Probably need to add -nodefaultlibs in comgr after this change.

Nov 20 2019, 1:51 PM
yaxunl added a reviewer for D59321: WIP: AMDGPU: Teach toolchain to link rocm device libs: scott.linder.
Nov 20 2019, 1:51 PM

Nov 14 2019

yaxunl added inline comments to D68578: [HIP] Fix device stub name.
Nov 14 2019, 2:55 PM
yaxunl added a comment to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.
In D70172#1745998, @rnk wrote:

Are we sure using both Itanium and MS C++ ABIs at the same time is really the best way forward here? What are the constraints on CUDA that require the Itanium ABI? I'm sure there are real reasons you can't just use the MS ABI as is, but I'm curious what they are. Was there some RFC or design showing that this is the right way forward?

I wonder if it would be more productive to add new, more expansive attributes, similar to __attribute__((ms_struct)), that tag class or function decls as MS or Itanium C++ ABI. CUDA could then leverage this as needed, and it would be much easier to construct test cases for MS/Itanium interop. This is an expansion in scope, but it seems like it could be generally useful, and if we're already going to enter the crazy world of multiple C++ ABIs in a single TU, we might as well bite the bullet and do it in a way that isn't specific to CUDA.

Nov 14 2019, 1:02 PM
yaxunl added inline comments to D68578: [HIP] Fix device stub name.
Nov 14 2019, 9:02 AM

Nov 13 2019

yaxunl added a comment to D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.

sorry I think I misunderstood the meaning of "blocking" so I put it back.

Nov 13 2019, 6:39 PM
yaxunl added 1 blocking reviewer(s) for D70172: [CUDA][HIP] Fix assertion due to dtor check on windows: rsmith.
Nov 13 2019, 6:39 PM
yaxunl removed 1 blocking reviewer(s) for D70172: [CUDA][HIP] Fix assertion due to dtor check on windows: rsmith.
Nov 13 2019, 11:23 AM
yaxunl updated the diff for D68578: [HIP] Fix device stub name.

Attempt to prefix the kernel stub name on the fly.

Nov 13 2019, 11:20 AM
yaxunl accepted D70191: [AMDGPU][HIP] Change default DWARF version to 4.

LGTM. Thanks.

Nov 13 2019, 10:13 AM · Restricted Project
yaxunl created D70172: [CUDA][HIP] Fix assertion due to dtor check on windows.
Nov 13 2019, 4:43 AM

Nov 7 2019

yaxunl added a comment to D68578: [HIP] Fix device stub name.

Distinguishing between multiple symbols associated with the same source-level declaration is the purpose of the GlobalDecl abstraction.

It seems GlobalDecl is just a wrapper for concrete Decl's

It's a Decl plus a discriminator which is required for certain kinds of declaration. See e.g. GlobalDecl(const CXXConstructorDecl *D, CXXCtorType Type). GlobalDecl asserts if you try to construct it using GlobalDecl(FunctionDecl*) with a constructor/destructor declaration; we could similarly make that forbid construction with a kernel and then require code to use a GlobalDecl constructor that passes down whether it's the kernel or the stub that's being requested.

John.

Nov 7 2019, 11:43 AM
yaxunl added a comment to D68578: [HIP] Fix device stub name.

Distinguishing between multiple symbols associated with the same source-level declaration is the purpose of the GlobalDecl abstraction.

Nov 7 2019, 11:07 AM

Nov 5 2019

yaxunl added a comment to D68578: [HIP] Fix device stub name.
In D68578#1700652, @tra wrote:

This patch proposes changing the source-level name for the stub. Unfortunately the way it attempt to implement it is by doing the renaming during mangling phase itself. This appears to be the wrong place to change source-level name.

Nov 5 2019, 11:52 AM
yaxunl accepted D63020: [HIP] Fix visibility for 'extern' device variables..

LGTM. Thanks! @scchan This may fix the undefined symbol in work item struct issue at -O0

Nov 5 2019, 10:57 AM · Restricted Project
yaxunl 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.

This should not be a concern. The coercing is only applied to the parameter itself. Within the function body, we still use the original struct A. The preparation in function prolog will copy that coerced argument into the original one (alloca-ed.) The modification of that parameter later will be applied to the original one due to the by-val nature.

A modified version of your code is compiled into the following code at O0:

define protected amdgpu_kernel void @_Z3foo1A(%struct.A.coerce %a.coerce) #0 {
entry:
  %a = alloca %struct.A, align 8, addrspace(5)
  %a1 = addrspacecast %struct.A addrspace(5)* %a to %struct.A*
  %x = alloca i32, align 4, addrspace(5)
  %x.ascast = addrspacecast i32 addrspace(5)* %x to i32*
  %agg.tmp = alloca %struct.A, align 8, addrspace(5)
  %agg.tmp.ascast = addrspacecast %struct.A addrspace(5)* %agg.tmp to %struct.A*
  %0 = bitcast %struct.A* %a1 to %struct.A.coerce*
  %1 = getelementptr inbounds %struct.A.coerce, %struct.A.coerce* %0, i32 0, i32 0
  %2 = extractvalue %struct.A.coerce %a.coerce, 0
  store i32 addrspace(1)* %2, i32 addrspace(1)** %1, align 8
  %3 = getelementptr inbounds %struct.A.coerce, %struct.A.coerce* %0, i32 0, i32 1
  %4 = extractvalue %struct.A.coerce %a.coerce, 1
  store i32 addrspace(1)* %4, i32 addrspace(1)** %3, align 8
  %p = getelementptr inbounds %struct.A, %struct.A* %a1, i32 0, i32 0
  store i32* %x.ascast, i32** %p, align 8
  %5 = bitcast %struct.A* %agg.tmp.ascast to i8*
  %6 = bitcast %struct.A* %a1 to i8*
  call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %5, i8* align 8 %6, i64 16, i1 false)
  %7 = getelementptr inbounds %struct.A, %struct.A* %agg.tmp.ascast, i32 0, i32 0
  %8 = load i32*, i32** %7, align 8
  %9 = getelementptr inbounds %struct.A, %struct.A* %agg.tmp.ascast, i32 0, i32 1
  %10 = load i32*, i32** %9, align 8
  call void @_Z1f1A(i32* %8, i32* %10) #3
  ret void
}

The modification of parameter a is applied the alloca-ed one.

Nov 5 2019, 10:56 AM · Restricted Project
yaxunl 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:

Nov 5 2019, 10:19 AM · Restricted Project

Nov 4 2019

yaxunl committed rG4264e7bbfdb3: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation (authored by yaxunl).
[CUDA][HIP] Disable emitting llvm.linker.options in device compilation
Nov 4 2019, 8:28 PM
yaxunl closed D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.
Nov 4 2019, 8:28 PM · Restricted Project
yaxunl added a comment to D69826: [hip] Enable pointer argument lowering through coercing type..

we need a test for byval struct and array. better to have a struct containing an array which contain another struct which contains a pointer. thanks.

Nov 4 2019, 2:17 PM · Restricted Project
yaxunl abandoned D69818: [HIP] Fix pointer type kernel arg for amdgpu.

Please review Michael's patch https://reviews.llvm.org/D69826 which supersedes this one. Thanks.

Nov 4 2019, 2:17 PM
yaxunl added a comment to D69818: [HIP] Fix pointer type kernel arg for amdgpu.

BTW Michael Liao will create another patch which will supersede this patch. That patch contains similar changes and also handles pointers in a byval struct or array.

Nov 4 2019, 2:11 PM
yaxunl updated the diff for D69818: [HIP] Fix pointer type kernel arg for amdgpu.

add a test for non-kernel function.

Nov 4 2019, 12:52 PM
yaxunl created D69818: [HIP] Fix pointer type kernel arg for amdgpu.
Nov 4 2019, 12:50 PM
yaxunl added inline comments to D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.
Nov 4 2019, 11:17 AM · Restricted Project

Nov 3 2019

yaxunl updated the diff for D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.

revised by Artem's comments. Also fix test since linker option only supported on windows.

Nov 3 2019, 2:19 PM · Restricted Project

Nov 1 2019

yaxunl added inline comments to D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.
Nov 1 2019, 5:02 PM · Restricted Project

Oct 31 2019

yaxunl added a comment to D69679: [AMDGPU] Add amdgpu-promote-pointer-kernargs pass.
In D69679#1729466, @tra wrote:

@arsenm has a point. We can do it in clang, and it seems to be a better long-term solution compared to patching-up the inputs' AS that we've done for NVPTX and, now, AMDGPU.

If we do want to change the IR-level calling convention for the kernels, clang would not be the only place that would need to adapt to that change. We will need to think about transitioning existing external users, too (e.g. XLA in TensorFlow & JAX, julia). We may need to keep the promote-pointers-to-global-AS pass around for a while until the LLVM users have a chance to change their code to pass pointers using correct address space.

Oct 31 2019, 7:47 PM · Restricted Project
yaxunl committed rGbb1616ba4726: [CodeGen] Fix invalid llvm.linker.options about pragma detect_mismatch (authored by yaxunl).
[CodeGen] Fix invalid llvm.linker.options about pragma detect_mismatch
Oct 31 2019, 7:29 PM
yaxunl closed D69678: [CodeGen] Fix invalid llvm.linker.options about pragma detect_mismatch.
Oct 31 2019, 7:29 PM · Restricted Project
yaxunl accepted D69666: clang: Fix assert on void pointer arithmetic with address_space.

LGTM

Oct 31 2019, 6:16 PM
yaxunl added inline comments to D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.
Oct 31 2019, 6:07 PM · Restricted Project
yaxunl added a comment to D69679: [AMDGPU] Add amdgpu-promote-pointer-kernargs pass.

Clang should just be directly emitting the arguments with the global address space in the first place. It already has support for coercing argument types per calling convention and this is no different

Oct 31 2019, 2:24 PM · Restricted Project
yaxunl added reviewers for D69679: [AMDGPU] Add amdgpu-promote-pointer-kernargs pass: rjmccall, tra.
Oct 31 2019, 2:24 PM · Restricted Project
yaxunl added a comment to D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.
In D57829#1729094, @tra wrote:

Could you, please, give us a bit more context and provide more info for the questions @rjmccall and I asked before?

Is the problem specifically because autolink is not supported on device side? Or is disabling autolink just a convoluted way to avoid calling EmitModuleLinkOptions()?
If it's the former, then we should just disable it unconditionally -- we already filter out some other flags (e.g. asan).
If it's the latter, then tweaking autolink option behavior is just covering the problem. Sooner or later EmitModuleLinkOptions() will be called by something else and we'll be back to where we are now.

Oct 31 2019, 1:10 PM · Restricted Project
yaxunl created D69678: [CodeGen] Fix invalid llvm.linker.options about pragma detect_mismatch.
Oct 31 2019, 12:56 PM · Restricted Project
yaxunl added a comment to D69666: clang: Fix assert on void pointer arithmetic with address_space.

Is the description reversed?

Oct 31 2019, 10:20 AM

Oct 30 2019

yaxunl updated the diff for D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.

use -fno-autolink to disable linker option for device compilation.

Oct 30 2019, 9:07 PM · Restricted Project
yaxunl reclaimed D57829: [CUDA][HIP] Disable emitting llvm.linker.options in device compilation.

we still need this

Oct 30 2019, 8:21 PM · Restricted Project

Oct 29 2019

yaxunl created D69582: Let clang driver support parallel jobs.
Oct 29 2019, 1:26 PM

Oct 22 2019

yaxunl committed rG1c98ff49a30b: Fix name of warn_ignored_hip_only_option (authored by yaxunl).
Fix name of warn_ignored_hip_only_option
Oct 22 2019, 1:45 PM