[all-commits] [llvm/llvm-project] b6ef36: [AMDGPU] Update WMMA intrinsics with explicit f16 ...

[Piotr Sobczak via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: b6ef36a1c427d07116fea84623b7caa37d8a7d7b
      https://github.com/llvm/llvm-project/commit/b6ef36a1c427d07116fea84623b7caa37d8a7d7b
  Author: Piotr Sobczak <Piotr.Sobczak at amd.com>
  Date:   2022-07-01 (Fri, 01 Jul 2022)

  Changed paths:
    M llvm/include/llvm/IR/IntrinsicsAMDGPU.td
    M llvm/lib/Target/AMDGPU/SIInstrInfo.td
    M llvm/lib/Target/AMDGPU/VOP3PInstructions.td
    M llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.wmma_32.ll
    M llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.wmma_64.ll
    M llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wmma_32.ll
    M llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wmma_64.ll
    M llvm/test/CodeGen/AMDGPU/wmma_multiple_32.ll
    M llvm/test/CodeGen/AMDGPU/wmma_multiple_64.ll

  Log Message:
  -----------
  [AMDGPU] Update WMMA intrinsics with explicit f16 types

Update intrinsics to use n x f16 and n x i16 instead
of 32-bit types. This may avoid the need for a bitcast
and is probably less confusing.

Depends on making v16f16 and v16i16 types legal.

Reviewed By: rampitec

Differential Revision: https://reviews.llvm.org/D128951


  Commit: 4a782252127761b60d33e74f9d9acb0aad6f742f
      https://github.com/llvm/llvm-project/commit/4a782252127761b60d33e74f9d9acb0aad6f742f
  Author: Piotr Sobczak <Piotr.Sobczak at amd.com>
  Date:   2022-07-01 (Fri, 01 Jul 2022)

  Changed paths:
    M clang/include/clang/Basic/BuiltinsAMDGPU.def
    M clang/lib/CodeGen/CGBuiltin.cpp
    A clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32.cl
    A clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl

  Log Message:
  -----------
  [AMDGPU] Add WMMA clang builtins

Add WMMA clang builtins and tests. Extra changes in code
are needed to handle function overloads.

WavefrontSize 32:
__builtin_amdgcn_wmma_f32_16x16x16_f16_w32
__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32
__builtin_amdgcn_wmma_f16_16x16x16_f16_w32
__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32
__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32
__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32

WavefrontSize 64:
__builtin_amdgcn_wmma_f32_16x16x16_f16_w64
__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64
__builtin_amdgcn_wmma_f16_16x16x16_f16_w64
__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64
__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64
__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64

Reviewed By: rampitec

Differential Revision: https://reviews.llvm.org/D128952


Compare: https://github.com/llvm/llvm-project/compare/92a3e1b5c9e7...4a7822521277