This is an archive of the discontinued LLVM Phabricator instance.

Differential D33276

[AMDGPU] add __builtin_amdgcn_s_getpc
ClosedPublic

Authored by timcorringham on May 17 2017, 7:02 AM.

Details

Reviewers
arsenm
Commits
rG702fe45bcd08: [AMDGPU] add __builtin_amdgcn_s_getpc
rC303861: [AMDGPU] add __builtin_amdgcn_s_getpc
rL303861: [AMDGPU] add __builtin_amdgcn_s_getpc
Summary

Added the builtin corresponding to the s_getpc intrinsic added in llvm D32862

Diff Detail

Event Timeline

timcorringham created this revision.May 17 2017, 7:02 AM
Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptMay 17 2017, 7:02 AM
timcorringham added a reviewer: arsenm.May 17 2017, 7:04 AM
timcorringham mentioned this in D32862: [AMDGPU] add intrinsic for s_getpc.May 17 2017, 7:48 AM
arsenm accepted this revision.May 17 2017, 12:53 PM

LGTM

This revision is now accepted and ready to land.May 17 2017, 12:53 PM
Closed by commit rL303861: [AMDGPU] add __builtin_amdgcn_s_getpc (authored by timcorringham). · Explain WhyMay 25 2017, 7:16 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
include/
clang/
Basic/
1 line
test/
CodeGenOpenCL/
7 lines

Diff 99294

include/clang/Basic/BuiltinsAMDGPU.def

Show All 30 Lines
BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc") BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc") BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc") BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Instruction builtins. // Instruction builtins.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n") BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n")
BUILTIN(__builtin_amdgcn_s_getpc, "LUi", "n")
BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n") BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n") BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n")
BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n") BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n")
BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n") BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n") BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
▲ Show 20 Lines • Show All 95 Lines • Show Last 20 Lines

test/CodeGenOpenCL/builtins-amdgcn.cl

Show First 20 Lines • Show All 474 Lines • ▼ Show 20 Lines
// CHECK-LABEL: @test_fmed3_f32 // CHECK-LABEL: @test_fmed3_f32
// CHECK: call float @llvm.amdgcn.fmed3.f32( // CHECK: call float @llvm.amdgcn.fmed3.f32(
void test_fmed3_f32(global float* out, float a, float b, float c) void test_fmed3_f32(global float* out, float a, float b, float c)
{ {
*out = __builtin_amdgcn_fmed3f(a, b, c); *out = __builtin_amdgcn_fmed3f(a, b, c);
} }
// CHECK-LABEL: @test_s_getpc
// CHECK: call i64 @llvm.amdgcn.s.getpc()
void test_s_getpc(global ulong* out)
{
*out = __builtin_amdgcn_s_getpc();
}
// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent } // CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
// CHECK-DAG: ![[EXEC]] = !{!"exec"} // CHECK-DAG: ![[EXEC]] = !{!"exec"}