maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

AMDGPU: Add builtins for recently added intrinsics 

llvm-svn: 262126
This commit is contained in:
Matt Arsenault 2016-02-27 09:54:43 +00:00
parent 360d244d5b
commit 39edcd0e1d
3 changed files with 41 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -40,8 +40,18 @@ BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
//===----------------------------------------------------------------------===//
// VI+ only builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n")
//===----------------------------------------------------------------------===//
// Legacy names with amdgpu prefix
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgpu_rsq, "dd", "nc")
BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc")
BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc")

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

@ -3,6 +3,8 @@
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
typedef unsigned long ulong;
// CHECK-LABEL: @test_div_scale_f64
// CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)
// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1
@ -169,6 +171,29 @@ void test_s_barrier()
__builtin_amdgcn_s_barrier();
}
// CHECK-LABEL: @test_s_memtime
// CHECK: call i64 @llvm.amdgcn.s.memtime()
void test_s_memtime(global ulong* out)
{
*out = __builtin_amdgcn_s_memtime();
}
// CHECK-LABEL: @test_s_memrealtime
// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
void test_s_memrealtime(global ulong* out)
{
*out = __builtin_amdgcn_s_memrealtime();
}
// CHECK-LABEL: @test_s_sleep
// CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
// CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
void test_s_sleep()
{
__builtin_amdgcn_s_sleep(1);
__builtin_amdgcn_s_sleep(15);
}
// 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) {

6
clang/test/SemaOpenCL/builtins-amdgcn.cl Normal file
View File

@ -0,0 +1,6 @@
// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -fsyntax-only -verify %s
void test_s_sleep(int x)
{
__builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
}