Page MenuHomePhabricator

JonChesterfield (Jon Chesterfield)
User

Projects

User does not belong to any projects.

User Details

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

Recent Activity

Today

JonChesterfield added inline comments to D95313: [WIP] Move part of nvptx devicertl under clang.
Sun, Jan 24, 10:41 AM · Restricted Project, Restricted Project
JonChesterfield requested review of D95313: [WIP] Move part of nvptx devicertl under clang.
Sun, Jan 24, 10:36 AM · Restricted Project, Restricted Project
JonChesterfield committed rGe5e448aafa76: [libomptarget][cuda] Fix build, change missed from D95274 (authored by JonChesterfield).
[libomptarget][cuda] Fix build, change missed from D95274
Sun, Jan 24, 10:30 AM
JonChesterfield committed rGc3074d48d38c: [libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics (authored by JonChesterfield).
[libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics
Sun, Jan 24, 2:59 AM
JonChesterfield closed D95294: [libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics.
Sun, Jan 24, 2:59 AM · Restricted Project
JonChesterfield added a comment to D95300: [OpenMP][deviceRTLs] Separate declaration of target dependent functions from `target_impl.h`.

'new header file target_impl.h' should be 'new header file target_interface.h' in the commit message

Sun, Jan 24, 2:53 AM · Restricted Project
JonChesterfield added inline comments to D95300: [OpenMP][deviceRTLs] Separate declaration of target dependent functions from `target_impl.h`.
Sun, Jan 24, 2:51 AM · Restricted Project

Yesterday

JonChesterfield added inline comments to D95298: [OpenMP][libomptarget][WIP] Refined the file structure of `deviceRTLs`.
Sat, Jan 23, 3:34 PM · Restricted Project
JonChesterfield added inline comments to D95294: [libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics.
Sat, Jan 23, 3:32 PM · Restricted Project
JonChesterfield added inline comments to D95298: [OpenMP][libomptarget][WIP] Refined the file structure of `deviceRTLs`.
Sat, Jan 23, 3:27 PM · Restricted Project
JonChesterfield added a comment to D95298: [OpenMP][libomptarget][WIP] Refined the file structure of `deviceRTLs`.

I really like about half of this change.

Sat, Jan 23, 3:26 PM · Restricted Project
JonChesterfield committed rGdc70c56be592: [libomptarget][amdgpu][nfc] Update comments (authored by JonChesterfield).
[libomptarget][amdgpu][nfc] Update comments
Sat, Jan 23, 2:54 PM
JonChesterfield closed D95295: [libomptarget][amdgpu][nfc] Update comments.
Sat, Jan 23, 2:54 PM · Restricted Project
JonChesterfield added inline comments to D95294: [libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics.
Sat, Jan 23, 2:27 PM · Restricted Project
JonChesterfield added inline comments to D95294: [libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics.
Sat, Jan 23, 1:30 PM · Restricted Project
JonChesterfield requested review of D95295: [libomptarget][amdgpu][nfc] Update comments.
Sat, Jan 23, 12:38 PM · Restricted Project
JonChesterfield committed rG78b0630b72a9: [libomptarget][cuda] Call v2 functions explicitly (authored by JonChesterfield).
[libomptarget][cuda] Call v2 functions explicitly
Sat, Jan 23, 12:33 PM
JonChesterfield closed D95274: [libomptarget][cuda] Call v2 functions explicitly.
Sat, Jan 23, 12:33 PM · Restricted Project
JonChesterfield requested review of D95294: [libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics.
Sat, Jan 23, 12:28 PM · Restricted Project
JonChesterfield added a comment to D95161: [Clang][OpenMP][NVPTX] Replace `libomptarget-nvptx-path` with `libomptarget-nvptx-bc-path`.

Still LGTM. A little surprised these are being run on windows, but good to be robust.

Sat, Jan 23, 11:18 AM · Restricted Project

Fri, Jan 22

JonChesterfield requested review of D95274: [libomptarget][cuda] Call v2 functions explicitly.
Fri, Jan 22, 4:29 PM · Restricted Project
JonChesterfield committed rG47e95e87a3e4: [libomptarget] Build cuda plugin without cuda installed locally (authored by JonChesterfield).
[libomptarget] Build cuda plugin without cuda installed locally
Fri, Jan 22, 4:15 PM
JonChesterfield closed D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Fri, Jan 22, 4:15 PM · Restricted Project
JonChesterfield added inline comments to D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Fri, Jan 22, 9:28 AM · Restricted Project
JonChesterfield updated the diff for D95155: [libomptarget] Build cuda plugin without cuda installed locally.
  • address review comments
Fri, Jan 22, 9:27 AM · Restricted Project
JonChesterfield added a comment to D77609: [OpenMP] Added the support for hidden helper task in RTL.

The information I've got on the possible race is:
When this patch is applied (by git's automerge, I think) to the rocm stack, a test located at:
https://github.com/ROCm-Developer-Tools/aomp/blob/master/test/smoke/devices/devices.c
fails in unpredictable fashion.

Fri, Jan 22, 8:33 AM · Restricted Project
JonChesterfield planned changes to D93135: [libomptarget][devicertl] Port amdgcn devicertl to openmp.

Most (possibly all except the amdgpu cmake) of this change is incorporated in D94745 or already committed to trunk. I'm going to leave this up as a reference for now but don't intend to land it.

Fri, Jan 22, 7:37 AM · Restricted Project
JonChesterfield added inline comments to D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Fri, Jan 22, 7:35 AM · Restricted Project
JonChesterfield committed rG9b19ecb8f1ec: [libomptarget][devicertl] Drop templated atomic functions (authored by JonChesterfield).
[libomptarget][devicertl] Drop templated atomic functions
Fri, Jan 22, 6:49 AM
JonChesterfield closed D95093: [libomptarget][devicertl] Drop templated atomic functions.
Fri, Jan 22, 6:49 AM · Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Fri, Jan 22, 4:16 AM · Restricted Project, Restricted Project
JonChesterfield added a comment to D95093: [libomptarget][devicertl] Drop templated atomic functions.

I'd like to call the clang/llvm intrinsics instead, plus replace unsigned long long with uint64_t. Easier to check the correctness of that after the call sites are limited to one file.

Fri, Jan 22, 4:08 AM · Restricted Project
JonChesterfield added inline comments to D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Fri, Jan 22, 3:12 AM · Restricted Project
JonChesterfield added inline comments to D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Fri, Jan 22, 3:11 AM · Restricted Project

Thu, Jan 21

JonChesterfield updated the diff for D95093: [libomptarget][devicertl] Drop templated atomic functions.
  • rebase
Thu, Jan 21, 4:20 PM · Restricted Project
JonChesterfield accepted D95161: [Clang][OpenMP][NVPTX] Replace `libomptarget-nvptx-path` with `libomptarget-nvptx-bc-path`.

LGTM, thanks!

Thu, Jan 21, 4:17 PM · Restricted Project
JonChesterfield retitled D95155: [libomptarget] Build cuda plugin without cuda installed locally from [libomptarget] WIP build plugin without cuda.h to [libomptarget] Build cuda plugin without cuda installed locally.
Thu, Jan 21, 3:53 PM · Restricted Project
JonChesterfield updated the diff for D95155: [libomptarget] Build cuda plugin without cuda installed locally.
  • fix copy/paste error, support higher arity functions
  • rebase, tidy
Thu, Jan 21, 3:51 PM · Restricted Project
JonChesterfield updated the summary of D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Thu, Jan 21, 3:38 PM · Restricted Project
JonChesterfield updated the diff for D95155: [libomptarget] Build cuda plugin without cuda installed locally.
  • fix copy/paste error, support higher arity functions
Thu, Jan 21, 3:32 PM · Restricted Project
JonChesterfield added a comment to D95155: [libomptarget] Build cuda plugin without cuda installed locally.

Interesting. Should we do this instead of the pure macro stuff? Is this going to break if the argument type doesn't match the declaration parameter type, e.g.,

void f(int);

f(0.0);
Thu, Jan 21, 3:25 PM · Restricted Project
JonChesterfield added a comment to D94855: [OpenMP] Add time profiling support in libomp.

This breaks the build for me. Compiling openmp via ENABLE_RUNTIMES without any extra options, the include path isn't set up so kmp_runtime.cpp fails to find the header.

Thu, Jan 21, 11:48 AM · Restricted Project
JonChesterfield added inline comments to D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Thu, Jan 21, 11:23 AM · Restricted Project
JonChesterfield requested review of D95155: [libomptarget] Build cuda plugin without cuda installed locally.
Thu, Jan 21, 11:21 AM · Restricted Project

Wed, Jan 20

JonChesterfield accepted D95085: [OpenMP][NVPTX] Added forward declaration for atomic operations.

Thanks. If you land this I'll rebase D95093 manually

Wed, Jan 20, 3:53 PM · Restricted Project
JonChesterfield added a comment to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.

Yeah... That's not gone well. Unless @jdoerfert can point us at a better way to spell these in openmp, let's add 'fix target atomics' to the todo pile and use clang intrinsics for the time being.

Wed, Jan 20, 3:10 PM · Restricted Project, Restricted Project
JonChesterfield added a comment to D95093: [libomptarget][devicertl] Drop templated atomic functions.

Most of the intrinsic functions are defined on a single type, e.g. if we wanted to use the nvvm encoded one for atomic_inc.

Wed, Jan 20, 2:57 PM · Restricted Project
JonChesterfield added a comment to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.

#pragma omp declare target

Wed, Jan 20, 2:51 PM · Restricted Project, Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Wed, Jan 20, 2:44 PM · Restricted Project, Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Wed, Jan 20, 2:38 PM · Restricted Project, Restricted Project
JonChesterfield added inline comments to D95093: [libomptarget][devicertl] Drop templated atomic functions.
Wed, Jan 20, 2:22 PM · Restricted Project
JonChesterfield requested review of D95093: [libomptarget][devicertl] Drop templated atomic functions.
Wed, Jan 20, 2:19 PM · Restricted Project
JonChesterfield added inline comments to D95085: [OpenMP][NVPTX] Added forward declaration for atomic operations.
Wed, Jan 20, 2:12 PM · Restricted Project
JonChesterfield added a comment to D95085: [OpenMP][NVPTX] Added forward declaration for atomic operations.

Possibly obsolete?

Wed, Jan 20, 1:14 PM · Restricted Project
JonChesterfield added a comment to D95062: [libomptarget][devicertl][nfc] Simplify target_atomic abstraction.

LG, can you replace it with clang or OpenMP atomics next?

Wed, Jan 20, 11:59 AM · Restricted Project
JonChesterfield committed rGfbc1dcb94655: [libomptarget][devicertl][nfc] Simplify target_atomic abstraction (authored by JonChesterfield).
[libomptarget][devicertl][nfc] Simplify target_atomic abstraction
Wed, Jan 20, 11:51 AM
JonChesterfield closed D95062: [libomptarget][devicertl][nfc] Simplify target_atomic abstraction.
Wed, Jan 20, 11:51 AM · Restricted Project
JonChesterfield committed rGea616f9026dc: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify (authored by JonChesterfield).
[libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify
Wed, Jan 20, 11:45 AM
JonChesterfield closed D95060: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify.
Wed, Jan 20, 11:45 AM · Restricted Project
JonChesterfield updated the diff for D95060: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify.
  • clang-format
Wed, Jan 20, 11:44 AM · Restricted Project
JonChesterfield added inline comments to D95060: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify.
Wed, Jan 20, 10:53 AM · Restricted Project
JonChesterfield added inline comments to D95060: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify.
Wed, Jan 20, 10:06 AM · Restricted Project
JonChesterfield requested review of D95062: [libomptarget][devicertl][nfc] Simplify target_atomic abstraction.
Wed, Jan 20, 10:01 AM · Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Wed, Jan 20, 9:47 AM · Restricted Project, Restricted Project
JonChesterfield accepted D95058: [OpenMP][NVPTX] Added forward declaration to pave the way for building deviceRTLs with OpenMP.
Wed, Jan 20, 9:46 AM · Restricted Project
JonChesterfield added inline comments to D95060: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify.
Wed, Jan 20, 9:45 AM · Restricted Project
JonChesterfield requested review of D95060: [libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify.
Wed, Jan 20, 9:41 AM · Restricted Project
JonChesterfield committed rGe069662deb1f: [libomptarget][devicertl] Wrap source in declare target pragmas (authored by JonChesterfield).
[libomptarget][devicertl] Wrap source in declare target pragmas
Wed, Jan 20, 7:51 AM
JonChesterfield closed D95048: [libomptarget][devicertl] Wrap source in declare target pragmas.
Wed, Jan 20, 7:51 AM · Restricted Project
JonChesterfield added inline comments to D95048: [libomptarget][devicertl] Wrap source in declare target pragmas.
Wed, Jan 20, 7:50 AM · Restricted Project
JonChesterfield updated the diff for D95048: [libomptarget][devicertl] Wrap source in declare target pragmas.
  • insert whitespace requested at review
Wed, Jan 20, 7:49 AM · Restricted Project
JonChesterfield added a comment to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.

Let's split this. Declare target in all the source files can go first, doesn't hurt anyone.

D95048

Wed, Jan 20, 7:20 AM · Restricted Project, Restricted Project
JonChesterfield requested review of D95048: [libomptarget][devicertl] Wrap source in declare target pragmas.
Wed, Jan 20, 7:19 AM · Restricted Project
JonChesterfield accepted D95013: [OpenMP][NVPTX] Replaced CUDA builtin vars with LLVM intrinsics.

Nice. Particularly like getting rid of these because they use field access syntax to make a function call, which clang does with a Microsoft extension.

Wed, Jan 20, 2:53 AM · Restricted Project

Tue, Jan 19

JonChesterfield added inline comments to D95007: [CUDA][HIP] Add -fuse-cuid.
Tue, Jan 19, 5:04 PM
JonChesterfield added a reviewer for D95007: [CUDA][HIP] Add -fuse-cuid: jdoerfert.

@jdoerfert this may be a candidate for deriving names of extracted target regions

Tue, Jan 19, 4:59 PM
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Tue, Jan 19, 4:39 PM · Restricted Project, Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Tue, Jan 19, 2:45 PM · Restricted Project, Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Tue, Jan 19, 2:04 PM · Restricted Project, Restricted Project
JonChesterfield added a comment to D94884: [Clang][OpenMP] Include header for CUDA builtin vars into OpenMP wrapper header.

I don't think introducing everything from the cuda namespace into openmp nvptx offloading is a feature. Inevitably people will call threadIdx.x instead of the openmp or clang equivalent, and this will mask missing functionality in openmp.

Tue, Jan 19, 1:55 PM · Restricted Project
JonChesterfield added a comment to D94871: [Clang][OpenMP] Fixed an issue that clang crashed when compiling OpenMP program in device only mode without host IR.

Nice! Thank you

Tue, Jan 19, 1:40 PM · Restricted Project, Restricted Project
JonChesterfield added inline comments to D94648: [amdgpu] Implement lower function LDS pass.
Tue, Jan 19, 8:27 AM · Restricted Project
JonChesterfield updated the diff for D94648: [amdgpu] Implement lower function LDS pass.
  • address some review comments
  • Run tests under new pass manager
Tue, Jan 19, 8:20 AM · Restricted Project
JonChesterfield added a comment to D94961: [OpenMP] Add OpenMP offloading toolchain for AMDGPU.

This patch was written, roughly, by:

  • copying the known-working openmp driver from rocm into the trunk source tree
  • deleting lots of stuff that didn't look necessary
  • deleting some stuff that is broadly necessary, but the specifics are up for debate
Tue, Jan 19, 4:12 AM · Restricted Project

Mon, Jan 18

JonChesterfield added a comment to D94648: [amdgpu] Implement lower function LDS pass.

Fixed the easy parts. Adding a call to the new pass manager by RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-function-lds < %s | FileCheck %s does't work - it looks like the pass manager runs the module pass on individual functions. Trying to find some documentation about the new pass manager to see where the plumbing error lies.

Mon, Jan 18, 9:02 AM · Restricted Project
JonChesterfield updated the diff for D94648: [amdgpu] Implement lower function LDS pass.
  • address some review comments
Mon, Jan 18, 9:00 AM · Restricted Project
JonChesterfield updated subscribers of D77609: [OpenMP] Added the support for hidden helper task in RTL.

@ronlieb tells me an out of tree offloading test (aomp/test/smoke/devices) started crashing (hangs/segv/fp exception) with this patch applied. That doesn't make sense to me since this doesn't appear to change the target offloading logic, but it might be a smoking gun for a lifetime management error somewhere in the above. Does anyone know if the host openmp runtime is expected to be clean under things like valgrind or thread sanitizer?

Mon, Jan 18, 6:06 AM · Restricted Project

Fri, Jan 15

JonChesterfield added a comment to D94648: [amdgpu] Implement lower function LDS pass.

Great review guys, thanks! With the exception of avoiding inline asm (which I'd like to, but am not sure what to do about) I'll fix the rest when I'm back in the office.

Fri, Jan 15, 11:43 AM · Restricted Project
JonChesterfield added inline comments to D94745: [OpenMP][WIP] Build the deviceRTLs with OpenMP instead of target dependent language - NOT FOR REVIEW.
Fri, Jan 15, 2:19 AM · Restricted Project, Restricted Project

Thu, Jan 14

JonChesterfield committed rG214387c2c694: [libomptarget][nvptx] Reduce calls to cuda header (authored by JonChesterfield).
[libomptarget][nvptx] Reduce calls to cuda header
Thu, Jan 14, 6:17 PM
JonChesterfield closed D94731: [libomptarget][nvptx] Reduce calls to cuda header.
Thu, Jan 14, 6:16 PM · Restricted Project
JonChesterfield added a comment to D94731: [libomptarget][nvptx] Reduce calls to cuda header.

Remaining calls could be replaced with builtins by duplicating parts of
the cuda wrapper infra. This would mean choosing which to call based
on architecture number (>= sm_70) instead of CUDA_VERSION.

Thu, Jan 14, 6:14 PM · Restricted Project
JonChesterfield updated the summary of D94731: [libomptarget][nvptx] Reduce calls to cuda header.
Thu, Jan 14, 6:14 PM · Restricted Project
JonChesterfield retitled D94731: [libomptarget][nvptx] Reduce calls to cuda header from [libomptarget][nvptx] Call builtins instead of cuda to [libomptarget][nvptx] Reduce calls to cuda header.
Thu, Jan 14, 6:11 PM · Restricted Project
JonChesterfield updated the diff for D94731: [libomptarget][nvptx] Reduce calls to cuda header.
  • reduce scope
Thu, Jan 14, 6:09 PM · Restricted Project
JonChesterfield added a comment to D94731: [libomptarget][nvptx] Reduce calls to cuda header.

That's interesting. Move some of the current target_impl.h into clang shipped headers.

Thu, Jan 14, 5:28 PM · Restricted Project
JonChesterfield added a comment to D94731: [libomptarget][nvptx] Reduce calls to cuda header.

Note that this is incremental (on the basis that it's already hard enough to review). Cuda.h is still used for CUDA_VERSION here, but also for the atomic functions and a few libc prototypes. I have the library compiling without it locally.

Thu, Jan 14, 5:08 PM · Restricted Project
JonChesterfield requested review of D94731: [libomptarget][nvptx] Reduce calls to cuda header.
Thu, Jan 14, 4:31 PM · Restricted Project
JonChesterfield committed rG6e7094c14b22: [libomptarget][nvptx][nfc] Move target_impl functions out of header (authored by JonChesterfield).
[libomptarget][nvptx][nfc] Move target_impl functions out of header
Thu, Jan 14, 4:20 PM
JonChesterfield closed D94728: [libomptarget][nvptx][nfc] Move target_impl functions out of header.
Thu, Jan 14, 4:20 PM · Restricted Project