[PATCH] D26863: [AMDGPU] Change frexp.exp builtin to return i16 for f16 input

[Konstantin Zhuravlyov via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL287390: [AMDGPU] Change frexp.exp builtin to return i16 for f16 input (authored by kzhuravl).

Changed prior to commit:
  https://reviews.llvm.org/D26863?vs=78546&id=78586#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D26863

Files:
  cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
  cfe/trunk/lib/CodeGen/CGBuiltin.cpp
  cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
  cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl


Index: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
===================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -55,7 +55,7 @@
 }
 
 // CHECK-LABEL: @test_frexp_exp_f16
-// CHECK: call i32 @llvm.amdgcn.frexp.exp.f16
+// CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16
 void test_frexp_exp_f16(global short* out, half a)
 {
   *out = __builtin_amdgcn_frexp_exph(a);
Index: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -166,14 +166,14 @@
 }
 
 // CHECK-LABEL: @test_frexp_exp_f32
-// CHECK: call i32 @llvm.amdgcn.frexp.exp.f32
+// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32
 void test_frexp_exp_f32(global int* out, float a)
 {
   *out = __builtin_amdgcn_frexp_expf(a);
 }
 
 // CHECK-LABEL: @test_frexp_exp_f64
-// CHECK: call i32 @llvm.amdgcn.frexp.exp.f64
+// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64
 void test_frexp_exp_f64(global int* out, double a)
 {
   *out = __builtin_amdgcn_frexp_exp(a);
Index: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp
@@ -8250,9 +8250,18 @@
   case AMDGPU::BI__builtin_amdgcn_frexp_manth:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_mant);
   case AMDGPU::BI__builtin_amdgcn_frexp_exp:
-  case AMDGPU::BI__builtin_amdgcn_frexp_expf:
-  case AMDGPU::BI__builtin_amdgcn_frexp_exph:
-    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
+  case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
+    Value *Src0 = EmitScalarExpr(E->getArg(0));
+    Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp,
+                                { Builder.getInt32Ty(), Src0->getType() });
+    return Builder.CreateCall(F, Src0);
+  }
+  case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
+    Value *Src0 = EmitScalarExpr(E->getArg(0));
+    Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp,
+                                { Builder.getInt16Ty(), Src0->getType() });
+    return Builder.CreateCall(F, Src0);
+  }
   case AMDGPU::BI__builtin_amdgcn_fract:
   case AMDGPU::BI__builtin_amdgcn_fractf:
   case AMDGPU::BI__builtin_amdgcn_fracth:
Index: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
@@ -92,7 +92,7 @@
 TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts")
-TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "ih", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D26863.78586.patch
Type: text/x-patch
Size: 3268 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20161118/57e81e90/attachment-0001.bin>