projects / platform / upstream / llvm.git / commit
? search:
summary | shortlog | log | commit | commitdiff | tree
(parent: b6ef36a) | patch
[AMDGPU] Add WMMA clang builtins
authorPiotr Sobczak <Piotr.Sobczak@amd.com>
Fri, 1 Jul 2022 06:18:09 +0000 (08:18 +0200)
committerPiotr Sobczak <Piotr.Sobczak@amd.com>
Fri, 1 Jul 2022 06:55:25 +0000 (08:55 +0200)
commit4a782252127761b60d33e74f9d9acb0aad6f742f
tree7c0979c82bffb031c64226c1527bb6cc2a3e4c70tree | snapshot
parentb6ef36a1c427d07116fea84623b7caa37d8a7d7bcommit | diff
[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
clang/include/clang/Basic/BuiltinsAMDGPU.def diff | blob | history
clang/lib/CodeGen/CGBuiltin.cpp diff | blob | history
clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32.cl [new file with mode: 0644] blob
clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.cl [new file with mode: 0644] blob
Domain: System / Toolchain;