[PATCH] D63366: AMDGPU: Add GWS instruction builtins
Matt Arsenault via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 14 15:25:33 PDT 2019
arsenm created this revision.
arsenm added reviewers: rampitec, b-sumner, yaxunl.
Herald added subscribers: t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely, kzhuravl.
https://reviews.llvm.org/D63366
Files:
include/clang/Basic/BuiltinsAMDGPU.def
test/CodeGenOpenCL/builtins-amdgcn.cl
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -548,6 +548,18 @@
*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 }
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -45,6 +45,8 @@
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")
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D63366.204864.patch
Type: text/x-patch
Size: 1458 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190614/42994fd4/attachment.bin>
More information about the cfe-commits
mailing list