yaxunl created this revision.
yaxunl added reviewers: b-sumner, arsenm, foad.
Herald added subscribers: StephenFan, kerbowa, hiraditya, tpr, dstuttard,
jvesely, kzhuravl.
Herald added a project: All.
yaxunl requested review of this revision.
Herald added a subscriber: wdng.
Herald added a project: LLVM.
https://reviews.llvm.org/D158367
Files:
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/lib/Basic/Targets/AMDGPU.cpp
clang/test/CodeGenOpenCL/amdgpu-features-readonly.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-gws-insts.cl
llvm/lib/TargetParser/TargetParser.cpp
Index: llvm/lib/TargetParser/TargetParser.cpp
===================================================================
--- llvm/lib/TargetParser/TargetParser.cpp
+++ llvm/lib/TargetParser/TargetParser.cpp
@@ -286,6 +286,8 @@
Features["gfx11-insts"] = true;
Features["atomic-fadd-rtn-insts"] = true;
Features["image-insts"] = true;
+ Features["gds"] = true;
+ Features["gws"] = true;
break;
case GK_GFX1036:
case GK_GFX1035:
@@ -311,6 +313,8 @@
Features["image-insts"] = true;
Features["s-memrealtime"] = true;
Features["s-memtime-inst"] = true;
+ Features["gds"] = true;
+ Features["gws"] = true;
break;
case GK_GFX1012:
case GK_GFX1011:
@@ -333,6 +337,8 @@
Features["image-insts"] = true;
Features["s-memrealtime"] = true;
Features["s-memtime-inst"] = true;
+ Features["gds"] = true;
+ Features["gws"] = true;
break;
case GK_GFX942:
case GK_GFX941:
@@ -362,6 +368,7 @@
Features["s-memrealtime"] = true;
Features["ci-insts"] = true;
Features["s-memtime-inst"] = true;
+ Features["gws"] = true;
break;
case GK_GFX90A:
Features["gfx90a-insts"] = true;
@@ -412,6 +419,8 @@
case GK_GFX600:
Features["image-insts"] = true;
Features["s-memtime-inst"] = true;
+ Features["gds"] = true;
+ Features["gws"] = true;
break;
case GK_NONE:
break;
Index: clang/test/CodeGenOpenCL/builtins-amdgcn-gws-insts.cl
===================================================================
--- /dev/null
+++ clang/test/CodeGenOpenCL/builtins-amdgcn-gws-insts.cl
@@ -0,0 +1,31 @@
+// 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);
+}
Index: clang/test/CodeGenOpenCL/amdgpu-features-readonly.cl
===================================================================
--- /dev/null
+++ clang/test/CodeGenOpenCL/amdgpu-features-readonly.cl
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -triple amdgcn -target-feature +gds -o /dev/null %s 2>&1 \
+// RUN: | FileCheck --check-prefix=GDS %s
+// RUN: %clang_cc1 -triple amdgcn -target-feature +gws -o /dev/null %s 2>&1 \
+// RUN: | FileCheck --check-prefix=GWS %s
+
+// GDS: warning: feature flag '+gds' is ignored since the feature is read only [-Winvalid-command-line-argument]
+// GWS: warning: feature flag '+gws' is ignored since the feature is read only [-Winvalid-command-line-argument]
+
+kernel void test() {}
Index: clang/lib/Basic/Targets/AMDGPU.cpp
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.cpp
+++ clang/lib/Basic/Targets/AMDGPU.cpp
@@ -244,7 +244,8 @@
MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
CUMode = !(GPUFeatures & llvm::AMDGPU::FEATURE_WGP);
- ReadOnlyFeatures.insert("image-insts");
+ for (auto F : {"image-insts", "gds", "gws"})
+ ReadOnlyFeatures.insert(F);
HalfArgsAndReturns = true;
}
Index: clang/include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -67,11 +67,6 @@
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")
@@ -172,6 +167,15 @@
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.
//===----------------------------------------------------------------------===//
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
