device-libs: Implement rsqrt by enabling contract on fdiv+sqrt intrinsic

This was already how the half version is implemented.

The test already showed this was preserved, but tighten up the checks to
make sure there aren't any additional instructions

Change-Id: I0dbb6198e1f8e3a0168c73007cb06a72da6c36d9

Author: arsenm

amd/device-libs/ocml/src/nativeF.cl

@@ -24,7 +24,8 @@ MATH_MANGLE(native_sqrt)(float x)
 CONSTATTR float
 MATH_MANGLE(native_rsqrt)(float x)
 {
-    return __builtin_amdgcn_rsqf(x);
+    #pragma clang fp contract(fast)
+    return 1.0f / __builtin_sqrtf(x);
 }
 
 CONSTATTR float

amd/device-libs/test/compile/native_rsqrt.cl

@@ -14,14 +14,24 @@ half __ocml_native_rsqrt_f16(half);
 // GFX700: v_sqrt_f32
 // GFX700: v_rcp_f32
 
-// GFX803: v_rsq_f16
+// GFX803: {{(flat|global|buffer)}}_load_{{(ushort|b16)}} [[VAL:v[0-9+]]],
+// GFX803-NOT: [[VAL]]
+// GFX803: v_rsq_f16{{(_e32)?}} [[RESULT:v[0-9]+]], [[VAL]]
+// GFX803-NOT: [[RESULT]]
+// GFX803: [[RESULT]]
+// GFX803-NOT: [[RESULT]]
 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-LABEL: {{^}}test_native_rsqrt_f32:
-// GCN: v_rsq_f32
+// GCN: {{(flat|global|buffer)}}_load_{{(dword|b32)}} [[VAL:v[0-9+]]],
+// GCN-NOT: [[VAL]]
+// GCN: v_rsq_f32{{(_e32)?}} [[RESULT:v[0-9]+]], [[VAL]]
+// GCN-NOT: [[RESULT]]
+// GCN: [[RESULT]]
+// GCN-NOT: [[RESULT]]
 kernel void test_native_rsqrt_f32(global float* restrict out, global float* restrict in) {
     int id = get_local_id(0);
     out[id] = native_rsqrt(in[id]);