Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-01-16 13:08:42 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

AMDGPU: Export workitem builtins

Browse Source 
Reviewers: tstellardAMD

Differential Revision: http://reviews.llvm.org/D20299

llvm-svn: 275030
This commit is contained in:
Jan Vesely 2016-07-10 22:38:04 +00:00
parent 617c962752
commit d7e03a5bd9
4 changed files with 137 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

28
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -17,6 +17,20 @@
#if defined(BUILTIN) && !defined(TARGET_BUILTIN)
# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
#endif
//===----------------------------------------------------------------------===//
// SI+ only builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*2", "nc")
BUILTIN(__builtin_amdgcn_implicitarg_ptr, "Uc*2", "nc")
BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workgroup_id_z, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
//===----------------------------------------------------------------------===//
// Instruction builtins.
@ -67,6 +81,20 @@ TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
//===----------------------------------------------------------------------===//
// R600-NI only builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_r600_implicitarg_ptr, "Uc*7", "nc")
BUILTIN(__builtin_r600_read_tgid_x, "Ui", "nc")
BUILTIN(__builtin_r600_read_tgid_y, "Ui", "nc")
BUILTIN(__builtin_r600_read_tgid_z, "Ui", "nc")
BUILTIN(__builtin_r600_read_tidig_x, "Ui", "nc")
BUILTIN(__builtin_r600_read_tidig_y, "Ui", "nc")
BUILTIN(__builtin_r600_read_tidig_z, "Ui", "nc")
//===----------------------------------------------------------------------===//
// Legacy names with amdgpu prefix
//===----------------------------------------------------------------------===//

28
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -26,6 +26,7 @@
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/InlineAsm.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/MDBuilder.h"
#include <sstream>
using namespace clang;
@ -331,6 +332,17 @@ static llvm::Value *EmitOverflowIntrinsic(CodeGenFunction &CGF,
return CGF.Builder.CreateExtractValue(Tmp, 0);
}
static Value *emitRangedBuiltin(CodeGenFunction &CGF,
unsigned IntrinsicID,
int low, int high) {
llvm::MDBuilder MDHelper(CGF.getLLVMContext());
llvm::MDNode *RNode = MDHelper.createRange(APInt(32, low), APInt(32, high));
Value *F = CGF.CGM.getIntrinsic(IntrinsicID, {});
llvm::Instruction *Call = CGF.Builder.CreateCall(F);
Call->setMetadata(llvm::LLVMContext::MD_range, RNode);
return Call;
}
namespace {
struct WidthAndSignedness {
unsigned Width;
@ -7670,6 +7682,22 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
return emitFPIntBuiltin(*this, E, Intrinsic::AMDGPU_ldexp);
}
// amdgcn workitem
case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0, 1024);
case AMDGPU::BI__builtin_amdgcn_workitem_id_y:
return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_y, 0, 1024);
case AMDGPU::BI__builtin_amdgcn_workitem_id_z:
return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 1024);
// r600 workitem
case AMDGPU::BI__builtin_r600_read_tidig_x:
return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_x, 0, 1024);
case AMDGPU::BI__builtin_r600_read_tidig_y:
return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_y, 0, 1024);
case AMDGPU::BI__builtin_r600_read_tidig_z:
return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0, 1024);
default:
return nullptr;
}

45
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@ -291,6 +291,49 @@ void test_legacy_ldexp_f64(global double* out, double a, int b)
*out = __builtin_amdgpu_ldexp(a, b);
}
// CHECK-LABEL: @test_kernarg_segment_ptr
// CHECK: call i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr()
void test_kernarg_segment_ptr(__attribute__((address_space(2))) unsigned char ** out)
{
*out = __builtin_amdgcn_kernarg_segment_ptr();
}
// CHECK-LABEL: @test_implicitarg_ptr
// CHECK: call i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr()
void test_implicitarg_ptr(__attribute__((address_space(2))) unsigned char ** out)
{
*out = __builtin_amdgcn_implicitarg_ptr();
}
// CHECK-LABEL: @test_get_group_id(
// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()
void test_get_group_id(int d, global int *out)
{
switch (d) {
case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;
case 1: *out = __builtin_amdgcn_workgroup_id_y(); break;
case 2: *out = __builtin_amdgcn_workgroup_id_z(); break;
default: *out = 0;
}
}
// CHECK-LABEL: @test_get_local_id(
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]]
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]]
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[WI_RANGE]]
void test_get_local_id(int d, global int *out)
{
switch (d) {
case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
default: *out = 0;
}
}
// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
// CHECK: ![[EXEC]] = !{!"exec"}
// CHECK-DAG: ![[EXEC]] = !{!"exec"}

37
clang/test/CodeGenOpenCL/builtins-r600.cl
View File

@ -32,3 +32,40 @@ void test_legacy_ldexp_f64(global double* out, double a, int b)
*out = __builtin_amdgpu_ldexp(a, b);
}
#endif
// CHECK-LABEL: @test_implicitarg_ptr
// CHECK: call i8 addrspace(7)* @llvm.r600.implicitarg.ptr()
void test_implicitarg_ptr(__attribute__((address_space(7))) unsigned char ** out)
{
*out = __builtin_r600_implicitarg_ptr();
}
// CHECK-LABEL: @test_get_group_id(
// CHECK: tail call i32 @llvm.r600.read.tgid.x()
// CHECK: tail call i32 @llvm.r600.read.tgid.y()
// CHECK: tail call i32 @llvm.r600.read.tgid.z()
void test_get_group_id(int d, global int *out)
{
switch (d) {
case 0: *out = __builtin_r600_read_tgid_x(); break;
case 1: *out = __builtin_r600_read_tgid_y(); break;
case 2: *out = __builtin_r600_read_tgid_z(); break;
default: *out = 0;
}
}
// CHECK-LABEL: @test_get_local_id(
// CHECK: tail call i32 @llvm.r600.read.tidig.x(), !range [[WI_RANGE:![0-9]*]]
// CHECK: tail call i32 @llvm.r600.read.tidig.y(), !range [[WI_RANGE]]
// CHECK: tail call i32 @llvm.r600.read.tidig.z(), !range [[WI_RANGE]]
void test_get_local_id(int d, global int *out)
{
switch (d) {
case 0: *out = __builtin_r600_read_tidig_x(); break;
case 1: *out = __builtin_r600_read_tidig_y(); break;
case 2: *out = __builtin_r600_read_tidig_z(); break;
default: *out = 0;
}
}
// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}