intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-31 07:27:33 +08:00
Code Issues Packages Projects Releases Wiki Activity

Revert rL363684 : AMDGPU: Add GWS instruction builtins

Browse Source 
........
Depends on rL363678 which was reverted at rL363797

llvm-svn: 363824
This commit is contained in:
Simon Pilgrim
2019-06-19 15:35:45 +00:00
parent a644b04b8c
commit 6828bc5614
2 changed files with 0 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@@ -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")

12
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@@ -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 }