diff options
| author | Simon Pilgrim <llvm-dev@redking.me.uk> | 2019-06-19 15:35:45 +0000 |
|---|---|---|
| committer | Simon Pilgrim <llvm-dev@redking.me.uk> | 2019-06-19 15:35:45 +0000 |
| commit | 6828bc561495172c8b26bd56ada6051f2c65bd59 (patch) | |
| tree | a80c29135cb15e27a5a2113f1e88b786fac6c472 /clang | |
| parent | a644b04b8cd80e5a7aa4185610fe01570b9827de (diff) | |
| download | bcm5719-llvm-6828bc561495172c8b26bd56ada6051f2c65bd59.tar.gz bcm5719-llvm-6828bc561495172c8b26bd56ada6051f2c65bd59.zip | |
Revert rL363684 : AMDGPU: Add GWS instruction builtins
........
Depends on rL363678 which was reverted at rL363797
llvm-svn: 363824
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 12 |
2 files changed, 0 insertions, 14 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index a8aadff770d..89f68fbe82f 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -45,8 +45,6 @@ BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n") BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n") -BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n") -BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n") // FIXME: Need to disallow constant address space. BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 27a91a18cc9..bd7fe78fc96 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -548,18 +548,6 @@ kernel void test_ds_consume_lds(global int* out, local int* ptr) { *out = __builtin_amdgcn_ds_consume(ptr); } -// CHECK-LABEL: @test_gws_init( -// CHECK: call void @llvm.amdgcn.ds.gws.init(i32 %value, i32 %id) -kernel void test_gws_init(uint value, uint id) { - __builtin_amdgcn_ds_gws_init(value, id); -} - -// CHECK-LABEL: @test_gws_barrier( -// CHECK: call void @llvm.amdgcn.ds.gws.barrier(i32 %value, i32 %id) -kernel void test_gws_barrier(uint value, uint id) { - __builtin_amdgcn_ds_gws_barrier(value, id); -} - // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024} // CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly } // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent } |
