Page MenuHomePhabricator

saiislam (Saiyedul Islam)
User

Projects

User does not belong to any projects.

User Details

User Since
Nov 15 2019, 2:59 AM (28 w, 1 d)

Recent Activity

Yesterday

saiislam added a comment to D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics.

Actually, the question really is about why inc/dec are needed as separate operations either as IR intrinsics or Clang builtins. Why not just expose a __builtin_amdgcn_atomicrmw that takes a scope, and map it to the LLVM atomicrmw? That would be way cleaner. The language can provide convenience functions for inc/dec that internally call the rmw builtin.

Fri, May 29, 10:53 AM · Restricted Project
saiislam retitled D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics from [AMDGPU] Expose llvm atomic inc/dec instructions as clang builtins for AMDGPU target to [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics.
Fri, May 29, 10:53 AM · Restricted Project
saiislam updated the diff for D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics.
  1. Updated title and description.
  2. Replaced pointer element type usage to i32 type for getIntrinsic.
  3. Volatile argument of the instrinsic now comes from pointer type of the first argument of the builtin.
  4. Updated all test cases wrt above change.
  5. Added test case for volatile pointer type.
Fri, May 29, 10:53 AM · Restricted Project
saiislam added parent revisions for D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics: D75917: Expose llvm fence instruction as clang intrinsic, D73076: [libomptarget] Implement most hip atomic functions in terms of intrinsics.
Fri, May 29, 8:06 AM · Restricted Project
saiislam added a child revision for D73076: [libomptarget] Implement most hip atomic functions in terms of intrinsics: D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics.
Fri, May 29, 8:06 AM · Restricted Project
saiislam added a child revision for D75917: Expose llvm fence instruction as clang intrinsic: D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics.
Fri, May 29, 8:06 AM · Restricted Project
saiislam created D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics.
Fri, May 29, 8:06 AM · Restricted Project

Wed, May 27

saiislam committed rG602d9b0afc77: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1 (authored by saiislam).
[OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1
Wed, May 27, 1:03 AM
saiislam closed D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.
Wed, May 27, 1:02 AM · Restricted Project, Restricted Project

Tue, May 26

saiislam added inline comments to D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.
Tue, May 26, 10:49 AM · Restricted Project, Restricted Project
saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Shifted test cases in openmp-offload-gpu.c for better visual segmentation.
Updated device function test case to be more accurate.

Tue, May 26, 5:22 AM · Restricted Project, Restricted Project

Fri, May 22

saiislam added inline comments to D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.
Fri, May 22, 6:23 AM · Restricted Project, Restricted Project
saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Added test case to show treatment of specific functions as builtins or functions on the device

Fri, May 22, 5:51 AM · Restricted Project, Restricted Project

Mon, May 18

saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Fixed typo. Added test case to check for device side functions.

Mon, May 18, 6:58 PM · Restricted Project, Restricted Project

Sun, May 17

saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Removed isOpenMPGPU() to avoid defining OpenMP compatibility of an architecture.
Reverting back to explicitly checking NVPTX and AMDGCN architectures. Also, split
handling of NVPTX's and AMDGCN's handling of getBuiltinID. For AMDGCN it now uses
OpenMPIsDevice LangOpt and returns 0 for every device library function, except for
printf and malloc.

Sun, May 17, 10:20 PM · Restricted Project, Restricted Project

Fri, May 15

saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Moved isGPU() from llvm's Triple.h to clang's TargetInfo. Renamed it to isOpenMPGPU()
to represent target's compatibility with OpenMP offloading and reduce its scope.

Fri, May 15, 3:02 AM · Restricted Project, Restricted Project

Thu, May 14

saiislam added a reviewer for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1: arsenm.
Thu, May 14, 9:10 AM · Restricted Project, Restricted Project
saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Added test cases. Added a wrapper isGPU() for isNVPTX()/isAMDGCN().

Thu, May 14, 5:20 AM · Restricted Project, Restricted Project
saiislam added inline comments to D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.
Thu, May 14, 5:20 AM · Restricted Project, Restricted Project

Mon, May 11

saiislam added reviewers for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1: sameerds, msearles, hliao.
Mon, May 11, 6:23 PM · Restricted Project, Restricted Project
saiislam committed rG117e5609e98b: [AMDGPU] Reserving VGPR for future SGPR Spill (authored by saiislam).
[AMDGPU] Reserving VGPR for future SGPR Spill
Mon, May 11, 5:51 PM
saiislam closed D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.
Mon, May 11, 5:50 PM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.

Fixed lint errors.

Mon, May 11, 5:18 PM · Restricted Project, Restricted Project
saiislam created D79754: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 1.
Mon, May 11, 5:18 PM · Restricted Project, Restricted Project
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Removed a break

Mon, May 11, 11:50 AM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Switched the nested loops in lowerShiftReservedVGPR() to move
sorting of unique live-ins to the outer loop. Improved cleanup of
reserved vgpr after lowering it to a lower value, and when it needs
to be removed.

Mon, May 11, 11:18 AM · Restricted Project, Unknown Object (Project)

Fri, May 8

saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Handled a test case which was failing in debug only.

Fri, May 8, 3:36 PM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Removed unsued frame index argument from vgpr reserving function. Added condition to reserve vgpr only if machine function has stack objects.

Fri, May 8, 1:26 PM · Restricted Project, Unknown Object (Project)

Thu, May 7

saiislam added inline comments to D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.
Thu, May 7, 10:47 AM · Restricted Project, Unknown Object (Project)
saiislam added a reviewer for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill: cdevadas.
Thu, May 7, 3:21 AM · Restricted Project, Unknown Object (Project)

Wed, May 6

saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Cleaned code for reserving vgpr.

Wed, May 6, 5:38 PM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Fixed failing test cases. Now VGPR reservation happens in an
independent function and doesn't overload allocateSGPRSpillToVGPR.
Updated test case to hard code usage of all but one VGPRs and use
it for SGPR spill.

Wed, May 6, 3:19 PM · Restricted Project, Unknown Object (Project)

Tue, May 5

saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Handled few seg faults in the lit-tests.

Tue, May 5, 5:20 AM · Restricted Project, Unknown Object (Project)

Sun, May 3

saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Removed commented code.

Sun, May 3, 9:12 PM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Rebased and added code for proper lower shifting of reserved VGPR.

Sun, May 3, 8:46 PM · Restricted Project, Unknown Object (Project)

Apr 24 2020

saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Updated description and added a failing test case for integer scope.

Apr 24 2020, 6:58 PM · Restricted Project
saiislam updated the summary of D75917: Expose llvm fence instruction as clang intrinsic.
Apr 24 2020, 5:53 PM · Restricted Project

Apr 23 2020

saiislam added inline comments to D75917: Expose llvm fence instruction as clang intrinsic.
Apr 23 2020, 6:27 AM · Restricted Project
saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Added check and test for sync scope to be a constant literal.

Apr 23 2020, 6:27 AM · Restricted Project

Apr 22 2020

saiislam updated the summary of D75917: Expose llvm fence instruction as clang intrinsic.
Apr 22 2020, 11:18 PM · Restricted Project
saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Removed documentation from clang doc. Squashed all changes into a single commit.

Apr 22 2020, 11:18 PM · Restricted Project
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Now highest available VGPR is reserved in the beginning and it is
switched with lowest available VGPR after RA.

Apr 22 2020, 3:15 PM · Restricted Project, Unknown Object (Project)
saiislam added inline comments to D75917: Expose llvm fence instruction as clang intrinsic.
Apr 22 2020, 8:38 AM · Restricted Project
saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Moved builtin-amdgcn-fence-failure.cpp from clang/test/CodeGenCXX/
to clang/test/Sema/ since it is checking an error.

Apr 22 2020, 8:38 AM · Restricted Project
saiislam updated the summary of D75917: Expose llvm fence instruction as clang intrinsic.
Apr 22 2020, 4:49 AM · Restricted Project
saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Removed stale commented code

Apr 22 2020, 4:49 AM · Restricted Project
saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Changed the builtin to be AMDGCN-specific

Apr 22 2020, 4:49 AM · Restricted Project

Apr 17 2020

saiislam updated the summary of D75917: Expose llvm fence instruction as clang intrinsic.
Apr 17 2020, 10:47 AM · Restricted Project
saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.
  1. Moved test case to clang/test/CodeGenCXX.
  2. Added a failing test case with invalid sync scope, which gets detected by implementation of fence instruction.
  3. Updated the change description of the builtin.
  4. Updated the clang documentation describing mapping of C++ memory-ordering to LLVM memory-ordering.
Apr 17 2020, 10:47 AM · Restricted Project

Apr 10 2020

saiislam added inline comments to D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.
Apr 10 2020, 10:05 PM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Removed reserved VGPR from RA list. Added a test case which
ensures spilling of SGPR into reserved VGPR.

Apr 10 2020, 3:44 AM · Restricted Project, Unknown Object (Project)

Apr 6 2020

saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Changes:

  1. Moved builtin definition with rest of atomic builtins
  2. Updated validity checking of memory order using exact mathches instead of range matching
  3. Added a sucessful test case which passes arbitrary string scope
  4. Corrected formatting
Apr 6 2020, 4:50 AM · Restricted Project
saiislam added inline comments to D75917: Expose llvm fence instruction as clang intrinsic.
Apr 6 2020, 3:45 AM · Restricted Project

Apr 5 2020

saiislam updated the diff for D75917: Expose llvm fence instruction as clang intrinsic.

Removed OpenCL specific dependencies

Apr 5 2020, 10:08 AM · Restricted Project

Mar 10 2020

saiislam created D75917: Expose llvm fence instruction as clang intrinsic.
Mar 10 2020, 6:24 AM · Restricted Project

Nov 18 2019

saiislam updated the summary of D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.
Nov 18 2019, 2:56 AM · Restricted Project, Unknown Object (Project)
saiislam updated the diff for D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.

Corrected formatting

Nov 18 2019, 2:56 AM · Restricted Project, Unknown Object (Project)
saiislam created D70379: [AMDGPU] Reserving VGPR for future SGPR Spill.
Nov 18 2019, 2:01 AM · Restricted Project, Unknown Object (Project)