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

Proposed by Erik Schnetter
Status: Merged
Merged at revision: 49
Proposed branch: lp://staging/~schnetter/pocl/main
Merge into: lp://staging/~pocl/pocl/trunk
Diff against target: 10089 lines (+7828/-739)
136 files modified
.bzrignore (+1/-1)
README (+1/-0)
acinclude.m4 (+0/-271)
configure.ac (+7/-6)
examples/Makefile.am (+1/-1)
examples/example1/example1.c (+2/-1)
examples/example1/example1.cl (+6/-6)
examples/example2/example2.c (+1/-1)
examples/trig/Makefile.am (+31/-0)
examples/trig/trig (+210/-0)
examples/trig/trig.c (+97/-0)
examples/trig/trig.cl (+13/-0)
examples/trig/trig_exec.c (+158/-0)
include/_kernel.h (+1255/-13)
lib/CL/clBuildProgram.c (+1/-1)
lib/CL/clCreateKernel.c (+2/-1)
lib/CL/clEnqueueNDRangeKernel.c (+1/-1)
lib/CL/devices/native/native.c (+1/-1)
lib/CL/devices/pthread/pthread.c (+0/-317)
lib/CL/devices/pthread/pthread.h (+0/-100)
lib/CL/pocl_cl.h (+1/-1)
lib/kernel/Makefile.am (+114/-2)
lib/kernel/abs.cl (+35/-0)
lib/kernel/abs_diff.cl (+26/-0)
lib/kernel/acos.cl (+26/-0)
lib/kernel/acosh.cl (+26/-0)
lib/kernel/acospi.cl (+26/-0)
lib/kernel/add_sat.cl (+55/-0)
lib/kernel/all.cl (+142/-0)
lib/kernel/any.cl (+142/-0)
lib/kernel/as_type.cl (+171/-0)
lib/kernel/asin.cl (+26/-0)
lib/kernel/asinh.cl (+26/-0)
lib/kernel/asinpi.cl (+26/-0)
lib/kernel/atan.cl (+26/-0)
lib/kernel/atan2.cl (+26/-0)
lib/kernel/atan2pi.cl (+26/-0)
lib/kernel/atanh.cl (+26/-0)
lib/kernel/atanpi.cl (+26/-0)
lib/kernel/bitselect.cl (+33/-0)
lib/kernel/cbrt.cl (+26/-0)
lib/kernel/ceil.cl (+154/-0)
lib/kernel/clamp.cl (+28/-0)
lib/kernel/clz.cl (+75/-0)
lib/kernel/convert_type.cl (+180/-0)
lib/kernel/copysign.cl (+130/-0)
lib/kernel/cos.cl (+26/-0)
lib/kernel/cosh.cl (+26/-0)
lib/kernel/cospi.cl (+27/-0)
lib/kernel/cross.cl (+46/-0)
lib/kernel/degrees.cl (+26/-0)
lib/kernel/distance.cl (+26/-0)
lib/kernel/dot.cl (+58/-10)
lib/kernel/erf.cl (+26/-0)
lib/kernel/erfc.cl (+26/-0)
lib/kernel/exp.cl (+26/-0)
lib/kernel/exp10.cl (+26/-0)
lib/kernel/exp2.cl (+26/-0)
lib/kernel/expm1.cl (+26/-0)
lib/kernel/fabs.cl (+131/-0)
lib/kernel/fast_distance.cl (+26/-0)
lib/kernel/fast_length.cl (+28/-0)
lib/kernel/fast_normalize.cl (+29/-0)
lib/kernel/fdim.cl (+26/-0)
lib/kernel/floor.cl (+152/-0)
lib/kernel/fma.cl (+27/-0)
lib/kernel/fmax.cl (+154/-0)
lib/kernel/fmin.cl (+154/-0)
lib/kernel/fmod.cl (+26/-0)
lib/kernel/fract.cl (+26/-0)
lib/kernel/hadd.cl (+29/-0)
lib/kernel/hypot.cl (+27/-0)
lib/kernel/ilogb.cl (+26/-0)
lib/kernel/isequal.cl (+29/-0)
lib/kernel/isfinite.cl (+28/-0)
lib/kernel/isgreater.cl (+28/-0)
lib/kernel/isgreaterequal.cl (+28/-0)
lib/kernel/isinf.cl (+28/-0)
lib/kernel/isless.cl (+28/-0)
lib/kernel/islessequal.cl (+28/-0)
lib/kernel/islessgreater.cl (+28/-0)
lib/kernel/isnan.cl (+28/-0)
lib/kernel/isnormal.cl (+28/-0)
lib/kernel/isnotequal.cl (+29/-0)
lib/kernel/isordered.cl (+29/-0)
lib/kernel/isunordered.cl (+28/-0)
lib/kernel/ldexp.cl (+27/-0)
lib/kernel/length.cl (+26/-0)
lib/kernel/lgamma.cl (+26/-0)
lib/kernel/log.cl (+26/-0)
lib/kernel/log10.cl (+26/-0)
lib/kernel/log1p.cl (+26/-0)
lib/kernel/log2.cl (+26/-0)
lib/kernel/logb.cl (+26/-0)
lib/kernel/mad.cl (+26/-0)
lib/kernel/mad24.cl (+2/-5)
lib/kernel/mad_hi.cl (+30/-0)
lib/kernel/mad_sat.cl (+94/-0)
lib/kernel/max.cl (+31/-0)
lib/kernel/maxmag.cl (+26/-0)
lib/kernel/min.cl (+31/-0)
lib/kernel/minmag.cl (+26/-0)
lib/kernel/mix.cl (+27/-0)
lib/kernel/mul24.cl (+26/-0)
lib/kernel/mul_hi.cl (+86/-0)
lib/kernel/nan.cl (+49/-0)
lib/kernel/nextafter.cl (+26/-0)
lib/kernel/normalize.cl (+26/-0)
lib/kernel/pow.cl (+26/-0)
lib/kernel/pown.cl (+31/-0)
lib/kernel/powr.cl (+26/-0)
lib/kernel/radians.cl (+26/-0)
lib/kernel/remainder.cl (+26/-0)
lib/kernel/rhadd.cl (+29/-0)
lib/kernel/rint.cl (+26/-0)
lib/kernel/rootn.cl (+35/-0)
lib/kernel/rotate.cl (+31/-0)
lib/kernel/round.cl (+26/-0)
lib/kernel/rsqrt.cl (+26/-0)
lib/kernel/select.cl (+34/-0)
lib/kernel/sign.cl (+31/-0)
lib/kernel/signbit.cl (+26/-0)
lib/kernel/sin.cl (+26/-0)
lib/kernel/sinh.cl (+26/-0)
lib/kernel/sinpi.cl (+26/-0)
lib/kernel/smoothstep.cl (+31/-0)
lib/kernel/sqrt.cl (+122/-0)
lib/kernel/step.cl (+38/-0)
lib/kernel/sub_sat.cl (+55/-0)
lib/kernel/tan.cl (+26/-0)
lib/kernel/tanh.cl (+26/-0)
lib/kernel/tanpi.cl (+26/-0)
lib/kernel/templates.h (+1185/-0)
lib/kernel/tgamma.cl (+26/-0)
lib/kernel/trunc.cl (+26/-0)
lib/kernel/upsample.cl (+70/-0)
To merge this branch: bzr merge lp://staging/~schnetter/pocl/main
Reviewer Review Type Date Requested Status
Carlos Sánchez de La Lama Approve
Erik Schnetter Needs Resubmitting
Review via email: mp+80344@code.staging.launchpad.net

Description of the change

This branch implements many functions for an OpenCL run-time library. Most functions are forwarded to libc, some are implemented directly when building for x86.

To post a comment you must log in.
Revision history for this message
Carlos Sánchez de La Lama (csanchezdll) wrote :

Hi erik,

lot of functions implemented, great! :)

Just some comments:

1) fmax.cl, fmin.cl and sqrt.cl gave me compilations errors, similar to this:
../../../../src/pocl.schnetter/lib/kernel/fmax.cl:35:11: error: invalid conversion between ext-vector type 'float4' and '__attribute__((__vector_size__(4 * sizeof(float)))) float'
  return ((float4)__builtin_ia32_maxss(*(float4*)&a, *(float4*)&b)).s0;
          ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
../../../../src/pocl.schnetter/lib/kernel/fmax.cl:91:11: error: invalid conversion between ext-vector type 'double2' and '__attribute__((__vector_size__(2 * sizeof(double)))) double'
  return ((double2)__builtin_ia32_maxsd(*(double2*)&a, *(double2*)&b)).s0;
          ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
2 errors generated.

2) Makefile in trig directory is not generated (missing the appropiate line at the end of configure.ac).

3) If I add that line trig test still do not compile for me:
/bin/bash ../../libtool --tag=CC --mode=link gcc -g -O2 -o trig trig.o ../../lib/CL/libCL.la
libtool: link: gcc -g -O2 -o .libs/trig trig.o ../../lib/CL/.libs/libCL.so
trig.o: In function `main':
/home/csanchez/prj/ttagpu/bld/pocl.schnetter/examples/trig/../../../../src/pocl.schnetter/examples/trig/trig.c:73: undefined reference to `sqrtf'
/home/csanchez/prj/ttagpu/bld/pocl.schnetter/examples/trig/../../../../src/pocl.schnetter/examples/trig/trig.c:78: undefined reference to `exec_trig_kernel'
/home/csanchez/prj/ttagpu/bld/pocl.schnetter/examples/trig/../../../../src/pocl.schnetter/examples/trig/trig.c:72: undefined reference to `sinf'
/home/csanchez/prj/ttagpu/bld/pocl.schnetter/examples/trig/../../../../src/pocl.schnetter/examples/trig/trig.c:74: undefined reference to `tanf'
/home/csanchez/prj/ttagpu/bld/pocl.schnetter/examples/trig/../../../../src/pocl.schnetter/examples/trig/trig.c:70: undefined reference to `cosf'
collect2: ld returned 1 exit status
make[3]: *** [trig] Error 1

Can you please check those? Are they due to some version mismatch among us?

Carlos

lp://staging/~schnetter/pocl/main updated
51. By Erik Schnetter

Call Clang built-in functions instead of libc functions.
Remove gamma() which is not part of OpenCL.

52. By Erik Schnetter

Correct Makefile errors in trig example

53. By Erik Schnetter

Correct quoting errors in configure.ac.
Correct syntax errors in CL/devices.
Re-add vanished acinclude.m4.

Revision history for this message
Erik Schnetter (schnetter) wrote :

Carlos

I just merged with the current trunk; this has syntax errors in CL/devices. Look for "+++", which should be "++" instead.

I corrected the errors in configure.ac and Makefile.am.

I don't know why you see the link errors for sqrtf and friends. I'm now using Clang __builtin functions instead of calling libm -- maybe that helps. Please look at the current tip of my branch (revision 53).

I really hope that the history of my branch doesn't pollute the main history. It contains quite a few stray paths. (I'm also learning how to merge with bzr.)

Revision history for this message
Erik Schnetter (schnetter) wrote :

Okay, so I resubmit my branch for merging.

review: Needs Resubmitting
lp://staging/~schnetter/pocl/main updated
54. By Erik Schnetter

Re-add vanished CL/devices/pthread

55. By Erik Schnetter

Correct some syntax errors.
Introduce some optimisations that LLVM doesn't perform.

Revision history for this message
Carlos Sánchez de La Lama (csanchezdll) wrote :

Saw the "+++x", yet it compiled with it :o (annoyed)

I think it was taking it as monary + (positive) and then pre increment, so it worked. Fixed, thanks.

Well the linking are not so reprising, -lm was not in the linking line. What surprises me is you not getting those.

I will recheck. And don't worry about stray paths in the history, is there just to track changes, if we start having multiple branches and similar we will care about that but for now it is not so important.

Revision history for this message
Erik Schnetter (schnetter) wrote :

The +++x didn't compile for me. I think it may be interpreted as ++ + x, which (rightfully) doesn't compile.

I'm on Mac OSX Lion, using LLVM 3.0 via MacPorts.

lp://staging/~schnetter/pocl/main updated
56. By Erik Schnetter

Clean up implementation of math functions.
Begin to implement integer functions.

57. By Erik Schnetter

Merge

Revision history for this message
Carlos Sánchez de La Lama (csanchezdll) wrote :

Still same fmax/fmin errors, and some news in nan.cl:

../../../../src/pocl.schnetter/lib/kernel/nan.cl:44:49: error: invalid conversion between ext-vector type 'vtype' (aka 'double16') and
      'unsigned long __attribute__((ext_vector_type(16)))'
                sizeof(stype)==4 /* float */ ? (vtype)((utype)(~FLT_MANT_MASK & as_uint ( NAN)) | ((utype)FLT_MANT_MASK & a)) :
                ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
../../../../src/pocl.schnetter/lib/kernel/templates.h:301:28: note: expanded from:
  IMPLEMENT_EXPR_V_U(NAME, EXPR, double16, double, ulong16)
                           ^
../../../../src/pocl.schnetter/lib/kernel/templates.h:287:12: note: expanded from:
    return EXPR; \
           ^~~~

I tried in my Mac os X Lion, same errors. I compiled LLVM/clang from svn source, but should not be a difference AFAIK.

[prj:ttagpu]bash-3.2$ clang --version
clang version 3.0 (branches/release_30 142628)
Target: x86_64-apple-darwin11.2.0
Thread model: posix

lp://staging/~schnetter/pocl/main updated
58. By Erik Schnetter

Improve some integer functions.
Add more integer functions.

Revision history for this message
Erik Schnetter (schnetter) wrote :

The compiler is correct -- my code tries to cast ulong16 to double16, which is not possible. I wonder why this doesn't generate an error for me.

I'll test my branch on a recent Ubuntu.

lp://staging/~schnetter/pocl/main updated
59. By Erik Schnetter

Implement integer functions

60. By Erik Schnetter

Merge

61. By Erik Schnetter

Add some common functions

62. By Erik Schnetter

Add geometry functions.

63. By Erik Schnetter

Add relational functions

64. By Erik Schnetter

Work around compiler problem (?) in fmax/fmin

65. By Erik Schnetter

Work around compiler problems (?) in fmax, fmin
Correct mad_sat

Revision history for this message
Carlos Sánchez de La Lama (csanchezdll) wrote :

Ok, it was still needed two minor fixes in trig/Makefile.am to link the math library and the pthreads if used. Now everything works, merged.

review: Approve

Preview Diff

[H/L] Next/Prev Comment, [J/K] Next/Prev File, [N/P] Next/Prev Hunk
The diff is not available at this time. You can reload the page or download it.