[PATCH] D17188: AMDGPU: Add sin/cos builtins

[Matt Arsenault via llvm-commits]

arsenm created this revision.
arsenm added a reviewer: tstellarAMD.
arsenm added a subscriber: llvm-commits.

http://reviews.llvm.org/D17188

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

Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -113,6 +113,27 @@
   *out = __builtin_amdgcn_rsq_clamp(a);
 }
 
+// CHECK-LABEL: @test_sin_f32
+// CHECK: call float @llvm.amdgcn.sin.f32
+void test_sin_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_sinf(a);
+}
+
+// CHECK-LABEL: @test_cos_f32
+// CHECK: call float @llvm.amdgcn.cos.f32
+void test_cos_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_cosf(a);
+}
+
+// CHECK-LABEL: @test_log_clamp_f32
+// CHECK: call float @llvm.amdgcn.log.clamp.f32
+void test_log_clamp_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_log_clampf(a);
+}
+
 // CHECK-LABEL: @test_ldexp_f32
 // CHECK: call float @llvm.amdgcn.ldexp.f32
 void test_ldexp_f32(global float* out, float a, int b)
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -7082,6 +7082,12 @@
   case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
   case AMDGPU::BI__builtin_amdgcn_rsq_clampf:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp);
+  case AMDGPU::BI__builtin_amdgcn_sinf:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin);
+  case AMDGPU::BI__builtin_amdgcn_cosf:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos);
+  case AMDGPU::BI__builtin_amdgcn_log_clampf:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log_clamp);
   case AMDGPU::BI__builtin_amdgcn_ldexp:
   case AMDGPU::BI__builtin_amdgcn_ldexpf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -29,6 +29,9 @@
 BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc")
 BUILTIN(__builtin_amdgcn_rsq_clampf, "ff", "nc")
+BUILTIN(__builtin_amdgcn_sinf, "ff", "nc")
+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_class, "bdi", "nc")


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17188.47769.patch
Type: text/x-patch
Size: 2375 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160212/41bc27e2/attachment.bin>