From 129719e62956d2d4c0166252a9b41ee55ead8099 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Sat, 13 Feb 2016 01:21:09 +0000 Subject: [PATCH] AMDGPU: Add sin/cos builtins git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@260783 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/BuiltinsAMDGPU.def | 3 +++ lib/CodeGen/CGBuiltin.cpp | 6 ++++++ test/CodeGenOpenCL/builtins-amdgcn.cl | 21 +++++++++++++++++++++ 3 files changed, 30 insertions(+) diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index ab7a01a7d0..180d7a3f9c 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -29,6 +29,9 @@ BUILTIN(__builtin_amdgcn_rsq, "dd", "nc") 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") diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp index 46fabfd70a..e384851112 100644 --- a/lib/CodeGen/CGBuiltin.cpp +++ b/lib/CodeGen/CGBuiltin.cpp @@ -7082,6 +7082,12 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, 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); diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 4dffc84e58..78741b6d91 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -113,6 +113,27 @@ void test_rsq_clamp_f64(global double* out, double a) *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) -- 2.40.0