[AMDGPU] Add target feature gws to clang

[arsenm]

llvm/lib/TargetParser/TargetParser.cpp

289

Gds feature is unused

[yaxunl]

llvm/lib/TargetParser/TargetParser.cpp

289

I am thinking to keep it in case we need it for newly added builtins or want to diagnose ops requires gds.

[arsenm]

llvm/lib/TargetParser/TargetParser.cpp

289

But it's easy to add back in if and when that happens

[yaxunl]

remove gds feature since it is not used

clang/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n")
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_sched_barrier, "vIi", "n")
 BUILTIN(__builtin_amdgcn_sched_group_barrier, "vIiIiIi", "n")
 BUILTIN(__builtin_amdgcn_iglp_opt, "vIi", "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")
-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")
 BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
 BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n")
 
 BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n")
 BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n")
 
 BUILTIN(__builtin_amdgcn_atomic_dec32, "UZiUZiD*UZiUicC*", "n")
 BUILTIN(__builtin_amdgcn_atomic_dec64, "UWiUWiD*UWiUicC*", "n")

 
 //===----------------------------------------------------------------------===//
 // Flat addressing builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_is_shared, "bvC*0", "nc")
 BUILTIN(__builtin_amdgcn_is_private, "bvC*0", "nc")
 
 //===----------------------------------------------------------------------===//
+// GWS builtins.
+//===----------------------------------------------------------------------===//
+TARGET_BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n", "gws")
+TARGET_BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n", "gws")
+TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n", "gws")
+TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n", "gws")
+TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n", "gws")
+
+//===----------------------------------------------------------------------===//
 // CI+ only builtins.
 //===----------------------------------------------------------------------===//
 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.

clang/lib/Basic/Targets/AMDGPU.cpp

if (getMaxPointerWidth() == 64) {
     LongWidth = LongAlign = 64;
     SizeType = UnsignedLong;
     PtrDiffType = SignedLong;
     IntPtrType = SignedLong;
   }
 
   MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
   CUMode = !(GPUFeatures & llvm::AMDGPU::FEATURE_WGP);
-  ReadOnlyFeatures.insert("image-insts");
+  for (auto F : {"image-insts", "gws"})
+    ReadOnlyFeatures.insert(F);
   HalfArgsAndReturns = true;
 }
 
 void AMDGPUTargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts) {
   TargetInfo::adjust(Diags, Opts);
   // ToDo: There are still a few places using default address space as private
   // address space in OpenCL, which needs to be cleaned up, then Opts.OpenCL
   // can be removed from the following line.

clang/test/CodeGenOpenCL/amdgpu-features-readonly.cl

+// RUN: %clang_cc1 -triple amdgcn -target-feature +gws -o /dev/null %s 2>&1 \
+// RUN:   | FileCheck --check-prefix=GWS %s
+
+// GWS: warning: feature flag '+gws' is ignored since the feature is read only [-Winvalid-command-line-argument]
+
+kernel void test() {}

clang/test/CodeGenOpenCL/builtins-amdgcn-gws-insts.cl

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90c -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx940 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1030 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned int uint;
+
+// CHECK-LABEL: define dso_local amdgpu_kernel void @test_builtins_amdgcn_gws_insts
+// CHECK-SAME: (i32 noundef [[A:%.*]], i32 noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.init(i32 [[A]], i32 [[B]])
+// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.barrier(i32 [[A]], i32 [[B]])
+// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.sema.v(i32 [[A]])
+// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.sema.br(i32 [[A]], i32 [[B]])
+// CHECK-NEXT:    tail call void @llvm.amdgcn.ds.gws.sema.p(i32 [[A]])
+// CHECK-NEXT:    ret void
+//
+kernel void test_builtins_amdgcn_gws_insts(uint a, uint b) {
+  __builtin_amdgcn_ds_gws_init(a, b);
+  __builtin_amdgcn_ds_gws_barrier(a, b);
+  __builtin_amdgcn_ds_gws_sema_v(a);
+  __builtin_amdgcn_ds_gws_sema_br(a, b);
+  __builtin_amdgcn_ds_gws_sema_p(a);
+}

flang/test/Lower/OpenMP/target_cpu_features.f90

 !REQUIRES: amdgpu-registered-target
 !RUN: %flang_fc1 -emit-fir -triple amdgcn-amd-amdhsa -target-cpu gfx908 -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s
 
 !===============================================================================
 ! Target_Enter Simple
 !===============================================================================
 
 !CHECK: omp.target = #omp.target<target_cpu = "gfx908",
 !CHECK-SAME: target_features = "+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,
 !CHECK-SAME: +dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,
-!CHECK-SAME: +gfx8-insts,+gfx9-insts,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,
+!CHECK-SAME: +gfx8-insts,+gfx9-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,
 !CHECK-SAME: +wavefrontsize64">
 !CHECK-LABEL: func.func @_QPomp_target_simple()
 subroutine omp_target_simple
   ! Directive needed to prevent subroutine from being filtered out when
   ! compiling for the device.
   !$omp declare target
 end subroutine omp_target_simple
 

llvm/lib/TargetParser/TargetParser.cpp

case GK_GFX1100:
       Features["dpp"] = true;
       Features["gfx8-insts"] = true;
       Features["gfx9-insts"] = true;
       Features["gfx10-insts"] = true;
       Features["gfx10-3-insts"] = true;
       Features["gfx11-insts"] = true;
       Features["atomic-fadd-rtn-insts"] = true;
       Features["image-insts"] = true;
+      Features["gws"] = true;
       break;
     case GK_GFX1036:
     case GK_GFX1035:
     case GK_GFX1034:
     case GK_GFX1033:
     case GK_GFX1032:
     case GK_GFX1031:
     case GK_GFX1030:
case GK_GFX1030:
       Features["dpp"] = true;
       Features["gfx8-insts"] = true;
       Features["gfx9-insts"] = true;
       Features["gfx10-insts"] = true;
       Features["gfx10-3-insts"] = true;
       Features["image-insts"] = true;
       Features["s-memrealtime"] = true;
       Features["s-memtime-inst"] = true;
+      Features["gws"] = true;
       break;
     case GK_GFX1012:
     case GK_GFX1011:
       Features["dot1-insts"] = true;
       Features["dot2-insts"] = true;
       Features["dot5-insts"] = true;
       Features["dot6-insts"] = true;
       Features["dot7-insts"] = true;
       Features["dot10-insts"] = true;
       [[fallthrough]];
     case GK_GFX1013:
     case GK_GFX1010:
       Features["dl-insts"] = true;
       Features["ci-insts"] = true;
       Features["16-bit-insts"] = true;
       Features["dpp"] = true;
       Features["gfx8-insts"] = true;
       Features["gfx9-insts"] = true;
       Features["gfx10-insts"] = true;
       Features["image-insts"] = true;
       Features["s-memrealtime"] = true;
       Features["s-memtime-inst"] = true;
+      Features["gws"] = true;
       break;
     case GK_GFX942:
     case GK_GFX941:
     case GK_GFX940:
       Features["gfx940-insts"] = true;
       Features["fp8-insts"] = true;
       Features["atomic-ds-pk-add-16-insts"] = true;
       Features["atomic-flat-pk-add-16-insts"] = true;
case GK_GFX940:
       Features["dot10-insts"] = true;
       Features["gfx9-insts"] = true;
       Features["gfx8-insts"] = true;
       Features["16-bit-insts"] = true;
       Features["dpp"] = true;
       Features["s-memrealtime"] = true;
       Features["ci-insts"] = true;
       Features["s-memtime-inst"] = true;
+      Features["gws"] = true;
       break;
     case GK_GFX90A:
       Features["gfx90a-insts"] = true;
       Features["atomic-buffer-global-pk-add-f16-insts"] = true;
       Features["atomic-fadd-rtn-insts"] = true;
       [[fallthrough]];
     case GK_GFX908:
       Features["dot3-insts"] = true;
if (T.isAMDGCN()) {
     case GK_GFX700:
       Features["ci-insts"] = true;
       [[fallthrough]];
     case GK_GFX602:
     case GK_GFX601:
     case GK_GFX600:
       Features["image-insts"] = true;
       Features["s-memtime-inst"] = true;
+      Features["gws"] = true;
       break;
     case GK_NONE:
       break;
     default:
       llvm_unreachable("Unhandled GPU!");
     }
   } else {
     if (GPU.empty())