llvm-project

Commit Graph

Author	SHA1	Message	Date
Jeroen Ketema	4f5a3d5d6f	Make ptx barrier work irrespective of the cl_mem_fence_flags This generates a "bar.sync 0” instruction, which not only causes the threads to wait, but does acts as a memory fence, as required by OpenCL. The fence does not differentiate between local and global memory. Unfortunately, there is no similar instruction which does not include a memory fence. Hence, we cannot optimize the case where neither CLK_LOCAL_MEM_FENCE nor CLK_GLOBAL_MEM_FENCE is passed. llvm-svn: 315228	2017-10-09 18:36:48 +00:00
Jan Vesely	7846c9b8f0	ptx: Fix builtin names after clang r274770 Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu> Acked-By: Aaron Watry <awatry@gmail.com> llvm-svn: 276423	2016-07-22 15:00:08 +00:00
Peter Collingbourne	a385c53413	PTX: move implementations of work-item and synchronisation functions to lib, and add header files in generic. Incorporates a patch by Tom Stellard! llvm-svn: 161313	2012-08-05 22:25:37 +00:00

Author

SHA1

Message

Date

Jeroen Ketema

4f5a3d5d6f

Make ptx barrier work irrespective of the cl_mem_fence_flags

This generates a "bar.sync 0” instruction, which not only causes the
threads to wait, but does acts as a memory fence, as required by
OpenCL. The fence does not differentiate between local and global
memory. Unfortunately, there is no similar instruction which does
not include a memory fence. Hence, we cannot optimize the case
where neither CLK_LOCAL_MEM_FENCE nor CLK_GLOBAL_MEM_FENCE is
passed.

llvm-svn: 315228

2017-10-09 18:36:48 +00:00

Jan Vesely

7846c9b8f0

ptx: Fix builtin names after clang r274770

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 276423

2016-07-22 15:00:08 +00:00

Peter Collingbourne

a385c53413

PTX: move implementations of work-item and synchronisation functions

to lib, and add header files in generic.  Incorporates a patch by
Tom Stellard!

llvm-svn: 161313

2012-08-05 22:25:37 +00:00

3 Commits