llvm-project/cross-project-tests/amdgpu/builtins-amdgcn-wmma-w32.cl
Diana Picus 58dada5f0a [AMDGPU] Add cross-project-tests for WMMA builtins
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
2023-02-17 09:20:46 +01:00

69 lines
2.5 KiB
Common Lisp

// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize32 -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 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 v16s __attribute__((ext_vector_type(16)));
#ifdef WMMA_GFX1100_TESTS
// Wave32
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_f16_w32:
// CHECK-GFX1100: v_wmma_f32_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
//
void test_amdgcn_wmma_f32_16x16x16_f16_w32(global v8f* out, v16h a, v16h b, v8f c)
{
*out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c);
}
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_bf16_w32:
// CHECK-GFX1100: v_wmma_f32_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
//
void test_amdgcn_wmma_f32_16x16x16_bf16_w32(global v8f* out, v16s a, v16s b, v8f c)
{
*out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a, b, c);
}
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f16_16x16x16_f16_w32:
// CHECK-GFX1100: v_wmma_f16_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
//
void test_amdgcn_wmma_f16_16x16x16_f16_w32(global v16h* out, v16h a, v16h b, v16h c)
{
*out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(a, b, c, true);
}
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_bf16_16x16x16_bf16_w32:
// CHECK-GFX1100: v_wmma_bf16_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
//
void test_amdgcn_wmma_bf16_16x16x16_bf16_w32(global v16s* out, v16s a, v16s b, v16s c)
{
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a, b, c, true);
}
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu8_w32:
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu8 v[{{.*}}], v[{{.*}} v[{{.*}} v[{{.*}}] neg_lo:[1,1,0]
//
void test_amdgcn_wmma_i32_16x16x16_iu8_w32(global v8i* out, v4i a, v4i b, v8i c)
{
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a, true, b, c, false);
}
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu4_w32:
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu4 v[{{.*}}, v[{{.*}} v[{{.*}} v[{{.*}} neg_lo:[1,1,0]
void test_amdgcn_wmma_i32_16x16x16_iu4_w32(global v8i* out, v2i a, v2i b, v8i c)
{
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(true, a, true, b, c, false);
}
#endif