From f46f41411b1fa1360fb13235df77e60c8b6e8992 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Wed, 19 Jun 2019 19:55:49 +0000 Subject: [PATCH] Reapply "r363684: AMDGPU: Add GWS instruction builtins" llvm-svn: 363871 --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 ++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 12 ++++++++++++ 2 files changed, 14 insertions(+) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 89f68fbe82f4..a8aadff770d5 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -45,6 +45,8 @@ 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 bd7fe78fc96f..27a91a18cc97 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -548,6 +548,18 @@ 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 }