From 452c4a719284d0384ecda8d391ec49aba012a0ff Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Sat, 27 Feb 2016 09:54:43 +0000 Subject: [PATCH] AMDGPU: Add builtins for recently added intrinsics git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@262126 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/BuiltinsAMDGPU.def | 10 ++++++++++ test/CodeGenOpenCL/builtins-amdgcn.cl | 25 +++++++++++++++++++++++++ test/SemaOpenCL/builtins-amdgcn.cl | 6 ++++++ 3 files changed, 41 insertions(+) create mode 100644 test/SemaOpenCL/builtins-amdgcn.cl diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index 180d7a3f9c..95b3d814a5 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -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") diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 78741b6d91..c8085311ba 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -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) { diff --git a/test/SemaOpenCL/builtins-amdgcn.cl b/test/SemaOpenCL/builtins-amdgcn.cl new file mode 100644 index 0000000000..bb342d058c --- /dev/null +++ b/test/SemaOpenCL/builtins-amdgcn.cl @@ -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}} +} -- 2.40.0