Skip to content

Commit 58dada5

Browse files
rovkarovka
committed
[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
1 parent a048d73 commit 58dada5

File tree

5 files changed

+151
-1
lines changed

5 files changed

+151
-1
lines changed

‎cross-project-tests/CMakeLists.txt

+8
Original file line numberDiff line numberDiff line change
@@ -86,6 +86,13 @@ add_lit_testsuite(check-intrinsic-headers "Running intrinsic header tests"
8686
DEPENDS ${CROSS_PROJECT_TEST_DEPS}
8787
)
8888

89+
# AMDGPU tests.
90+
add_lit_testsuite(check-cross-amdgpu "Running AMDGPU cross-project tests"
91+
${CMAKE_CURRENT_BINARY_DIR}/amdgpu
92+
EXCLUDE_FROM_CHECK_ALL
93+
DEPENDS clang
94+
)
95+
8996
# Add check-cross-project-* targets.
9097
add_lit_testsuites(CROSS_PROJECT ${CMAKE_CURRENT_SOURCE_DIR}
9198
DEPENDS ${CROSS_PROJECT_TEST_DEPS}
@@ -94,3 +101,4 @@ add_lit_testsuites(CROSS_PROJECT ${CMAKE_CURRENT_SOURCE_DIR}
94101
set_target_properties(check-cross-project PROPERTIES FOLDER "Tests")
95102
set_target_properties(check-debuginfo PROPERTIES FOLDER "Tests")
96103
set_target_properties(check-intrinsic-headers PROPERTIES FOLDER "Tests")
104+
set_target_properties(check-cross-amdgpu PROPERTIES FOLDER "Tests")

‎cross-project-tests/amdgpu/builtins-amdgcn-wmma-w32.cl

+68
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,68 @@
1+
// 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
2+
3+
typedef float v4f __attribute__((ext_vector_type(4)));
4+
typedef float v8f __attribute__((ext_vector_type(8)));
5+
typedef half v16h __attribute__((ext_vector_type(16)));
6+
typedef int v2i __attribute__((ext_vector_type(2)));
7+
typedef int v4i __attribute__((ext_vector_type(4)));
8+
typedef int v8i __attribute__((ext_vector_type(8)));
9+
typedef short v16s __attribute__((ext_vector_type(16)));
10+
11+
#ifdef WMMA_GFX1100_TESTS
12+
13+
// Wave32
14+
15+
16+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_f16_w32:
17+
// CHECK-GFX1100: v_wmma_f32_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
18+
//
19+
void test_amdgcn_wmma_f32_16x16x16_f16_w32(global v8f* out, v16h a, v16h b, v8f c)
20+
{
21+
*out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c);
22+
}
23+
24+
25+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_bf16_w32:
26+
// CHECK-GFX1100: v_wmma_f32_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
27+
//
28+
void test_amdgcn_wmma_f32_16x16x16_bf16_w32(global v8f* out, v16s a, v16s b, v8f c)
29+
{
30+
*out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a, b, c);
31+
}
32+
33+
34+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f16_16x16x16_f16_w32:
35+
// CHECK-GFX1100: v_wmma_f16_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
36+
//
37+
void test_amdgcn_wmma_f16_16x16x16_f16_w32(global v16h* out, v16h a, v16h b, v16h c)
38+
{
39+
*out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(a, b, c, true);
40+
}
41+
42+
43+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_bf16_16x16x16_bf16_w32:
44+
// CHECK-GFX1100: v_wmma_bf16_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
45+
//
46+
void test_amdgcn_wmma_bf16_16x16x16_bf16_w32(global v16s* out, v16s a, v16s b, v16s c)
47+
{
48+
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a, b, c, true);
49+
}
50+
51+
52+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu8_w32:
53+
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu8 v[{{.*}}], v[{{.*}} v[{{.*}} v[{{.*}}] neg_lo:[1,1,0]
54+
//
55+
void test_amdgcn_wmma_i32_16x16x16_iu8_w32(global v8i* out, v4i a, v4i b, v8i c)
56+
{
57+
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a, true, b, c, false);
58+
}
59+
60+
61+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu4_w32:
62+
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu4 v[{{.*}}, v[{{.*}} v[{{.*}} v[{{.*}} neg_lo:[1,1,0]
63+
void test_amdgcn_wmma_i32_16x16x16_iu4_w32(global v8i* out, v2i a, v2i b, v8i c)
64+
{
65+
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(true, a, true, b, c, false);
66+
}
67+
68+
#endif

‎cross-project-tests/amdgpu/builtins-amdgcn-wmma-w64.cl

+72
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,72 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// 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
3+
4+
typedef float v4f __attribute__((ext_vector_type(4)));
5+
typedef float v8f __attribute__((ext_vector_type(8)));
6+
typedef half v8h __attribute__((ext_vector_type(8)));
7+
typedef half v16h __attribute__((ext_vector_type(16)));
8+
typedef int v2i __attribute__((ext_vector_type(2)));
9+
typedef int v4i __attribute__((ext_vector_type(4)));
10+
typedef int v8i __attribute__((ext_vector_type(8)));
11+
typedef short v8s __attribute__((ext_vector_type(8)));
12+
typedef short v16s __attribute__((ext_vector_type(16)));
13+
14+
#ifdef WMMA_GFX1100_TESTS
15+
16+
// Wave64
17+
18+
19+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_f16_w64:
20+
// CHECK-GFX1100: v_wmma_f32_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
21+
//
22+
void test_amdgcn_wmma_f32_16x16x16_f16_w64(global v4f* out, v16h a, v16h b, v4f c)
23+
{
24+
*out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w64(a, b, c);
25+
}
26+
27+
28+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_bf16_w64:
29+
// CHECK-GFX1100: v_wmma_f32_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
30+
//
31+
void test_amdgcn_wmma_f32_16x16x16_bf16_w64(global v4f* out, v16s a, v16s b, v4f c)
32+
{
33+
*out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64(a, b, c);
34+
}
35+
36+
37+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_f16_16x16x16_f16_w64:
38+
// CHECK-GFX1100: v_wmma_f16_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
39+
//
40+
void test_amdgcn_wmma_f16_16x16x16_f16_w64(global v8h* out, v16h a, v16h b, v8h c)
41+
{
42+
*out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w64(a, b, c, true);
43+
}
44+
45+
46+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_bf16_16x16x16_bf16_w64:
47+
// CHECK-GFX1100: v_wmma_bf16_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
48+
//
49+
void test_amdgcn_wmma_bf16_16x16x16_bf16_w64(global v8s* out, v16s a, v16s b, v8s c)
50+
{
51+
*out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64(a, b, c, true);
52+
}
53+
54+
55+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu8_w64:
56+
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu8 v[{{.*}}], v[{{.*}} v[{{.*}} v[{{.*}}] neg_lo:[1,1,0]
57+
//
58+
void test_amdgcn_wmma_i32_16x16x16_iu8_w64(global v4i* out, v4i a, v4i b, v4i c)
59+
{
60+
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64(true, a, true, b, c, false);
61+
}
62+
63+
64+
// CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu4_w64:
65+
// CHECK-GFX1100: v_wmma_i32_16x16x16_iu4 v[{{.*}} v[{{.*}} v[{{.*}} v[{{.*}}neg_lo:[1,1,0]
66+
//
67+
void test_amdgcn_wmma_i32_16x16x16_iu4_w64(global v4i* out, v2i a, v2i b, v4i c)
68+
{
69+
*out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64(true, a, true, b, c, false);
70+
}
71+
72+
#endif

‎cross-project-tests/amdgpu/lit.local.cfg

+2
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,2 @@
1+
if 'clang' not in config.available_features or 'AMDGPU' not in config.targets_to_build:
2+
config.unsupported = True

‎cross-project-tests/lit.cfg.py

+1-1
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@
2222
config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell)
2323

2424
# suffixes: A list of file extensions to treat as test files.
25-
config.suffixes = ['.c', '.cpp', '.m']
25+
config.suffixes = ['.c', '.cl', '.cpp', '.m']
2626

2727
# excludes: A list of directories to exclude from the testsuite. The 'Inputs'
2828
# subdirectories contain auxiliary inputs for various tests in their parent

0 commit comments

Comments
 (0)