This is an archive of the discontinued LLVM Phabricator instance.

Differential D18338

AMDGPU: Add frexp_exp builtins
ClosedPublic

Authored by arsenm on Mar 21 2016, 3:25 PM.

Details

Reviewers
• tstellarAMD

Diff Detail

Event Timeline

arsenm updated this revision to Diff 51237.Mar 21 2016, 3:25 PM
arsenm retitled this revision from to AMDGPU: Add frexp_exp intrinsic.
arsenm updated this object.
arsenm added a reviewer: • tstellarAMD.
arsenm retitled this revision from AMDGPU: Add frexp_exp intrinsic to AMDGPU: Add frexp_exp builtins.
arsenm added a subscriber: llvm-commits.
arsenm edited subscribers, added: cfe-commits; removed: llvm-commits.
• tstellarAMD accepted this revision.Mar 30 2016, 6:42 AM
• tstellarAMD edited edge metadata.

LGTM.

This revision is now accepted and ready to land.Mar 30 2016, 6:42 AM
arsenm closed this revision.Mar 30 2016, 4:03 PM

r264960

Revision Contents

PathSize
include/
clang/
Basic/
4 lines
lib/
CodeGen/
8 lines
test/
CodeGenOpenCL/
28 lines

Diff 51237

include/clang/Basic/BuiltinsAMDGPU.def

Show All 28 Lines
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_sinf, "ff", "nc")
BUILTIN(__builtin_amdgcn_cosf, "ff", "nc") BUILTIN(__builtin_amdgcn_cosf, "ff", "nc")
BUILTIN(__builtin_amdgcn_log_clampf, "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_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_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")
BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n") BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n") BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
Show All 16 Lines

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,098 Lines • ▼ Show 20 Lines case AMDGPU::BI__builtin_amdgcn_sinf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin); return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin);
case AMDGPU::BI__builtin_amdgcn_cosf: case AMDGPU::BI__builtin_amdgcn_cosf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos); return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos);
case AMDGPU::BI__builtin_amdgcn_log_clampf: case AMDGPU::BI__builtin_amdgcn_log_clampf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log_clamp); 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_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_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
case AMDGPU::BI__builtin_amdgpu_rsq: case AMDGPU::BI__builtin_amdgpu_rsq:
case AMDGPU::BI__builtin_amdgpu_rsqf: { case AMDGPU::BI__builtin_amdgpu_rsqf: {
if (getTarget().getTriple().getArch() == Triple::amdgcn) if (getTarget().getTriple().getArch() == Triple::amdgcn)
▲ Show 20 Lines • Show All 350 Lines • Show Last 20 Lines

test/CodeGenOpenCL/builtins-amdgcn.cl

Show First 20 Lines • Show All 144 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_ldexp_f64 // CHECK-LABEL: @test_ldexp_f64
// CHECK: call double @llvm.amdgcn.ldexp.f64 // CHECK: call double @llvm.amdgcn.ldexp.f64
void test_ldexp_f64(global double* out, double a, int b) void test_ldexp_f64(global double* out, double a, int b)
{ {
*out = __builtin_amdgcn_ldexp(a, b); *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-LABEL: @test_class_f32
// CHECK: call i1 @llvm.amdgcn.class.f32 // CHECK: call i1 @llvm.amdgcn.class.f32
void test_class_f32(global float* out, float a, int b) void test_class_f32(global float* out, float a, int b)
{ {
*out = __builtin_amdgcn_classf(a, b); *out = __builtin_amdgcn_classf(a, b);
} }
// CHECK-LABEL: @test_class_f64 // CHECK-LABEL: @test_class_f64
▲ Show 20 Lines • Show All 89 Lines • Show Last 20 Lines