Page MenuHomePhabricator

JonChesterfield (Jon Chesterfield)
User

Projects

User does not belong to any projects.

User Details

User Since
Aug 11 2015, 1:08 PM (304 w, 4 d)

Recent Activity

Yesterday

JonChesterfield accepted D103813: [AMDGPU][Libomptarget] Drop dead code related to g_atl_machine.

LG. My preference would be as commented above, but the vector containing the contents of the other two works too.

Fri, Jun 11, 1:03 PM · Restricted Project
JonChesterfield added inline comments to D103813: [AMDGPU][Libomptarget] Drop dead code related to g_atl_machine.
Fri, Jun 11, 5:24 AM · Restricted Project

Mon, Jun 7

JonChesterfield added a comment to D103817: [AMDGPU] Introduce command line switch to control super aligning of LDS..

Thanks!

Mon, Jun 7, 3:28 PM · Restricted Project
Herald added a reviewer for D103813: [AMDGPU][Libomptarget] Drop dead code related to g_atl_machine: jdoerfert.
Mon, Jun 7, 7:37 AM · Restricted Project
JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

Yes, I see that you agree with the wording you committed. We can't fix it now.

Mon, Jun 7, 6:41 AM · Restricted Project
JonChesterfield added inline comments to D103225: [AMDGPU] Replace non-kernel function uses of LDS globals by pointers..
Mon, Jun 7, 6:16 AM · Restricted Project
JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

I note that you did not change the commit message to reflect the contents of the patch as requested.

The commit message says, what the the patch is doing - it increases the alignment of LDS if it is under-aligned in order to reduce the probability of unlinged access. So, I do not think, any change to commit message is required.

Mon, Jun 7, 6:11 AM · Restricted Project
JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

I note that you did not change the commit message to reflect the contents of the patch as requested.

Mon, Jun 7, 5:54 AM · Restricted Project
JonChesterfield accepted D103795: [AMDGPU][Libomptarget] Remove atlc global.

Nice, thanks. The task support was nominally deleted some time ago, iirc the init() part survived because the variable was used elsewhere as well. Good to see it go.

Mon, Jun 7, 3:48 AM · Restricted Project

Fri, Jun 4

JonChesterfield accepted D103545: [NFC][libomptarget] Reduce the dependency on libelf.

Thanks!

Fri, Jun 4, 12:09 PM · Restricted Project
JonChesterfield accepted D103600: [AMDGPU][Libomptarget] Rework logic for locating kernarg pools.

That worked out cleanly, thanks. The temporary vector used to indicate failure with .empty is slightly unusual but obviously works.

Fri, Jun 4, 5:25 AM · Restricted Project
JonChesterfield added a comment to D103600: [AMDGPU][Libomptarget] Rework logic for locating kernarg pools.

Took a while to make sense of this. I think you've preserved the existing strategy of building a collection of kernarg memory pools, while moving the global into a class member and adding some more tests on whether a pool should be added to the collection.

Fri, Jun 4, 2:12 AM · Restricted Project
JonChesterfield added a comment to D103431: [AMDGPU] Fix missing lowering of LDS used in global scope..

If the current logic of shouldLowerLDSToStruct() is more robust compere to earlier ones, then, why cannot we use it, instead of patching-up earlier ones?

Fri, Jun 4, 2:00 AM · Restricted Project
JonChesterfield updated subscribers of D103545: [NFC][libomptarget] Reduce the dependency on libelf.

Looks functionally correct. Always in favour of moving stuff out of headers where we can, though as an unrelated change we might want to start building plugins with flto or similar to ensure that remains free.

Fri, Jun 4, 1:52 AM · Restricted Project
JonChesterfield added a comment to D103546: Added ELFObjectFileBase::checkMagic() for checking ELF magic word..

git-clang-format HEAD^ may be useful for reformatting the already changed parts. LGTM, @jhenderson?

Fri, Jun 4, 1:44 AM · Restricted Project
JonChesterfield added a comment to D103671: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

Please continue the review in the existing place (i.e. reopen it) instead of making a new one

Fri, Jun 4, 12:40 AM · Restricted Project
JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

If we're going with this anyway, let's be explicit in the title and commit message that we are over-aligning data. We aren't 'fixing' things, or doing them 'properly', we are intentionally burning memory in the hope of improving runtime.

Fri, Jun 4, 12:39 AM · Restricted Project

Wed, Jun 2

JonChesterfield accepted D103509: [AMDGPU][Libomptarget][NFC] Remove bunch of dead structs.
Wed, Jun 2, 3:08 AM · Restricted Project
JonChesterfield accepted D103508: [AMDGPU][Libomptarget][NFC] Remove atmi_place_t.

LG, thanks!

Wed, Jun 2, 3:03 AM · Restricted Project

Tue, Jun 1

JonChesterfield added a comment to D102883: [OpenMP] Define named constants for interop's foreign runtime ID.

Does anyone have a reference for what we're supposed to do with one of these things?

Tue, Jun 1, 8:18 AM · Restricted Project
JonChesterfield added a comment to D103431: [AMDGPU] Fix missing lowering of LDS used in global scope..

For that test, I note that commenting out the llvm.used statement results in the variable being lowered. I.e. the following works fine:

Tue, Jun 1, 7:14 AM · Restricted Project
JonChesterfield added a comment to D103431: [AMDGPU] Fix missing lowering of LDS used in global scope..

Is this the proposed fix for the missing lowering of the LDS in the following?

Tue, Jun 1, 6:55 AM · Restricted Project
JonChesterfield accepted D103443: [AMDGPU][Libomptarget] Remove g_atmi_machine global.

Iirc it did more things in ATMI. Nice simplification here.

Tue, Jun 1, 3:03 AM · Restricted Project

Mon, May 31

JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

I'm not very convinced by this given the recent enthusiasm for decreasing LDS usage. Types are naturally aligned by default, so I think the only time they are aligned by less than that is when the programmer asked for it (or when the vectorizer is involved, but that's disabled on amdgpu afaik).

Mon, May 31, 12:16 PM · Restricted Project
JonChesterfield accepted D103389: [AMDGPU][Libomptarget][NFC] Split host and device malloc.

LGTM, thanks. Cleans up the call sites without changing the calls into HSA. Thanks for dropping the extern C shim too.

Mon, May 31, 2:28 AM · Restricted Project

Fri, May 28

JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

Not really the right place for this comment, but better odds of it being seen than in other bug trackers. Mahesha reports:

Fri, May 28, 12:34 PM · Restricted Project

Thu, May 27

JonChesterfield added a comment to D103261: [AMDGPU] Increase alignment of LDS globals if necessary before LDS lowering..

What's the use case for this? It will increase memory use if it increases the alignment of any variables

Thu, May 27, 12:09 PM · Restricted Project
JonChesterfield added a comment to rGbb8622094d77: AMDGPU: Don't handle kernarg.segment.ptr in functions.

We could lower to the equivalent of

void * kernarg_segment_pointer()
{
  __attribute__((address_space(4))) void* p = __builtin_amdgcn_dispatch_ptr();
  void * res;
  __builtin_memcpy(&res, (char*)p + (320/8), 8);
  return res;
}

instead of zero when not in a kernel function. Which is what I've just replaced a call to the builtin with, after a fairly lengthy process of trying to guess why the GPU was reporting memory corruption.

Thu, May 27, 11:18 AM
JonChesterfield accepted D103239: [AMDGPU][Libomptarget][NFC] Remove atmi_mem_place_t.

Nice, like that a lot. I hadn't noticed that we are passing device id == 0 around to mean a CPU device, that seems a path to considerable confusion. Probably better to split the allocators into gpu (which takes a device id) and cpu (which doesn't) in some later patch.

Thu, May 27, 4:40 AM · Restricted Project
JonChesterfield accepted D103222: Fixed bug using string stream by including missing header.
Thu, May 27, 2:47 AM · Restricted Project
JonChesterfield requested changes to D103222: Fixed bug using string stream by including missing header.

Didn't realize we were using string stream. Probably don't want to, but thanks for the fix in the meantime

Thu, May 27, 2:47 AM · Restricted Project

Wed, May 26

JonChesterfield committed rG2fdf8bbd19c3: [libomptarget][nfc][amdgpu] Factor out setting upper bounds (authored by JonChesterfield).
[libomptarget][nfc][amdgpu] Factor out setting upper bounds
Wed, May 26, 11:58 AM
JonChesterfield closed D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds.
Wed, May 26, 11:58 AM · Restricted Project
JonChesterfield committed rGc5c1ec7945ff: [libomptarget][nfc][amdgpu] Refactor uses of KernelInfoTable (authored by JonChesterfield).
[libomptarget][nfc][amdgpu] Refactor uses of KernelInfoTable
Wed, May 26, 11:27 AM
JonChesterfield closed D103093: [libomptarget][nfc][amdgpu] Refactor uses of KernelInfoTable.
Wed, May 26, 11:27 AM · Restricted Project
JonChesterfield accepted D103190: [AMDGPU] Fix kernel LDS lowering for constants.

Ouch. Good to catch that quickly, thanks

Wed, May 26, 11:26 AM · Restricted Project
JonChesterfield added inline comments to D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds.
Wed, May 26, 11:24 AM · Restricted Project
JonChesterfield updated the diff for D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds.
  • twist code around clang-tidy
Wed, May 26, 11:23 AM · Restricted Project
JonChesterfield added inline comments to D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds.
Wed, May 26, 11:18 AM · Restricted Project
JonChesterfield committed rG07f59baad634: [libomptarget][nfc][amdgpu] Remove atmi_status_t type (authored by JonChesterfield).
[libomptarget][nfc][amdgpu] Remove atmi_status_t type
Wed, May 26, 9:02 AM
JonChesterfield closed D103115: [libomptarget][nfc][amdgpu] Remove atmi_status_t type.
Wed, May 26, 9:02 AM · Restricted Project
JonChesterfield updated the diff for D103115: [libomptarget][nfc][amdgpu] Remove atmi_status_t type.
  • rebase
Wed, May 26, 8:38 AM · Restricted Project
JonChesterfield added a comment to D103139: [AMDGPU] Replace non-kernel function uses of LDS globals by pointers..

Use different passes for different things, and optimisations are different from lowering. Those should not be contentious statements in a multipass compiler. Strictly I'm not sure this is an optimisation as it burns cycles to reduce memory, and that's only useful if the saved memory improves occupancy.

Wed, May 26, 7:05 AM · Restricted Project
JonChesterfield added a comment to D103139: [AMDGPU] Replace non-kernel function uses of LDS globals by pointers..

I see that. However, we are working on a multipass compiler with a clearly delimited format used between passes.

Wed, May 26, 6:03 AM · Restricted Project
JonChesterfield requested changes to D103139: [AMDGPU] Replace non-kernel function uses of LDS globals by pointers..

Same objection to previous diffs with this code in. You're adding loads of complicated code to an existing pass.

Wed, May 26, 4:45 AM · Restricted Project
JonChesterfield accepted D102692: [AMDGPU][Libomptarget] Move Kernel/Symbol info tables to RTLDeviceInfoTy.

LG, thanks!

Wed, May 26, 2:38 AM · Restricted Project

Tue, May 25

JonChesterfield updated the summary of D101935: [clang] Search runtimes build tree for openmp runtime.
Tue, May 25, 3:18 PM · Restricted Project
JonChesterfield updated the summary of D101960: [openmp] Drop requirement on library path environment variables.
Tue, May 25, 3:18 PM · Restricted Project, Restricted Project
JonChesterfield updated the summary of D102043: [libomptarget] Look for plugins next to libomptarget.so.
Tue, May 25, 3:17 PM · Restricted Project
JonChesterfield updated the summary of D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds.
Tue, May 25, 3:17 PM · Restricted Project
JonChesterfield updated the summary of D103093: [libomptarget][nfc][amdgpu] Refactor uses of KernelInfoTable.
Tue, May 25, 3:17 PM · Restricted Project
JonChesterfield updated the summary of D103115: [libomptarget][nfc][amdgpu] Remove atmi_status_t type.
Tue, May 25, 3:17 PM · Restricted Project
JonChesterfield added a comment to D103058: [libomptarget][nfc] Move hostcall required test to rtl.

Thanks for pointing it out! I was under the impression that the contents of 'Summary' as shown in Phabricator turns into the commit message.

Tue, May 25, 3:16 PM · Restricted Project
JonChesterfield committed rGdf005fa364ae: [libomptarget][nfc] Move hostcall required test to rtl (authored by JonChesterfield).
[libomptarget][nfc] Move hostcall required test to rtl
Tue, May 25, 2:47 PM
JonChesterfield closed D103058: [libomptarget][nfc] Move hostcall required test to rtl.
Tue, May 25, 2:47 PM · Restricted Project
JonChesterfield accepted D103058: [libomptarget][nfc] Move hostcall required test to rtl.
Tue, May 25, 2:43 PM · Restricted Project
JonChesterfield added a comment to D102692: [AMDGPU][Libomptarget] Move Kernel/Symbol info tables to RTLDeviceInfoTy.

passing by const& is much better than passing by &, I think we should do that in this patch

Tue, May 25, 2:23 PM · Restricted Project
JonChesterfield added a comment to D103115: [libomptarget][nfc][amdgpu] Remove atmi_status_t type.

Deleted the references to ATMI_STATUS_UNKNOWN manually, then

for f in `find . -type f`; do sed -i 's$ATMI_STATUS_SUCCESS$HSA_STATUS_SUCCESS$g' $f; sed -i 's$ATMI_STATUS_ERROR$HSA_STATUS_ERROR$g' $f; sed -i 's$atmi_status_t$hsa_status_t$g' $f; done

and clang-format

Tue, May 25, 2:05 PM · Restricted Project
JonChesterfield requested review of D103115: [libomptarget][nfc][amdgpu] Remove atmi_status_t type.
Tue, May 25, 1:56 PM · Restricted Project
JonChesterfield requested review of D103093: [libomptarget][nfc][amdgpu] Refactor uses of KernelInfoTable.
Tue, May 25, 8:49 AM · Restricted Project
Herald added a reviewer for D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds: jdoerfert.
Tue, May 25, 8:40 AM · Restricted Project
JonChesterfield requested review of D103090: [libomptarget][nfc][amdgpu] Factor out setting upper bounds.
Tue, May 25, 8:38 AM · Restricted Project
JonChesterfield added inline comments to D102692: [AMDGPU][Libomptarget] Move Kernel/Symbol info tables to RTLDeviceInfoTy.
Tue, May 25, 8:06 AM · Restricted Project
JonChesterfield added a comment to D102692: [AMDGPU][Libomptarget] Move Kernel/Symbol info tables to RTLDeviceInfoTy.

except for the lost error check above, I think this is now a non-functional change. LG with that error path restored

Tue, May 25, 8:00 AM · Restricted Project
JonChesterfield added a comment to D103058: [libomptarget][nfc] Move hostcall required test to rtl.

Was reviewed on aomp, will wait a little while before landing this to match.

Tue, May 25, 5:48 AM · Restricted Project
JonChesterfield accepted D103078: [AMDGPU][Libomptarget] Mark lambda_by_value test as XFAIL.
Tue, May 25, 4:37 AM · Restricted Project
JonChesterfield added a comment to D103058: [libomptarget][nfc] Move hostcall required test to rtl.

Side point on races - I think it is possible for one thread to start loading a binary, get immediately suspended, and a second thread to then launch a kernel that is supposed to come from that binary. The locking scheme in this plugin probably doesn't handle that. Hopefully libomptarget does have locking that makes that work, otherwise the cuda plugin is probably also racy.

Tue, May 25, 1:52 AM · Restricted Project
JonChesterfield added a comment to D103030: [libomptarget][nfc] Accept callable for hsa iterate_symbols.

Looks good. My opinion is that we should keep these wrappers in a separate file, but that could be a later patch.

Tue, May 25, 1:34 AM · Restricted Project
JonChesterfield committed rG75492e20fb7c: [libomptarget][nfc] Accept callable for hsa iterate_symbols (authored by JonChesterfield).
[libomptarget][nfc] Accept callable for hsa iterate_symbols
Tue, May 25, 1:29 AM
JonChesterfield closed D103030: [libomptarget][nfc] Accept callable for hsa iterate_symbols.
Tue, May 25, 1:29 AM · Restricted Project
JonChesterfield accepted D102847: [AMDGPU][Libomptarget] Delete g_atmi_initialized.

Great, thanks!

Tue, May 25, 1:22 AM · Restricted Project
JonChesterfield accepted D103075: [AMDGPU][Libomptarget] Inline atmi_init/atmi_finalize.

Nice, thanks!

Tue, May 25, 1:21 AM · Restricted Project
JonChesterfield added a comment to D103059: [libomptarget] [amdgpu] Added LDS usage to the kernel trace.

Refactoring separately is good, but only works in practice if it's done first. Otherwise the refactor usually gets forgotten, or replaced with something higher priority, and that is the path to entropy.

Tue, May 25, 1:12 AM · Restricted Project

Mon, May 24

JonChesterfield accepted D103059: [libomptarget] [amdgpu] Added LDS usage to the kernel trace.

no objection to the functional change, would be good to refactor the local weirdness in passing

Mon, May 24, 6:55 PM · Restricted Project
JonChesterfield added inline comments to D103059: [libomptarget] [amdgpu] Added LDS usage to the kernel trace.
Mon, May 24, 6:54 PM · Restricted Project
Herald added a reviewer for D103058: [libomptarget][nfc] Move hostcall required test to rtl: jdoerfert.

Hostcall exists as a libraries in rocm, aomp and my github. Trunk has hooks - mostly weak symbols - to allow the host plugin source to work with or without it present.

Mon, May 24, 5:36 PM · Restricted Project
JonChesterfield requested review of D103058: [libomptarget][nfc] Move hostcall required test to rtl.
Mon, May 24, 5:17 PM · Restricted Project
JonChesterfield accepted D103037: [libomptarget] [amdgpu] Fix copy-paste error setting NumThreads for a corner case..

Could change these to std::min

Mon, May 24, 1:19 PM · Restricted Project
JonChesterfield added inline comments to D102692: [AMDGPU][Libomptarget] Move Kernel/Symbol info tables to RTLDeviceInfoTy.
Mon, May 24, 9:05 AM · Restricted Project
JonChesterfield added inline comments to D103030: [libomptarget][nfc] Accept callable for hsa iterate_symbols.
Mon, May 24, 8:49 AM · Restricted Project
JonChesterfield updated the diff for D103030: [libomptarget][nfc] Accept callable for hsa iterate_symbols.
  • drop dead executable argument
Mon, May 24, 8:48 AM · Restricted Project
JonChesterfield requested review of D103030: [libomptarget][nfc] Accept callable for hsa iterate_symbols.
Mon, May 24, 8:43 AM · Restricted Project

Fri, May 21

JonChesterfield added a comment to D102954: [AMDGPU] Lower kernel LDS into a sorted structure.

Couple of minor comments above, overall approach looks sound to me

Fri, May 21, 4:48 PM · Restricted Project
JonChesterfield committed rGd54712ab4deb: [libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation (authored by JonChesterfield).
[libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation
Fri, May 21, 8:09 AM
JonChesterfield closed D102499: [libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation.
Fri, May 21, 8:09 AM · Restricted Project
JonChesterfield accepted D102691: [AMDGPU][Libomptarget] Remove global KernelNameMap.

Like it, thanks!

Fri, May 21, 6:59 AM · Restricted Project
JonChesterfield added a comment to D102691: [AMDGPU][Libomptarget] Remove global KernelNameMap.

Couple of style comments above. I think this works as-is, so we have confirmation that the map of strings can go. Which is great!

Fri, May 21, 4:15 AM · Restricted Project

Thu, May 20

JonChesterfield added a comment to D102884: [AMDGPU] Request module used variables from LDS lowering as internal.

Ah, apologies. I misread/misremembered this. The global variable looked up here is llvm.used / llvm.compiler.used, and not the LDS variable that is going to be removed.

Thu, May 20, 6:16 PM · Restricted Project
JonChesterfield accepted D102882: [AMDGPU] Fix module LDS selection.

That is much prettier, thank you.

Thu, May 20, 3:45 PM · Restricted Project
JonChesterfield added a comment to D102499: [libomptarget][amdgpu] Mark alloc, free weak to facilitate local experimentation.

I'm expecting to have a few in flight locally while trying different designs. I don't expect this to be a long lasting patch.

Thu, May 20, 3:30 PM · Restricted Project
JonChesterfield added a comment to D94648: [amdgpu] Implement lower function LDS pass.

Hi Jon,

For the 0 argument kernel, did you put debug statement within SITargetLowering::LowerFormalArguments() and test whether it will hit or not? My experimentation shows that it is indeed hit for 0 arg kernels too. So it is not problem with 0 arg kernel.

Thu, May 20, 3:01 PM · Restricted Project
JonChesterfield committed rG68b88ae6701a: [libomptarget] Improve dlwrap compile time error diagnostic (authored by JonChesterfield).
[libomptarget] Improve dlwrap compile time error diagnostic
Thu, May 20, 12:34 PM
JonChesterfield closed D102858: [libomptarget] Improve dlwrap compile time error diagnostic.
Thu, May 20, 12:33 PM · Restricted Project
JonChesterfield requested review of D102858: [libomptarget] Improve dlwrap compile time error diagnostic.
Thu, May 20, 11:33 AM · Restricted Project
JonChesterfield added inline comments to D102691: [AMDGPU][Libomptarget] Remove global KernelNameMap.
Thu, May 20, 9:19 AM · Restricted Project
JonChesterfield committed rGd18fb09c6939: [libomptarget][amdgpu] Remove majority of fatal errors (authored by JonChesterfield).
[libomptarget][amdgpu] Remove majority of fatal errors
Thu, May 20, 8:27 AM
JonChesterfield closed D102346: [libomptarget][amdgpu] Remove majority of fatal errors.
Thu, May 20, 8:27 AM · Restricted Project
JonChesterfield added inline comments to D102847: [AMDGPU][Libomptarget] Delete g_atmi_initialized.
Thu, May 20, 8:26 AM · Restricted Project
JonChesterfield added inline comments to D102600: [AMDGPU][Libomptarget] Rename & move g_executables to private.
Thu, May 20, 8:23 AM · Restricted Project
JonChesterfield added a comment to D102691: [AMDGPU][Libomptarget] Remove global KernelNameMap.

Taken a closer look at this. Cut down example msgpack,

".name" : "__omp_vmul_l8",
".symbol" : "__omp_vmul_l8.kd",
Thu, May 20, 8:19 AM · Restricted Project