aboutsummaryrefslogtreecommitdiff
path: root/test/CodeGenOpenCL
diff options
context:
space:
mode:
authorDaniil Fukalov <daniil.fukalov@amd.com>2018-05-21 16:18:07 +0000
committerDaniil Fukalov <daniil.fukalov@amd.com>2018-05-21 16:18:07 +0000
commita3f8c02f04ee414654417aebb6f848321c8d68ec (patch)
tree87ea02f70ae5bf79b7e31207f574bfe15c1e31e5 /test/CodeGenOpenCL
parent6b450481a2d4d66aa0b146a1d7fc1084c555488e (diff)
[AMDGPU] fixes for lds f32 builtins
1. added restrictions to memory scope, order and volatile parameters 2. added custom processing for these builtins - currently is not used code, needed to switch off GCCBuiltin link to the builtins (ongoing change to llvm tree) 3. builtins renamed as requested Differential Revision: https://reviews.llvm.org/D43281 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@332848 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenOpenCL')
-rw-r--r--test/CodeGenOpenCL/builtins-amdgcn-vi.cl12
1 files changed, 6 insertions, 6 deletions
diff --git a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index caa76e2374..afa312cfcb 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -91,18 +91,18 @@ void test_mov_dpp(global int* out, int src)
// CHECK-LABEL: @test_ds_fadd
// CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fadd(__attribute__((address_space(3))) float *out, float src) {
- *out = __builtin_amdgcn_ds_fadd(out, src, 0, 0, false);
+void test_ds_faddf(local float *out, float src) {
+ *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false);
}
// CHECK-LABEL: @test_ds_fmin
// CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fmin(__attribute__((address_space(3))) float *out, float src) {
- *out = __builtin_amdgcn_ds_fmin(out, src, 0, 0, false);
+void test_ds_fminf(local float *out, float src) {
+ *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
}
// CHECK-LABEL: @test_ds_fmax
// CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fmax(__attribute__((address_space(3))) float *out, float src) {
- *out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, false);
+void test_ds_fmaxf(local float *out, float src) {
+ *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
}