[Libclc-dev] Missing barrier.cl in SOURCE list

Jin Wang wangjin.jane at gmail.com
Wed Sep 5 10:51:40 PDT 2012


Hi,

Recently I am using clang and libclc to compile OpenCL source code to PTX.
For every "barrier(CL_LOCAL_MEM_FENCE)" I wrote in OpenCL source, I got a
function entry call instead of ".bar sync 0" as expected (shown as follows):

.func barrier
(
.reg .b32 barrier_param_0
)
;

// Callseq Start 0
 {
.reg .b32 temp_param_reg;
// <end>}
 .reg .b32 param0;
mov.b32 param0, %r4;
 call.uni
barrier,
(
 param0
);
 //{
}// Callseq End 0


I believe this is caused by incorrect implementation of barrier function in
libclc. But then I found there is an implementation of barrier in
"libclc/ptx-nvidiacl/lib/synchronizatoin/barrier.cl". However, this file is
not included in "libclc/ptx-nvidiacl/lib/SOURCE". After I add a new line
"synchronization/barrier.cl" to the SOURCE file, everything works just
fine. Could you please take a look and make necessary changes? Thank you!

Best,

Jin Wang
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/libclc-dev/attachments/20120905/7a7ee593/attachment-0002.html>


More information about the Libclc-dev mailing list