[libclc] 05/92: amdgcn: Implement {read_, write_, }mem_fence builtin

[Andreas Boll]

This is an automated email from the git hooks/post-receive script.

aboll-guest pushed a commit to branch master
in repository libclc.

commit 6633bab2dadce648338e41cd78e4663395850408
Author: Jan Vesely <jan.vesely at rutgers.edu>
Date:   Wed Aug 16 17:08:56 2017 +0000

    amdgcn: Implement {read_,write_,}mem_fence builtin
    
    v2: add more detailed comment about waitcnt instruction
    
    Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
    Reviewed-by: Aaron Watry <awatry at gmail.com>
    Tested-by: Aaron Watry <awatry at gmail.com>
    
    git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@311021 91177308-0d34-0410-b5e6-96231b3b80d8
---
 amdgcn/lib/SOURCES                                 |  2 ++
 amdgcn/lib/mem_fence/fence.cl                      | 39 ++++++++++++++++++++++
 amdgcn/lib/mem_fence/waitcnt.ll                    | 11 ++++++
 generic/include/clc/clc.h                          |  3 ++
 .../clc/explicit_fence/explicit_memory_fence.h     |  3 ++
 5 files changed, 58 insertions(+)

diff --git a/amdgcn/lib/SOURCES b/amdgcn/lib/SOURCES
index 1ff5fd1..24f5949 100644
--- a/amdgcn/lib/SOURCES
+++ b/amdgcn/lib/SOURCES
@@ -1,4 +1,6 @@
 math/ldexp.cl
+mem_fence/fence.cl
+mem_fence/waitcnt.ll
 synchronization/barrier_impl.ll
 workitem/get_global_offset.cl
 workitem/get_group_id.cl
diff --git a/amdgcn/lib/mem_fence/fence.cl b/amdgcn/lib/mem_fence/fence.cl
new file mode 100644
index 0000000..408ffc3
--- /dev/null
+++ b/amdgcn/lib/mem_fence/fence.cl
@@ -0,0 +1,39 @@
+#include <clc/clc.h>
+
+void __clc_amdgcn_s_waitcnt(unsigned flags);
+
+// s_waitcnt takes 16bit argument with a combined number of maximum allowed
+// pending operations:
+// [12:8] LGKM -- LDS, GDS, Konstant (SMRD), Messages
+// [7] -- undefined
+// [6:4] -- exports, GDS, and mem write
+// [3:0] -- vector memory operations
+
+// Newer clang supports __builtin_amdgcn_s_waitcnt
+#if __clang_major__ >= 5
+#  define __waitcnt(x) __builtin_amdgcn_s_waitcnt(x)
+#else
+#  define __waitcnt(x) __clc_amdgcn_s_waitcnt(x)
+#endif
+
+_CLC_DEF void mem_fence(cl_mem_fence_flags flags)
+{
+	if (flags & CLK_GLOBAL_MEM_FENCE) {
+		// scalar loads are counted with LGKM but we don't know whether
+		// the compiler turned any loads to scalar
+		__waitcnt(0);
+	} else if (flags & CLK_LOCAL_MEM_FENCE)
+		__waitcnt(0xff); // LGKM is [12:8]
+}
+#undef __waitcnt
+
+// We don't have separate mechanism for read and write fences
+_CLC_DEF void read_mem_fence(cl_mem_fence_flags flags)
+{
+	mem_fence(flags);
+}
+
+_CLC_DEF void write_mem_fence(cl_mem_fence_flags flags)
+{
+	mem_fence(flags);
+}
diff --git a/amdgcn/lib/mem_fence/waitcnt.ll b/amdgcn/lib/mem_fence/waitcnt.ll
new file mode 100644
index 0000000..8be7f18
--- /dev/null
+++ b/amdgcn/lib/mem_fence/waitcnt.ll
@@ -0,0 +1,11 @@
+declare void @llvm.amdgcn.s.waitcnt(i32) #0
+
+; Export waitcnt intrinsic for clang < 5
+define void @__clc_amdgcn_s_waitcnt(i32 %flags) #1 {
+entry:
+  tail call void @llvm.amdgcn.s.waitcnt(i32 %flags)
+  ret void
+}
+
+attributes #0 = { nounwind }
+attributes #1 = { nounwind alwaysinline }
diff --git a/generic/include/clc/clc.h b/generic/include/clc/clc.h
index 5130632..deb9d70 100644
--- a/generic/include/clc/clc.h
+++ b/generic/include/clc/clc.h
@@ -179,6 +179,9 @@
 #include <clc/synchronization/cl_mem_fence_flags.h>
 #include <clc/synchronization/barrier.h>
 
+/* 6.11.9 Explicit Memory Fence Functions */
+#include <clc/explicit_fence/explicit_memory_fence.h>
+
 /* 6.11.10 Async Copy and Prefetch Functions */
 #include <clc/async/async_work_group_copy.h>
 #include <clc/async/async_work_group_strided_copy.h>
diff --git a/generic/include/clc/explicit_fence/explicit_memory_fence.h b/generic/include/clc/explicit_fence/explicit_memory_fence.h
new file mode 100644
index 0000000..8e046b1
--- /dev/null
+++ b/generic/include/clc/explicit_fence/explicit_memory_fence.h
@@ -0,0 +1,3 @@
+_CLC_DECL void mem_fence(cl_mem_fence_flags flags);
+_CLC_DECL void read_mem_fence(cl_mem_fence_flags flags);
+_CLC_DECL void write_mem_fence(cl_mem_fence_flags flags);

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/libclc.git