Author: arsenm
Date: Wed Mar 30 17:57:40 2016
New Revision: 264960
URL: http://llvm.org/viewvc/llvm-project?rev=264960&view=rev
Log:
AMDGPU: Add frexp_mant + frexp_exp builtins
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=264960&r1=264959&r2=264960&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Wed Mar 30 17:57:40 2016
@@ -34,6 +34,10 @@ BUILTIN(__builtin_amdgcn_cosf, "ff", "nc
BUILTIN(__builtin_amdgcn_log_clampf, "ff", "nc")
BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc")
BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc")
+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_class, "bdi", "nc")
BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=264960&r1=264959&r2=264960&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Wed Mar 30 17:57:40 2016
@@ -7105,6 +7105,14 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
case AMDGPU::BI__builtin_amdgcn_ldexp:
case AMDGPU::BI__builtin_amdgcn_ldexpf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
+ case AMDGPU::BI__builtin_amdgcn_frexp_mant:
+ case AMDGPU::BI__builtin_amdgcn_frexp_mantf: {
+ return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_mant);
+ }
+ case AMDGPU::BI__builtin_amdgcn_frexp_exp:
+ case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
+ return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
+ }
case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=264960&r1=264959&r2=264960&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Wed Mar 30 17:57:40 2016
@@ -150,6 +150,34 @@ void test_ldexp_f64(global double* out,
*out = __builtin_amdgcn_ldexp(a, b);
}
+// CHECK-LABEL: @test_frexp_mant_f32
+// CHECK: call float @llvm.amdgcn.frexp.mant.f32
+void test_frexp_mant_f32(global float* out, float a)
+{
+ *out = __builtin_amdgcn_frexp_mantf(a);
+}
+
+// CHECK-LABEL: @test_frexp_mant_f64
+// CHECK: call double @llvm.amdgcn.frexp.mant.f64
+void test_frexp_mant_f64(global double* out, double a)
+{
+ *out = __builtin_amdgcn_frexp_mant(a);
+}
+
+// CHECK-LABEL: @test_frexp_exp_f32
+// CHECK: call i32 @llvm.amdgcn.frexp.exp.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
+void test_frexp_exp_f64(global int* out, double a)
+{
+ *out = __builtin_amdgcn_frexp_exp(a);
+}
+
// CHECK-LABEL: @test_class_f32
// CHECK: call i1 @llvm.amdgcn.class.f32
void test_class_f32(global float* out, float a, int b)
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
