This is an archive of the discontinued LLVM Phabricator instance.

Differential D17188

AMDGPU: Add sin/cos builtins
ClosedPublic

Authored by arsenm on Feb 11 2016, 11:44 PM.

Details

Reviewers
• tstellarAMD

Diff Detail

Event Timeline

arsenm updated this revision to Diff 47769.Feb 11 2016, 11:44 PM
arsenm retitled this revision from to AMDGPU: Add sin/cos builtins.
arsenm updated this object.
arsenm added a reviewer: • tstellarAMD.
arsenm added a parent revision: D17187: AMDGPU: Add intrinsics for sin/cos.
arsenm added a subscriber: llvm-commits.
arsenm edited subscribers, added: cfe-commits; removed: llvm-commits.
• tstellarAMD accepted this revision.Feb 12 2016, 11:48 AM
• tstellarAMD edited edge metadata.

LGTM.

This revision is now accepted and ready to land.Feb 12 2016, 11:48 AM
arsenm closed this revision.Feb 12 2016, 5:25 PM

r260783

Revision Contents

PathSize
include/
clang/
Basic/
3 lines
lib/
CodeGen/
6 lines
test/
CodeGenOpenCL/
21 lines

Diff 47769

include/clang/Basic/BuiltinsAMDGPU.def

Show All 23 Lines
BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc") BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc")
BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc") BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc")
BUILTIN(__builtin_amdgcn_rcp, "dd", "nc") BUILTIN(__builtin_amdgcn_rcp, "dd", "nc")
BUILTIN(__builtin_amdgcn_rcpf, "ff", "nc") BUILTIN(__builtin_amdgcn_rcpf, "ff", "nc")
BUILTIN(__builtin_amdgcn_rsq, "dd", "nc") BUILTIN(__builtin_amdgcn_rsq, "dd", "nc")
BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc") BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc")
BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc") BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc")
BUILTIN(__builtin_amdgcn_rsq_clampf, "ff", "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_ldexp, "ddi", "nc")
BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc") BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc")
BUILTIN(__builtin_amdgcn_class, "bdi", "nc") BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
BUILTIN(__builtin_amdgcn_classf, "bfi", "nc") BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc") BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc") BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc") BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc") BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
// Legacy names with amdgpu prefix // Legacy names with amdgpu prefix
BUILTIN(__builtin_amdgpu_rsq, "dd", "nc") BUILTIN(__builtin_amdgpu_rsq, "dd", "nc")
BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc") BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc")
BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc") BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc")
BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc") BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
#undef BUILTIN #undef BUILTIN

lib/CodeGen/CGBuiltin.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 7,076 Lines • ▼ Show 20 LinesValue *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_rcpf: case AMDGPU::BI__builtin_amdgcn_rcpf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rcp); return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rcp);
case AMDGPU::BI__builtin_amdgcn_rsq: case AMDGPU::BI__builtin_amdgcn_rsq:
case AMDGPU::BI__builtin_amdgcn_rsqf: case AMDGPU::BI__builtin_amdgcn_rsqf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq); return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq);
case AMDGPU::BI__builtin_amdgcn_rsq_clamp: case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
case AMDGPU::BI__builtin_amdgcn_rsq_clampf: case AMDGPU::BI__builtin_amdgcn_rsq_clampf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp); 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_ldexp:
case AMDGPU::BI__builtin_amdgcn_ldexpf: case AMDGPU::BI__builtin_amdgcn_ldexpf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp); return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
case AMDGPU::BI__builtin_amdgcn_class: case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf: case AMDGPU::BI__builtin_amdgcn_classf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class); return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
// Legacy amdgpu prefix // Legacy amdgpu prefix
▲ Show 20 Lines • Show All 353 Lines • Show Last 20 Lines

test/CodeGenOpenCL/builtins-amdgcn.cl

Show First 20 Lines • Show All 107 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_rsq_clamp_f64 // CHECK-LABEL: @test_rsq_clamp_f64
// CHECK: call double @llvm.amdgcn.rsq.clamp.f64 // CHECK: call double @llvm.amdgcn.rsq.clamp.f64
void test_rsq_clamp_f64(global double* out, double a) void test_rsq_clamp_f64(global double* out, double a)
{ {
*out = __builtin_amdgcn_rsq_clamp(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-LABEL: @test_ldexp_f32
// CHECK: call float @llvm.amdgcn.ldexp.f32 // CHECK: call float @llvm.amdgcn.ldexp.f32
void test_ldexp_f32(global float* out, float a, int b) void test_ldexp_f32(global float* out, float a, int b)
{ {
*out = __builtin_amdgcn_ldexpf(a, b); *out = __builtin_amdgcn_ldexpf(a, b);
} }
// CHECK-LABEL: @test_ldexp_f64 // CHECK-LABEL: @test_ldexp_f64
▲ Show 20 Lines • Show All 80 Lines • Show Last 20 Lines