mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-25 17:36:06 +00:00
58dada5f0a
Add a few tests to make sure we get the expected instruction for the WMMA builtins (and generally that our builtins and intrinsics are on the same page and won't blow up). Differential Revision: https://reviews.llvm.org/D144176
73 lines
2.7 KiB
Common Lisp
73 lines
2.7 KiB
Common Lisp
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
|
|
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize64 -DWMMA_GFX1100_TESTS -S -o - %s | FileCheck %s --check-prefix=CHECK-GFX1100
|
|
|
|
typedef float v4f __attribute__((ext_vector_type(4)));
|
|
typedef float v8f __attribute__((ext_vector_type(8)));
|
|
typedef half v8h __attribute__((ext_vector_type(8)));
|
|
typedef half v16h __attribute__((ext_vector_type(16)));
|
|
typedef int v2i __attribute__((ext_vector_type(2)));
|
|
typedef int v4i __attribute__((ext_vector_type(4)));
|
|
typedef int v8i __attribute__((ext_vector_type(8)));
|
|
typedef short v8s __attribute__((ext_vector_type(8)));
|
|
typedef short v16s __attribute__((ext_vector_type(16)));
|
|
|
|
#ifdef WMMA_GFX1100_TESTS
|
|
|
|
// Wave64
|
|
|
|
|
|
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_f16_w64:
|
|
// CHECK-GFX1100: v_wmma_f32_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
|
|
//
|
|
void test_amdgcn_wmma_f32_16x16x16_f16_w64(global v4f* out, v16h a, v16h b, v4f c)
|
|
{
|
|
*out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w64(a, b, c);
|
|
}
|
|
|
|
|
|
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_bf16_w64:
|
|
// CHECK-GFX1100: v_wmma_f32_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
|
|
//
|
|
void test_amdgcn_wmma_f32_16x16x16_bf16_w64(global v4f* out, v16s a, v16s b, v4f c)
|
|
{
|
|
*out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64(a, b, c);
|
|
}
|
|
|
|
|
|
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f16_16x16x16_f16_w64:
|
|
// CHECK-GFX1100: v_wmma_f16_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
|
|
//
|
|
void test_amdgcn_wmma_f16_16x16x16_f16_w64(global v8h* out, v16h a, v16h b, v8h c)
|
|
{
|
|
*out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w64(a, b, c, true);
|
|
}
|
|
|
|
|
|
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_bf16_16x16x16_bf16_w64:
|
|
// CHECK-GFX1100: v_wmma_bf16_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
|
|
//
|
|
void test_amdgcn_wmma_bf16_16x16x16_bf16_w64(global v8s* out, v16s a, v16s b, v8s c)
|
|
{
|
|
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64(a, b, c, true);
|
|
}
|
|
|
|
|
|
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu8_w64:
|
|
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu8 v[{{.*}}], v[{{.*}} v[{{.*}} v[{{.*}}] neg_lo:[1,1,0]
|
|
//
|
|
void test_amdgcn_wmma_i32_16x16x16_iu8_w64(global v4i* out, v4i a, v4i b, v4i c)
|
|
{
|
|
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64(true, a, true, b, c, false);
|
|
}
|
|
|
|
|
|
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu4_w64:
|
|
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu4 v[{{.*}} v[{{.*}} v[{{.*}} v[{{.*}}neg_lo:[1,1,0]
|
|
//
|
|
void test_amdgcn_wmma_i32_16x16x16_iu4_w64(global v4i* out, v2i a, v2i b, v4i c)
|
|
{
|
|
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64(true, a, true, b, c, false);
|
|
}
|
|
|
|
#endif
|