[PATCH] D17515: AMDGPU: Add builtins for recently added intrinsics

Matt Arsenault via cfe-commits Mon, 22 Feb 2016 09:34:56 -0800

arsenm created this revision.
arsenm added a reviewer: tstellarAMD.
arsenm added a subscriber: cfe-commits.


http://reviews.llvm.org/D17515

Files:
  include/clang/Basic/BuiltinsAMDGPU.def
  test/CodeGenOpenCL/builtins-amdgcn.cl
  test/SemaOpenCL/builtins-amdgcn.cl

Index: test/SemaOpenCL/builtins-amdgcn.cl
===================================================================
--- /dev/null
+++ 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}}
+}
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ 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 @@
   __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) {
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -40,8 +40,18 @@
 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")

Index: test/SemaOpenCL/builtins-amdgcn.cl
===================================================================
--- /dev/null
+++ 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}}
+}
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ 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 @@
   __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) {
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -40,8 +40,18 @@
 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")

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

[PATCH] D17515: AMDGPU: Add builtins for recently added intrinsics

Reply via email to