Author: arsenm Date: Thu Jun 20 14:33:57 2019 New Revision: 363986 URL: http://llvm.org/viewvc/llvm-project?rev=363986&view=rev Log: AMDGPU: Add DS GWS sema builtins
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=363986&r1=363985&r2=363986&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Thu Jun 20 14:33:57 2019 @@ -47,6 +47,9 @@ BUILTIN(__builtin_amdgcn_s_dcache_inv, " BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n") BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n") +BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n") +BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n") +BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n") // FIXME: Need to disallow constant address space. BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") @@ -108,6 +111,7 @@ BUILTIN(__builtin_amdgcn_ds_consume, "ii //===----------------------------------------------------------------------===// TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts") TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts") +TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_release_all, "vUi", "n", "ci-insts") //===----------------------------------------------------------------------===// // Interpolation builtins. Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl?rev=363986&r1=363985&r2=363986&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl Thu Jun 20 14:33:57 2019 @@ -3,6 +3,8 @@ // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s +typedef unsigned int uint; + // CHECK-LABEL: @test_s_dcache_inv_vol // CHECK: call void @llvm.amdgcn.s.dcache.inv.vol( void test_s_dcache_inv_vol() @@ -17,3 +19,9 @@ void test_buffer_wbinvl1_vol() __builtin_amdgcn_buffer_wbinvl1_vol(); } +// CHECK-LABEL: @test_gws_sema_release_all( +// CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %id) +void test_gws_sema_release_all(uint id) +{ + __builtin_amdgcn_ds_gws_sema_release_all(id); +} Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=363986&r1=363985&r2=363986&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Thu Jun 20 14:33:57 2019 @@ -560,6 +560,24 @@ kernel void test_gws_barrier(uint value, __builtin_amdgcn_ds_gws_barrier(value, id); } +// CHECK-LABEL: @test_gws_sema_v( +// CHECK: call void @llvm.amdgcn.ds.gws.sema.v(i32 %id) +kernel void test_gws_sema_v(uint id) { + __builtin_amdgcn_ds_gws_sema_v(id); +} + +// CHECK-LABEL: @test_gws_sema_br( +// CHECK: call void @llvm.amdgcn.ds.gws.sema.br(i32 %value, i32 %id) +kernel void test_gws_sema_br(uint value, uint id) { + __builtin_amdgcn_ds_gws_sema_br(value, id); +} + +// CHECK-LABEL: @test_gws_sema_p( +// CHECK: call void @llvm.amdgcn.ds.gws.sema.p(i32 %id) +kernel void test_gws_sema_p(uint id) { + __builtin_amdgcn_ds_gws_sema_p(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 } Modified: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl?rev=363986&r1=363985&r2=363986&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl (original) +++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl Thu Jun 20 14:33:57 2019 @@ -1,8 +1,9 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s -void test_ci_biltins() +void test_ci_builtins() { __builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}} __builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}} + __builtin_amdgcn_ds_gws_sema_release_all(0); // expected-error {{'__builtin_amdgcn_ds_gws_sema_release_all' needs target feature ci-insts}} } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits