1. Nov 11, 2023
  2. Nov 10, 2023
    • Nikita Popov's avatar
      [InstCombine] Avoid use of shift constant expressions (NFCI) · b43b2a64
      Nikita Popov authored
      Use the constant folding API instead. As we're working on
      ImmConstants, these folds are guaranteed to succeed.
      b43b2a64
    • Guray Ozen's avatar
      [mlir] Add sm_90a GEMM test 128x128x128 (F32 += F16 * F16) (#69913) · 51916f0c
      Guray Ozen authored
      This PR adds a test that performs GEMM 128x128x128 (F32 += F16 * F16).
      It uses `sm_90a` features in NVGPU dialect.
      
      Simplified algorithm is as follows:
      
      **Prologue** 
      ```
      mgroup = mbarriers.init x 2
      tma.load ... shmem_buffer_lhs<0 x 128 x 64>
      tma.load ... shmem_buffer_rhs<0 x 64 x 64>
      tma.load ... shmem_buffer_rhs<0 x 64 x 64>
      mbarrier.expect_tx 32768
      tma.load ... shmem_buffer_lhs<1 x 128 x 64>
      tma.load ... shmem_buffer_rhs<1 x 64 x 64>
      tma.load ... shmem_buffer_rhs<1 x 64 x 64>
      mbarrier.expect_tx 32768
      ```
      **Mainloop**
      ```
      matrixD = 
       for(i = 0;...2) {   
         mbarrier.try_wait [i]
         lhs = shmem_buffer_lhs<pipe x 128 x 64>
         rhs = shmem_buffer_rhs<pipe x 64 x 128>
         yield nvgpu.warpgroup.mma (lhs, rhs)
      
      //   Expanded : nvgpu.warpgroup.mma [128][128]+=[128][64]*[64][128]
      //                  wgmma.m64n128k16(A[0:64][0:16]  *  B[0:16][0:128])
      //                  wgmma.m64n128k16(A[0:64][16:32] *  B[16:32][0:128])
      //                  wgmma.m64n128k16(A[0:64][32:48] *  B[32:48][0:128])
      //                  wgmma.m64n128k16(A[0:64][48:64] *  B[48:64][0:128])
      //                  wgmma.m64n128k16(A[64:128][0:16]  *  B[0:16][0:128])
      //                  wgmma.m64n128k16(A[64:128][16:32] *  B[16:32][0:128])
      //                  wgmma.m64n128k16(A[64:128][32:48] *  B[32:48][0:128])
      //                  wgmma.m64n128k16(A[64:128][48:64] *  B[48:64][0:128])
      ```
      
      **Epilogue** 
      ```
      //reg->shmem
      warpgroup.mma.store matrixD, shmem
      //shmem->glbmem
      parallel-for(i=0;...128)
       parallel-for(j=0;...128)
         store shmem, globalmem
      ```
      51916f0c
    • Guray Ozen's avatar
      [mlir] Add sm_90a GEMM test 128x128x128 (F32 =F16*F16) with predicate (#70028) · a00caad6
      Guray Ozen authored
      PR #69913 added a GEMM test (128x128x128 F32 += F16 * F16) with
      if-statement. This PR adds the same test using predicates in PTX.
      Predicate support is enabled using _BasicPtxBuilderInterface_
      `(nvgpu.opcode ..., predicate = %pred)`.
      
      The predicate condition is computed in `Step 2. [GPU] Elect fastest
      thread in CTA` inspired by cutlass. It is as follows:
      ```
      lane_predicate = nvvm.elect.sync
      warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0)
      warp_idx_in_warp_group = warp_idx % 4
      predicate = (lane_predicate & warp_idx_in_warp_group)
      ```
      
      Depends on #70027 #69934 #69935 #69584
      a00caad6
    • Anna Thomas's avatar
      4ba50a78
    • Guray Ozen's avatar
      [mlir] Fix sm90 test for new verifier · f4d59522
      Guray Ozen authored
      #70923 improved verifier. The verifier caught that the tensor map type in the tma descriptor in this test isn't correct. The program was working correctly anway since the offset is calculated correctly.
      
      This work fixes the test.
      f4d59522
    • Guillaume Chatelet's avatar
      [libc] Update configure.rst after config.json modification (#71942) · c07f73e7
      Guillaume Chatelet authored
      The update is automatically generated from `config/config.json`.
      c07f73e7
    • Joseph Huber's avatar
      [AMDGPU] Call the `FINI_ARRAY` destructors in the correct order (#71815) · c1d5865a
      Joseph Huber authored
      Summary:
      The AMDGPU backend uses the linker-provided INIT_ARRAY and FINI_ARRAY
      sections to call all the global constructors in a single kernel.
      Previously this mistakenly used the same iteration logic for both
      arrays. The destructors stored in FINI_ARRAY are stored in the same
      order as
      the ones in the INIT_ARRAY section so we need to traverse it in reverse
      order.
      c1d5865a
    • Joseph Huber's avatar
      [NVPTX] Allow the ctor/dtor lowering pass to emit kernels (#71549) · af8ebfdc
      Joseph Huber authored
      Summary:
      This pass emits the new "nvptx$device$init" and "nvptx$device$fini"
      kernels that are callable by the device. This intends to mimic the
      method of lowering for AMDGPU where we emit `amdgcn.device.init` and
      `amdgcn.device.fini` respectively. These kernels simply iterate a symbol
      called `__init_array_start/stop` and `__fini_array_start/stop`.
      Normally, the linker provides these symbols automatically. In the AMDGPU
      case we only need call the kernel and we call the ctors / dtors.
      However, for NVPTX we require the user initializes these variables to
      the associated globals that we already emit as a part of this pass.
      
      The motivation behind this change is to move away from OpenMP's handling
      of ctors / dtors. I would much prefer that the backend / runtime handles
      this. That allows us to handle ctors / dtors in a language agnostic way,
      
      This approach requires that the runtime initializes the associated
      globals. They are marked `weak` so we can emit this per-TU. The kernel
      itself is `weak_odr` as it is copied exactly.
      
      One downside is that any module containing these kernels elicitis the
      "stack size cannot be statically determined warning" every time from
      `nvlink` which is annoying but inconsequential for functionality. It
      would be nice if there were a way to silence this warning however.
      af8ebfdc
    • Nikita Popov's avatar
      [Clang] Add missing REQUIRES to tests (NFC) · 3a9cc17c
      Nikita Popov authored
      3a9cc17c
    • Nikita Popov's avatar
      [IR] Mark lshr and ashr constant expressions as undesirable · 82f68a99
      Nikita Popov authored
      These will no longer be created by default during constant folding.
      82f68a99
    • Nikita Popov's avatar
      [Clang] Generate test checks (NFC) · cd7ba9f3
      Nikita Popov authored
      cd7ba9f3
    • Nikita Popov's avatar
      [InstCombine] Require immediate constant in canEvaluateShifted() · 707bb421
      Nikita Popov authored
      Otherwise we risk infinite loops when shift constant expressions
      are no longer supported.
      707bb421
    • Quinn Dawkins's avatar
      [mlir][vector] Add distribution pattern for vector.create_mask (#71619) · d4d28914
      Quinn Dawkins authored
      This is the last step needed for basic support for distributing masked
      vector code. The lane id gets delinearized based on the distributed mask
      shape and then compared against the original mask sizes to compute the
      bounds for the distributed mask. Note that the distribution of masks is
      implicit on the shape specified by the warp op. As a result, it is the
      responsibility of the consumer of the mask to ensure the distributed
      mask will match its own distribution semantics.
      d4d28914
    • Jan Svoboda's avatar
      [clang][deps] Skip writing `DIAG_PRAGMA_MAPPINGS` record (#70874) · 22c68511
      Jan Svoboda authored
      Following up on #69975, this patch skips writing `DIAG_PRAGMA_MAPPINGS`
      as well. Deserialization of this PCM record is still showing up in
      profiles, since it needs to be VBR-decoded for every transitively loaded
      PCM file.
      
      The scanner doesn't make any guarantees about diagnostic accuracy (and
      it even disables all warnings), so skipping this record should be safe.
      22c68511
    • Nikita Popov's avatar
      [InstCombine] Avoid uses of ConstantExpr::getLShr() · 8391f405
      Nikita Popov authored
      Use the constant folding API instead.
      8391f405
    • Nikita Popov's avatar
      [InstCombine] Avoid some uses of ConstantExpr::getLShr() (NFC) · eb5199e8
      Nikita Popov authored
      Use the constant folding API instead. As we're working on
      ImmConstant, it is guaranteed to succeed.
      eb5199e8
    • diggerlin's avatar
      [Driver][LTO] Copy fix empty stats filename to AIX (#71738) · d51dd892
      diggerlin authored
      copy implement of https://github.com/llvm/llvm-project/pull/71359
      
       to
      AIX.cpp and add test scenario
      
      ---------
      
      Co-authored-by: default avatarzhijian <zhijian@ca.ibm.com>
      d51dd892
    • Sam Tebbs's avatar
      [Clang][SME2] Add single and multi min and max builtins (#71688) · b008d66c
      Sam Tebbs authored
      Adds the following SME2 builtins:
      
          svmin_single_(s8|s16|s32|s64)_x(2|4)
          svmin_single_(u8|u16|u32|u64)_x(2|4)
          svmin_single_(f16|f32|f64)_x(2|4)
          svmin_(s8|s16|s32|s64)_x(2|4)
          svmin_(u8|u16|u32|u64)_x(2|4)
          svmin_(f16|f32|f64)_x(2|4)
          svmax_single_(s8|s16|s32|s64)_x(2|4)
          svmax_single_(u8|u16|u32|u64)_x(2|4)
          svmax_single_(f16|f32|f64)_x(2|4)
          svmax_(s8|s16|s32|s64)_x(2|4)
          svmax_(u8|u16|u32|u64)_x(2|4)
          svmax_(f16|f32|f64)_x(2|4)
      
      See
      [https://github.com/ARM-software/acle/pull/217](https://github.com/ARM-software/acle/pull/217)
      
      Patch by: Kerry McLaughlin <kerry.mclaughlin@arm.com>
      b008d66c
    • Quinn Dawkins's avatar