From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Sat, 28 May 2016 00:43:27 +0000 (+0000)
Subject: AMDGPU: Add fract builtin
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=ddf13e224e3a784768022ba717b2ae1fb65ded73;p=clang

AMDGPU: Add fract builtin

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@271080 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 04a498fdee..d712789366 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -38,6 +38,8 @@ BUILTIN(__builtin_amdgcn_frexp_mant, "dd", "nc")
 BUILTIN(__builtin_amdgcn_frexp_mantf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_frexp_exp, "id", "nc")
 BUILTIN(__builtin_amdgcn_frexp_expf, "if", "nc")
+BUILTIN(__builtin_amdgcn_fract, "dd", "nc")
+BUILTIN(__builtin_amdgcn_fractf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
 BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
 BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp
index 2f96549355..728b1d0a7d 100644
--- a/lib/CodeGen/CGBuiltin.cpp
+++ b/lib/CodeGen/CGBuiltin.cpp
@@ -7129,6 +7129,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
   }
+  case AMDGPU::BI__builtin_amdgcn_fract:
+  case AMDGPU::BI__builtin_amdgcn_fractf:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
   case AMDGPU::BI__builtin_amdgcn_class:
   case AMDGPU::BI__builtin_amdgcn_classf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index 8ebc3ae3ab..11448c3744 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -178,6 +178,20 @@ void test_frexp_exp_f64(global int* out, double a)
   *out = __builtin_amdgcn_frexp_exp(a);
 }
 
+// CHECK-LABEL: @test_fract_f32
+// CHECK: call float @llvm.amdgcn.fract.f32
+void test_fract_f32(global int* out, float a)
+{
+  *out = __builtin_amdgcn_fractf(a);
+}
+
+// CHECK-LABEL: @test_fract_f64
+// CHECK: call double @llvm.amdgcn.fract.f64
+void test_fract_f64(global int* out, double a)
+{
+  *out = __builtin_amdgcn_fract(a);
+}
+
 // CHECK-LABEL: @test_class_f32
 // CHECK: call i1 @llvm.amdgcn.class.f32
 void test_class_f32(global float* out, float a, int b)