blob: c0559e9468f6256de8b2e5d3443f06cbe60e1432 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
|
Modified from https://github.com/ROCm/llvm-project/commit/794ebeffcafbf6f4d86cb1bfd7a5a0d1d30f1fc7
From 794ebeffcafbf6f4d86cb1bfd7a5a0d1d30f1fc7 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Thu, 4 Jan 2024 20:40:30 +0700
Subject: [PATCH] device-libs: Fix native_rsqrt test failures for f16-as-f32
case
At some point the code quality regressed for these.
Change-Id: Ib5c4c1d093a6056dd6213ed6eaf157ad806101fb
---
amd/device-libs/test/compile/native_rsqrt.cl | 13 +++++++++----
1 file changed, 9 insertions(+), 4 deletions(-)
diff --git a/test/compile/native_rsqrt.cl b/test/compile/native_rsqrt.cl
index 94ef375e4d522f..291cc6f32c3351 100644
--- a/test/compile/native_rsqrt.cl
+++ b/test/compile/native_rsqrt.cl
@@ -6,16 +6,21 @@
half __ocml_native_rsqrt_f16(half);
-// GCN: {{^}}test_native_rsqrt_f16:
-// GFX600: v_rsq_f32
-// GFX700: v_rsq_f32
+// FIXME: Promoted case using full expansion
+// GCN-LABEL: {{^}}test_native_rsqrt_f16:
+// GFX600: v_sqrt_f32
+// GFX600: v_rcp_f32
+
+// GFX700: v_sqrt_f32
+// GFX700: v_rcp_f32
+
// GFX803: v_rsq_f16
kernel void test_native_rsqrt_f16(global half* restrict out, global half* restrict in) {
int id = get_local_id(0);
out[id] = __ocml_native_rsqrt_f16(in[id]);
}
-// GCN: {{^}}test_native_rsqrt_f32:
+// GCN-LABEL: {{^}}test_native_rsqrt_f32:
// GCN: v_rsq_f32
kernel void test_native_rsqrt_f32(global float* restrict out, global float* restrict in) {
int id = get_local_id(0);
|
