Author: jvesely
Date: Wed Aug 16 10:08:56 2017
New Revision: 311021
URL: http://llvm.org/viewvc/llvm-project?rev=311021&view=rev
Log:
amdgcn: Implement {read_,write_,}mem_fence builtin
v2: add more detailed comment about waitcnt instruction
Signed-off-by: Jan Vesely <jan.ves...@rutgers.edu>
Reviewed-by: Aaron Watry <awa...@gmail.com>
Tested-by: Aaron Watry <awa...@gmail.com>
Added:
libclc/trunk/amdgcn/lib/mem_fence/
libclc/trunk/amdgcn/lib/mem_fence/fence.cl
libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll
libclc/trunk/generic/include/clc/explicit_fence/
libclc/trunk/generic/include/clc/explicit_fence/explicit_memory_fence.h
Modified:
libclc/trunk/amdgcn/lib/SOURCES
libclc/trunk/generic/include/clc/clc.h
Modified: libclc/trunk/amdgcn/lib/SOURCES
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/SOURCES?rev=311021&r1=311020&r2=311021&view=diff
==============================================================================
--- libclc/trunk/amdgcn/lib/SOURCES (original)
+++ libclc/trunk/amdgcn/lib/SOURCES Wed Aug 16 10:08:56 2017
@@ -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
Added: libclc/trunk/amdgcn/lib/mem_fence/fence.cl
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/mem_fence/fence.cl?rev=311021&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/mem_fence/fence.cl (added)
+++ libclc/trunk/amdgcn/lib/mem_fence/fence.cl Wed Aug 16 10:08:56 2017
@@ -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);
+}
Added: libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll?rev=311021&view=auto
==============================================================================
--- libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll (added)
+++ libclc/trunk/amdgcn/lib/mem_fence/waitcnt.ll Wed Aug 16 10:08:56 2017
@@ -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 }
Modified: libclc/trunk/generic/include/clc/clc.h
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/clc.h?rev=311021&r1=311020&r2=311021&view=diff
==============================================================================
--- libclc/trunk/generic/include/clc/clc.h (original)
+++ libclc/trunk/generic/include/clc/clc.h Wed Aug 16 10:08:56 2017
@@ -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>
Added: libclc/trunk/generic/include/clc/explicit_fence/explicit_memory_fence.h
URL:
http://llvm.org/viewvc/llvm-project/libclc/trunk/generic/include/clc/explicit_fence/explicit_memory_fence.h?rev=311021&view=auto
==============================================================================
--- libclc/trunk/generic/include/clc/explicit_fence/explicit_memory_fence.h
(added)
+++ libclc/trunk/generic/include/clc/explicit_fence/explicit_memory_fence.h Wed
Aug 16 10:08:56 2017
@@ -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);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
