pocl

Merge lp://staging/~schnetter/pocl/main into lp://staging/~pocl/pocl/trunk

main
Merge into trunk

Proposed by Erik Schnetter on 2011-12-06

Status:	Merged
Merge reported by:	Pekka Jääskeläinen
Merged at revision:	not available
Proposed branch:	lp://staging/~schnetter/pocl/main
Merge into:	lp://staging/~pocl/pocl/trunk
Diff against target:	3735 lines (+2544/-357) 56 files modified .bzrignore (+4/-0) README (+10/-5) configure.ac (+2/-0) examples/Makefile.am (+1/-1) examples/barriers/barriers.c (+1/-1) examples/example1/example1.c (+3/-1) examples/example1/example1_exec.c (+1/-1) examples/example2/example2.c (+1/-1) examples/example2a/example2a.c (+1/-1) examples/forloops/forloops.c (+1/-1) examples/kernel/Makefile.am (+31/-0) examples/kernel/kernel.c (+93/-0) examples/kernel/test_bitselect.cl (+1176/-0) examples/kernel/test_fabs.cl (+178/-0) examples/kernel/test_rotate.cl (+204/-0) examples/loopbarriers/loopbarriers.c (+1/-1) examples/run_all.sh (+1/-1) examples/scalarwave/Makefile.am (+31/-0) examples/scalarwave/scalarwave.c (+233/-0) examples/scalarwave/scalarwave.cl (+94/-0) examples/trig/trig_exec.c (+1/-1) include/_kernel.h (+11/-10) lib/CL/Makefile.am (+1/-0) lib/CL/clCreateBuffer.c (+7/-2) lib/CL/clCreateKernel.c (+10/-8) lib/CL/clEnqueueNDRangeKernel.c (+29/-1) lib/CL/clFinish.c (+31/-0) lib/CL/clGetDeviceIDs.c (+4/-5) lib/CL/clReleaseKernel.c (+6/-1) lib/CL/clReleaseMemObject.c (+1/-1) lib/CL/clSetKernelArg.c (+7/-1) lib/CL/devices/native/native.c (+6/-43) lib/CL/devices/native/native.h (+3/-3) lib/CL/devices/pthread/pthread.c (+27/-54) lib/CL/devices/pthread/pthread.h (+7/-7) lib/CL/pocl_cl.h (+3/-3) lib/kernel/Makefile.am (+5/-3) lib/kernel/arm/Makefile.am (+5/-3) lib/kernel/clz.cl (+1/-1) lib/kernel/get_global_id.c (+8/-4) lib/kernel/get_global_offset.c (+39/-0) lib/kernel/get_global_size.c (+1/-1) lib/kernel/get_group_id.c (+1/-1) lib/kernel/get_local_id.c (+4/-4) lib/kernel/get_local_size.c (+39/-0) lib/kernel/get_num_groups.c (+4/-4) lib/kernel/get_work_dim.c (+31/-0) lib/kernel/popcount.cl (+35/-0) lib/kernel/signbit.cl (+34/-3) lib/kernel/sources.mk (+4/-0) lib/kernel/tce/Makefile.am (+5/-3) lib/kernel/templates.h (+1/-2) lib/kernel/x86_64/Makefile.am (+5/-3) lib/kernel/x86_64/copysign.cl (+0/-169) lib/llvmopencl/Workgroup.cc (+70/-2) tests/testsuite.at (+31/-0)
To merge this branch:	bzr merge lp://staging/~schnetter/pocl/main
Related bugs:	Link a bug report

Reviewer	Date Requested	Status
Pekka Jääskeläinen	2011-12-06	Approve on 2011-12-14
Erik Schnetter		Needs Resubmitting on 2011-12-14
Review via email: mp+84682@code.staging.launchpad.net

Description of the change

This branch passes the same test cases as the trunk: all on Ubuntu, and four failures on Mac OSX.

This branch also introduces a new test case scalarwave, which fails everywhere because of a real problem (that I have not debugged yet). If you want, I can deactivate this new test case, to not introduce a failing test case into the trunk.

Revision history for this message

Pekka Jääskeläinen (pekka-jaaskelainen) wrote on 2011-12-06:

get_global_offset() implementation seems wrong. I think it needs the separate values from clEnqueueNDRangeKernel() which are not stored yet in the context struct, IIRC.

In my opinion having failing tests in the test suite that reproduce (reported) bugs is fine but they should be marked with XFAIL (with a bug id maybe?) or similar until they are fixed so we won't get confused what is a regression and what a known bug reproduction (or if you are into TDD, just unimplemented features).

Other than these, seems good to merge for me.

review: Needs Fixing

Revision history for this message

Erik Schnetter (schnetter) wrote on 2011-12-07:

Indeed; I misunderstood the meaning of global_offset.

-erik

2011/12/6 Pekka Jääskeläinen <email address hidden>

> Review: Needs Fixing
>
> get_global_offset() implementation seems wrong. I think it needs the
> separate values from clEnqueueNDRangeKernel() which are not stored yet in
> the context struct, IIRC.
>
> In my opinion having failing tests in the test suite that reproduce
> (reported) bugs is fine but they should be marked with XFAIL (with a bug id
> maybe?) or similar until they are fixed so we won't get confused what is a
> regression and what a known bug reproduction (or if you are into TDD, just
> unimplemented features).
>
> Other than these, seems good to merge for me.
> --
> https://code.launchpad.net/~schnetter/pocl/main/+merge/84682
> You are the owner of lp:~schnetter/pocl/main.
>

--
Erik Schnetter <email address hidden> http://www.cct.lsu.edu/~eschnett/

lp://staging/~schnetter/pocl/main updated on 2011-12-07

105. By Erik Schnetter on 2011-12-07

Correct implementation of get_global_offset

Introduce global variables to hold the work_dim and global_offset.
Set these variables.
Correct error in setting num_groups.

Disable scalarwave test case (could have used XFAIL instead, but I
don't know how to do this).

106. By Erik Schnetter on 2011-12-07

Ignore two auto-generated files

Revision history for this message

Erik Schnetter (schnetter) wrote on 2011-12-07:

The current head of the branch corrects the global offset issue.

I didn't implement XFAIL since I didn't know how to do this.

review: Needs Resubmitting

lp://staging/~schnetter/pocl/main updated on 2011-12-07

107. By Erik Schnetter on 2011-12-07: Implement get_work_dim() properly
108. By Erik Schnetter on 2011-12-07: More error checking when setting kernel arguments

Revision history for this message

Pekka Jääskeläinen (pekka-jaaskelainen) wrote on 2011-12-07:

The matrix transpose test case seems to crash here. Some double free? Valgrind should help debugging it.

6: example2: matrix transpose *** glibc detected *** /home/visit0r/src/pocl/examples/example2/.libs/lt-example2: free(): invalid next size (fast): 0x0000000000f05cc0 ***
======= Backtrace: =========
/lib/libc.so.6(+0x71ad6)[0x2b2e98a19ad6]
/lib/libc.so.6(cfree+0x6c)[0x2b2e98a1e84c]
/home/visit0r/src/pocl/lib/CL/.libs/libCL.so.0(+0x45d7)[0x2b2e985605d7]
/lib/libpthread.so.0(+0x68ba)[0x2b2e987928ba]
/lib/libc.so.6(clone+0x6d)[0x2b2e98a7702d]

lp://staging/~schnetter/pocl/main updated on 2011-12-08

109. By Erik Schnetter on 2011-12-07

Add debug output to scalarwave example

110. By Erik Schnetter on 2011-12-07

Merge from trunk

111. By Erik Schnetter on 2011-12-07

Properly free partially allocated memory object if there is an error

112. By Erik Schnetter on 2011-12-07

Properly remove pointers from host_buffers list

Remove pointers from host_buffers list when the buffer is freed.
Insert pointers into host_buffers list more efficiently.

113. By Erik Schnetter on 2011-12-08

Correct memory allocation in kernel handling

Insert kernels into linked list more efficiently.

Temporarily: Don't deallocate kernels (ever), because this would lead
to dangling pointers. Needs more fixing.

114. By Erik Schnetter on 2011-12-08

Correct malloc problems in handling memory objects

Don't use a linked list to distinguish between user-supplied and
malloc'd pointers. Instead, use the OpenCL memory flags to distinguish
between these.

Increase size for /proc/cpuinfo buffer (necessary on my system).

115. By Erik Schnetter on 2011-12-08

Merge from trunk

116. By Erik Schnetter on 2011-12-08

Merge from trunk

Revision history for this message

Erik Schnetter (schnetter) wrote on 2011-12-09:

The matrix transpose fails because pocl doesn't recognise the local argument as local, and thus doesn't allocate memory for it. This seems Mac OSX specific, and is now bug 901519.

lp://staging/~schnetter/pocl/main updated on 2011-12-14

117. By Erik Schnetter on 2011-12-09: Merge from trunk: Correct many errors on Mac OSX
118. By Erik Schnetter on 2011-12-13: Merge from trunk
119. By Erik Schnetter on 2011-12-13: Recommend --enable-debug when configuring
120. By Erik Schnetter on 2011-12-14: Simplify code
121. By Erik Schnetter on 2011-12-14: Correct signbit implementation
122. By Erik Schnetter on 2011-12-14: Add printf prototype
123. By Erik Schnetter on 2011-12-14: Add beginnings of testsuite
124. By Erik Schnetter on 2011-12-14: Add missing #ifdef cl_khr_fp64
125. By Erik Schnetter on 2011-12-14: Ignore autogenerated file
126. By Erik Schnetter on 2011-12-14: Add all source files
127. By Erik Schnetter on 2011-12-14: "kernel" example: Small improvements
128. By Erik Schnetter on 2011-12-14: Correct Scalarwave example. Simplify code.
129. By Erik Schnetter on 2011-12-14: Activate Scalarwave test case

Revision history for this message

Erik Schnetter (schnetter) wrote on 2011-12-14:

My branch is now in a good state. All test cases pass, and my new scalar wave example (solving a PDE) is a new test case that also passes.

I also have the beginnings of test cases for the kernel library, but this is still disable because most tests fail (for known reasons, e.g. 3-element vectors don't work, and shift counts are not handled correctly; bug reports have been submitted to llvm).

review: Needs Resubmitting

Revision history for this message

Pekka Jääskeläinen (pekka-jaaskelainen) wrote on 2011-12-14:

I merged this but please fix the memory leak in clReleaseKernel().

review: Approve