eliben (Eli Bendersky)
User

Projects

User does not belong to any projects.

User Details

User Since
Jan 7 2013, 9:27 AM (305 w, 3 d)

Recent Activity

Feb 3 2017

eliben removed a reviewer for D28571: RuntimeDyldELF/AArch64: Implement basic support for PIC relocations: eliben.
Feb 3 2017, 9:49 AM

Aug 19 2016

eliben accepted D22530: [lanai] Make Lanai backend non-experimental.
Aug 19 2016, 9:43 AM

Aug 18 2016

eliben added a comment to D23693: [lanai] Add ISA document to CompilerWritersInfo.

lgtm

Aug 18 2016, 2:09 PM

Jul 29 2016

eliben committed rL277204: Add a REQUIRES: assert on a Lanai test that uses a -debug-only flag.
Add a REQUIRES: assert on a Lanai test that uses a -debug-only flag
Jul 29 2016, 12:43 PM

Jul 7 2016

eliben accepted D22072: [lanai] Use peephole optimizer to generate more conditional ALU operations..
Jul 7 2016, 4:41 PM
eliben added a comment to D22072: [lanai] Use peephole optimizer to generate more conditional ALU operations..

Since most of the actual peephole code is taken from the ARM backend, I left less comments on that part.

Jul 7 2016, 9:11 AM

May 6 2016

eliben added a comment to D19924: Representing bundle locked groups as fragments in MCAssembler.

I love how much more localized the blundle handling code is with this patch.

I have never used the feature myself, I I don't think I am comfortable being the main reviewer. Eli, what do you think?

May 6 2016, 11:40 AM

Feb 16 2016

eliben added inline comments to D17002: [lanai] Add Lanai backend to clang driver.
Feb 16 2016, 9:52 AM
eliben updated subscribers of D17008: [lanai] Add ELF enum value and relocations.
Feb 16 2016, 9:49 AM

Feb 9 2016

eliben added a comment to D17002: [lanai] Add Lanai backend to clang driver.

Clang-level tests?

Feb 9 2016, 1:23 PM

Feb 3 2016

eliben accepted D16874: [NVPTX] Disable performance optimizations when OptLevel==None.
Feb 3 2016, 6:27 PM

Sep 25 2015

eliben added a reviewer for D13181: [CUDA] 32-bit NVPTX: long should be 32bit: jingyue.
Sep 25 2015, 3:25 PM
eliben accepted D13171: [CUDA] Added a wrapper header for inclusion of stock CUDA headers..

lgtm

Sep 25 2015, 1:23 PM
eliben added inline comments to D13171: [CUDA] Added a wrapper header for inclusion of stock CUDA headers..
Sep 25 2015, 12:14 PM

Sep 1 2015

eliben accepted D12453: [CUDA] Allow function overloads based on host/device attributes..

The CUDA parts look very good. Someone else should approve the overloading-related logic

Sep 1 2015, 12:26 PM
eliben committed rL246573: Fix typo in test.
Fix typo in test
Sep 1 2015, 11:57 AM

Aug 31 2015

eliben added inline comments to D12453: [CUDA] Allow function overloads based on host/device attributes..
Aug 31 2015, 3:59 PM
eliben accepted D12414: [NVPTX] add an NVPTX-specific alias analysis.

lgtm

Aug 31 2015, 3:39 PM

Aug 20 2015

eliben accepted D12058: [NVPTX] truncating 64-bit to 32-bit is free.

lgtm

Aug 20 2015, 1:52 PM
eliben added inline comments to D12058: [NVPTX] truncating 64-bit to 32-bit is free.
Aug 20 2015, 11:58 AM

Aug 19 2015

eliben accepted D11950: [CUDA] Check register names on appropriate side of cuda compilation only..

lgtm

Aug 19 2015, 2:53 PM
eliben accepted D11694: [CUDA] Added stubs for __nvvm_atom_add*_d().

lgtm

Aug 19 2015, 1:51 PM
eliben accepted D12122: [CUDA] Add appropriate host/device attribute to target-specific builtins..

lgtm

Aug 19 2015, 8:09 AM

Aug 18 2015

eliben added inline comments to D12122: [CUDA] Add appropriate host/device attribute to target-specific builtins..
Aug 18 2015, 4:05 PM

Aug 12 2015

eliben accepted D11993: [CUDA] Make sure we emit all templated __global__ functions on device side. Again..

lgtm

Aug 12 2015, 2:44 PM

Aug 11 2015

eliben accepted D11926: [NVPTX] Use 32-bit divides instead of 64-bit divides where possible.

lgtm

Aug 11 2015, 1:22 PM
eliben added inline comments to D11950: [CUDA] Check register names on appropriate side of cuda compilation only..
Aug 11 2015, 12:30 PM

Aug 5 2015

eliben added inline comments to D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 11:40 AM

Aug 1 2015

eliben added a comment to D11622: [NVPTX] allow register copy between float and int.

This looks reasonable for now. Longer term, I'd like to experiment with doing away with the typed registers completely and just use .bXX for all registers. Typing is nice for readability, but code gen can be cleaner without it. They all compile down to the same set of registers at the sass level, and I'd like to match the hardware as much as possible (given that we can't just generate sass).

Aug 1 2015, 8:19 AM

Jul 17 2015

eliben committed rL242542: Use inbounds GEPs for memcpy and memset lowering.
Use inbounds GEPs for memcpy and memset lowering
Jul 17 2015, 9:42 AM

Jul 16 2015

eliben accepted D11281: Doxygen: Enable autobrief feature, matching llvm config/coding standards..

lgtm

Jul 16 2015, 4:03 PM
eliben committed rL242439: Streamline the coding style in NVPTXLowerAggrCopies.
Streamline the coding style in NVPTXLowerAggrCopies
Jul 16 2015, 1:42 PM
eliben accepted D11201: [NVPTX] enable SpeculativeExecution in NVPTX.

lgtm

Jul 16 2015, 9:58 AM
eliben committed rL242413: Correct lowering of memmove in NVPTX.
Correct lowering of memmove in NVPTX
Jul 16 2015, 9:27 AM
eliben closed D11220: Correct lowering of memmove in NVPTX by committing rL242413: Correct lowering of memmove in NVPTX.
Jul 16 2015, 9:27 AM

Jul 15 2015

eliben updated the diff for D11220: Correct lowering of memmove in NVPTX.

Updated with comments

Jul 15 2015, 12:34 PM
eliben added a comment to D11220: Correct lowering of memmove in NVPTX.

Thanks for the comments

Jul 15 2015, 12:33 PM
eliben added a comment to D11220: Correct lowering of memmove in NVPTX.

Hi,

My name is Okwan Kwon, and I have two comments.

  1. Use getRawDest() and getRawSource() instead. getDest() and getSource() will strip off any casts to give the original pointer. When the original pointer type is not i8 *, then it will get an assertion.
Jul 15 2015, 12:33 PM
eliben added a comment to D11220: Correct lowering of memmove in NVPTX.

The algorithm looks good. Thanks for working on this!

Though I'm wondering if this shouldn't be moved to the CodeGen library. Other targets may be able to benefit from this, like AMDGPU.

Jul 15 2015, 8:11 AM
eliben retitled D11220: Correct lowering of memmove in NVPTX from to Correct lowering of memmove in NVPTX.
Jul 15 2015, 7:59 AM

Jul 10 2015

eliben committed rL241914: Actually support volatile memcpys in NVPTX lowering.
Actually support volatile memcpys in NVPTX lowering
Jul 10 2015, 8:40 AM
eliben closed D11091: Actually support volatile memcpys in NVPTX lowering by committing rL241914: Actually support volatile memcpys in NVPTX lowering.
Jul 10 2015, 8:40 AM
eliben added a comment to D11091: Actually support volatile memcpys in NVPTX lowering.

LGTM

More of a feature request, llvm.memcpy now takes alignments as arguments. We should emit larger loads/stores (e.g. st.u32 or even vector store) when alignments permit. Worth a TODO?

Jul 10 2015, 8:29 AM

Jul 9 2015

eliben updated subscribers of D11091: Actually support volatile memcpys in NVPTX lowering.
Jul 9 2015, 4:28 PM
eliben retitled D11091: Actually support volatile memcpys in NVPTX lowering from to Actually support volatile memcpys in NVPTX lowering.
Jul 9 2015, 4:27 PM
eliben committed rL241875: Replace index-loops by range-based loops.
Replace index-loops by range-based loops
Jul 9 2015, 4:06 PM

Jul 8 2015

eliben committed rL241736: Add tests for the NVPTXLowerAggrCopies pass..
Add tests for the NVPTXLowerAggrCopies pass.
Jul 8 2015, 2:29 PM
eliben committed rL241687: Cosmetic cleanups - NFC.
Cosmetic cleanups - NFC
Jul 8 2015, 9:33 AM

Jun 24 2015

eliben accepted D10666: [CUDA] Implemented __nvvm_atom_*_gen_* builtins..

I'd move the NVPTX builtin emission code to its own file unless others have strong objections, but looks good otherwise.

Jun 24 2015, 3:50 PM
eliben added inline comments to D10666: [CUDA] Implemented __nvvm_atom_*_gen_* builtins..
Jun 24 2015, 3:46 PM
eliben added inline comments to D10666: [CUDA] Implemented __nvvm_atom_*_gen_* builtins..
Jun 24 2015, 8:08 AM
eliben accepted D10664: Fixed a typo in __nvvm_atom_min_gen_l() type string. .

lgtm

Jun 24 2015, 7:56 AM

Jun 23 2015

eliben accepted D10663: Added missing test case for llvm.nvvm.sqrt.f NVPTX intrinsic.

LGTM - looks pretty straightforward

Jun 23 2015, 10:52 AM

Jun 18 2015

eliben updated subscribers of D10549: Add NVPTXPeephole pass to reduce unnecessary address cast.
Jun 18 2015, 4:14 PM

Jun 12 2015

eliben committed rL239646: Fix returning error message in LLVMLinkModules.
Fix returning error message in LLVMLinkModules
Jun 12 2015, 4:31 PM
eliben closed D10241: Fix returning error message in LLVMLinkModules by committing rL239646: Fix returning error message in LLVMLinkModules.
Jun 12 2015, 4:31 PM

Jun 10 2015

eliben accepted D10241: Fix returning error message in LLVMLinkModules.

LGTM

Jun 10 2015, 8:43 AM

Jun 9 2015

eliben committed rL239411: Add more wrappers for symbol APIs to the C API..
Add more wrappers for symbol APIs to the C API.
Jun 9 2015, 9:02 AM
eliben closed D10222: Augment the C API by committing rL239411: Add more wrappers for symbol APIs to the C API..
Jun 9 2015, 9:01 AM
eliben added a comment to D10222: Augment the C API.

Do you have commit rights, Antoine?

Jun 9 2015, 8:35 AM

Jun 8 2015

eliben accepted D10322: [NVPTX] run SROA after NVPTXFavorNonGenericAddrSpaces.

LGTM

Jun 8 2015, 2:52 PM

Jun 3 2015

eliben accepted D10222: Augment the C API.

lgtm

Jun 3 2015, 4:50 PM

May 29 2015

eliben accepted D10132: [docs] fix the declarations of the llvm.nvvm.ptr.gen.to.* intrinsics.

lgtm

May 29 2015, 2:35 PM

May 6 2015

eliben accepted D9507: [cuda] Include GPU binary into host object file and generate init/deinit code..

[assuming this just copies stuff from D8463 w/o modifications]

May 6 2015, 8:44 AM

Apr 28 2015

eliben added a comment to D9327: [cuda] Preserve TLS storage class of host variable during device-side compilation..

[with a small nit]

Apr 28 2015, 11:33 AM
eliben accepted D9327: [cuda] Preserve TLS storage class of host variable during device-side compilation..
Apr 28 2015, 11:33 AM
eliben added a comment to D9327: [cuda] Preserve TLS storage class of host variable during device-side compilation..

Is there any way to create a test that would catch this?

Apr 28 2015, 10:19 AM

Apr 27 2015

eliben accepted D9270: [cuda] treat file scope __asm as __host__ and ignore it during device-side compilation..

LGTM

Apr 27 2015, 9:48 AM
eliben accepted D9269: [cuda] Ignore "TLS unsupported by target" errors for host variables during device-side CUDA compilation. .

LGTM

Apr 27 2015, 9:47 AM

Apr 17 2015

eliben added inline comments to D9064: [CUDA] Support for built-in cuda variables (threadIdx and its friends)..
Apr 17 2015, 4:08 PM
eliben accepted D9066: [NVPTX] enable NaryReassociate in NVPTX.

LGTM

Apr 17 2015, 9:25 AM

Apr 15 2015

eliben committed rL235049: Create a frontend flag to disable CUDA cross-target call checks.
Create a frontend flag to disable CUDA cross-target call checks
Apr 15 2015, 3:30 PM
eliben closed D9036: Create a frontend flag to disable CUDA cross-target call checks by committing rL235049: Create a frontend flag to disable CUDA cross-target call checks.
Apr 15 2015, 3:30 PM
eliben added a comment to D9036: Create a frontend flag to disable CUDA cross-target call checks.
In D9036#156719, @rnk wrote:

As discussed, this flag actually affects overload resolution results, so it should probably be a language option as it is now. Maybe we shouldn't be having this target check affect overload results, but that's another big issue that we don't need to run down right this instant.

lgtm

Apr 15 2015, 2:57 PM
eliben added a comment to D9036: Create a frontend flag to disable CUDA cross-target call checks.
In D9036#156680, @tra wrote:

LGTM.
On a side note, do you still need fcuda_allow_host_calls_from_host_device ?

Apr 15 2015, 1:19 PM
eliben added inline comments to D9036: Create a frontend flag to disable CUDA cross-target call checks.
Apr 15 2015, 1:08 PM
eliben retitled D9036: Create a frontend flag to disable CUDA cross-target call checks from to Create a frontend flag to disable CUDA cross-target call checks.
Apr 15 2015, 12:25 PM
eliben accepted D8983: [SLSR] handle candidate form (B + i * S).

LGTM

Apr 15 2015, 7:57 AM

Apr 13 2015

eliben added a comment to D8985: [CUDA] Allow using integral non-type template parameters as launch_bounds attribute arguments..

The change looks fine. I'm not an expert on the template instantiation code, so you'll need to get an OK from Richard or someone else more familiar with this part.

Apr 13 2015, 9:40 AM

Apr 8 2015

eliben added a comment to D8463: End-to-end CUDA compilation..

non-driver parts LGTM

Apr 8 2015, 8:29 AM

Apr 6 2015

eliben added a comment to D8463: End-to-end CUDA compilation..

Where are the tests for emitting ctors/dtors, registering kernels, etc?

Apr 6 2015, 3:36 PM
eliben added a comment to D8463: End-to-end CUDA compilation..

A couple of replies to comments; will do another pass on the new revision

Apr 6 2015, 3:02 PM

Apr 3 2015

eliben accepted D8072: [MC] Write padding into fragments when -mc-relax-all flag is used.

LGTM

Apr 3 2015, 2:38 PM
eliben added a comment to D8463: End-to-end CUDA compilation..

Thanks for the updated description.

Apr 3 2015, 2:30 PM

Apr 2 2015

eliben added a comment to D8785: Fix typo and improve a couple sentences.

Committed it for you, Douglas, since you don't have commit rights yet.
Thanks for the patch

Apr 2 2015, 8:24 AM
eliben committed rL233920: Fix typo and reword in LangRef.
Fix typo and reword in LangRef
Apr 2 2015, 8:23 AM
eliben closed D8785: Fix typo and improve a couple sentences by committing rL233920: Fix typo and reword in LangRef.
Apr 2 2015, 8:22 AM

Apr 1 2015

eliben accepted D8785: Fix typo and improve a couple sentences.

lgtm

Apr 1 2015, 3:35 PM
eliben committed rL233836: Set the type of ptrdiff_t to signed on NVPTX targets.
Set the type of ptrdiff_t to signed on NVPTX targets
Apr 1 2015, 1:32 PM
eliben committed rL233830: Set NVPTX64 target's size_t to match other 64-bit targets.
Set NVPTX64 target's size_t to match other 64-bit targets
Apr 1 2015, 11:32 AM

Mar 31 2015

eliben committed rL233715: Add sm_37 target to Clang for NVPTX.
Add sm_37 target to Clang for NVPTX
Mar 31 2015, 10:06 AM

Mar 25 2015

eliben added a comment to D8575: Fix addrspace when emitting constructors of static local variables.

I think it's alright to handle the dtor crash in a separate change (if it looks like a separate issue). Just open a bug about it

Mar 25 2015, 5:20 AM

Mar 24 2015

eliben added a comment to D8575: Fix addrspace when emitting constructors of static local variables.

Oh another something - do we have tests that other special members work in this case? E.g. if StructWithCtor also had a copy ctor or dtor, for example

Mar 24 2015, 3:31 PM
eliben accepted D8575: Fix addrspace when emitting constructors of static local variables.

LGTM - the comment w/ the example really helps. Thanks

Mar 24 2015, 1:02 PM
eliben added a comment to D8575: Fix addrspace when emitting constructors of static local variables.

In general this looks good; there are quite a few layers of calls to emit the global var initialization, and I'm not familiar enough to say if this is inserted at the right layer. Hopefully one of the other reviewers will be able to clarify this.

Mar 24 2015, 6:51 AM

Mar 23 2015

eliben added inline comments to D8575: Fix addrspace when emitting constructors of static local variables.
Mar 23 2015, 11:56 PM
eliben committed rL233043: Cleanup: no need to pass DefinitionKind into ParseCXXInlineMethodDef.
Cleanup: no need to pass DefinitionKind into ParseCXXInlineMethodDef
Mar 23 2015, 4:52 PM
eliben closed D5541: CUDA: mark implicit intrinsics properly.
Mar 23 2015, 3:27 PM
eliben added a comment to D8466: Fix clang-tidy to account of correct source location of defaulted/deleted members.

Thanks for the review!

Mar 23 2015, 3:16 PM
eliben committed rL233032: Fix clang-tidy to not assume wrong source locations for defaulted members..
Fix clang-tidy to not assume wrong source locations for defaulted members.
Mar 23 2015, 3:16 PM