5 files changed, 39 insertions, 20 deletions
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl
new file mode 100644
index 00000000000..89c3e490ecd
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-error.cl
@@ -0,0 +1,18 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s
+
+// FIXME: We only get one error if the functions are the other order in the
+// file.
+
+typedef unsigned long ulong;
+
+ulong test_s_memrealtime()
+{
+  return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+}
+
+void test_s_sleep(int x)
+{
+  __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
+}
+
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
new file mode 100644
index 00000000000..cda87a8e1ed
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned long ulong;
+
+
+// CHECK-LABEL: @test_s_memrealtime
+// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
+void test_s_memrealtime(global ulong* out)
+{
+  *out = __builtin_amdgcn_s_memrealtime();
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 11448c37442..36164f898bc 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -220,13 +220,6 @@ 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)
diff --git a/clang/test/CodeGenOpenCL/builtins-r600.cl b/clang/test/CodeGenOpenCL/builtins-r600.cl
index 6369f3d83db..9ebcb1fe9d0 100644
--- a/clang/test/CodeGenOpenCL/builtins-r600.cl
+++ b/clang/test/CodeGenOpenCL/builtins-r600.cl
@@ -1,7 +1,5 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu rv670 -S -emit-llvm -o - %s | FileCheck %s
-
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s
 
 // CHECK-LABEL: @test_rsq_f32
 // CHECK: call float @llvm.r600.rsq.f32
@@ -10,12 +8,14 @@ void test_rsq_f32(global float* out, float a)
   *out = __builtin_amdgpu_rsqf(a);
 }
 
-// CHECK-LABEL: @test_rsq_f64
-// CHECK: call double @llvm.r600.rsq.f64
+#if cl_khr_fp64
+// XCHECK-LABEL: @test_rsq_f64
+// XCHECK: call double @llvm.r600.rsq.f64
 void test_rsq_f64(global double* out, double a)
 {
   *out = __builtin_amdgpu_rsq(a);
 }
+#endif
 
 // CHECK-LABEL: @test_legacy_ldexp_f32
 // CHECK: call float @llvm.AMDGPU.ldexp.f32
@@ -24,9 +24,11 @@ void test_legacy_ldexp_f32(global float* out, float a, int b)
   *out = __builtin_amdgpu_ldexpf(a, b);
 }
 
-// CHECK-LABEL: @test_legacy_ldexp_f64
-// CHECK: call double @llvm.AMDGPU.ldexp.f64
+#if cl_khr_fp64
+// XCHECK-LABEL: @test_legacy_ldexp_f64
+// XCHECK: call double @llvm.AMDGPU.ldexp.f64
 void test_legacy_ldexp_f64(global double* out, double a, int b)
 {
   *out = __builtin_amdgpu_ldexp(a, b);
 }
+#endif
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn.cl b/clang/test/SemaOpenCL/builtins-amdgcn.cl
deleted file mode 100644
index bb342d058cb..00000000000
--- a/clang/test/SemaOpenCL/builtins-amdgcn.cl
+++ /dev/null
@@ -1,6 +0,0 @@
-// 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}}
-}