This is an archive of the discontinued LLVM Phabricator instance.

Differential D16498

AMDGPU: Add amdgcn cube builtins
ClosedPublic

Authored by arsenm on Jan 22 2016, 5:20 PM.

Details

Reviewers
• tstellarAMD

Diff Detail

Event Timeline

arsenm updated this revision to Diff 45777.Jan 22 2016, 5:20 PM
arsenm retitled this revision from to AMDGPU: Add amdgcn cube builtins.
arsenm updated this object.
arsenm added a reviewer: • tstellarAMD.
arsenm added a parent revision: D16497: AMDGPU: Add new amdgcn intrinsics for cube instructions.
arsenm added a subscriber: cfe-commits.
• tstellarAMD accepted this revision.Jan 25 2016, 7:16 PM
• tstellarAMD edited edge metadata.

LGTM.

This revision is now accepted and ready to land.Jan 25 2016, 7:16 PM
arsenm closed this revision.Jan 26 2016, 6:22 PM

r258794

Revision Contents

PathSize
include/
clang/
Basic/
4 lines
test/
CodeGenOpenCL/
24 lines

Diff 45777

include/clang/Basic/BuiltinsAMDGPU.def

Show All 27 Lines
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_clamped, "dd", "nc") BUILTIN(__builtin_amdgcn_rsq_clamped, "dd", "nc")
BUILTIN(__builtin_amdgcn_rsq_clampedf, "ff", "nc") BUILTIN(__builtin_amdgcn_rsq_clampedf, "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_cubesc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubetc, "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

test/CodeGenOpenCL/builtins-amdgcn.cl

Show First 20 Lines • Show All 142 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_s_barrier // CHECK-LABEL: @test_s_barrier
// CHECK: call void @llvm.amdgcn.s.barrier( // CHECK: call void @llvm.amdgcn.s.barrier(
void test_s_barrier() void test_s_barrier()
{ {
__builtin_amdgcn_s_barrier(); __builtin_amdgcn_s_barrier();
} }
// CHECK-LABEL: @test_cubeid(
// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
void test_cubeid(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubeid(a, b, c);
}
// CHECK-LABEL: @test_cubesc(
// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
void test_cubesc(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubesc(a, b, c);
}
// CHECK-LABEL: @test_cubetc(
// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
void test_cubetc(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubetc(a, b, c);
}
// CHECK-LABEL: @test_cubema(
// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
void test_cubema(global float* out, float a, float b, float c) {
*out = __builtin_amdgcn_cubema(a, b, c);
}
// Legacy intrinsics with AMDGPU prefix // Legacy intrinsics with AMDGPU prefix
// CHECK-LABEL: @test_legacy_rsq_f32 // CHECK-LABEL: @test_legacy_rsq_f32
// CHECK: call float @llvm.amdgcn.rsq.f32 // CHECK: call float @llvm.amdgcn.rsq.f32
void test_legacy_rsq_f32(global float* out, float a) void test_legacy_rsq_f32(global float* out, float a)
{ {
*out = __builtin_amdgpu_rsqf(a); *out = __builtin_amdgpu_rsqf(a);
} }
Show All 21 Lines