hughperkins

Allright; now linker errors!

[  1%] Building CXX object CMakeFiles/patch_hostside.dir/src/patch_hostside.cpp.o
[  2%] Building CXX object CMakeFiles/patch_hostside.dir/src/struct_clone.cpp.o
[  3%] Building CXX object CMakeFiles/patch_hostside.dir/src/mutations.cpp.o
[  4%] Building CXX object CMakeFiles/patch_hostside.dir/src/readIR.cpp.o
[  5%] Building CXX object CMakeFiles/patch_hostside.dir/third_party/argparsecpp/argparsecpp.cpp.o
[  6%] Building CXX object CMakeFiles/patch_hostside.dir/src/type_dumper.cpp.o
[  7%] Building CXX object CMakeFiles/patch_hostside.dir/src/GlobalNames.cpp.o
[  8%] Building CXX object CMakeFiles/patch_hostside.dir/src/EasyCL/util/easycl_stringhelper.cpp.o

hughperkins

Allright; now linker errors!

[  1%] Building CXX object CMakeFiles/patch_hostside.dir/src/patch_hostside.cpp.o
[  2%] Building CXX object CMakeFiles/patch_hostside.dir/src/struct_clone.cpp.o
[  3%] Building CXX object CMakeFiles/patch_hostside.dir/src/mutations.cpp.o
[  4%] Building CXX object CMakeFiles/patch_hostside.dir/src/readIR.cpp.o
[  5%] Building CXX object CMakeFiles/patch_hostside.dir/third_party/argparsecpp/argparsecpp.cpp.o
[  6%] Building CXX object CMakeFiles/patch_hostside.dir/src/type_dumper.cpp.o
[  7%] Building CXX object CMakeFiles/patch_hostside.dir/src/GlobalNames.cpp.o
[  8%] Building CXX object CMakeFiles/patch_hostside.dir/src/EasyCL/util/easycl_stringhelper.cpp.o

hughperkins

b [[ 1.27999997  1.96000004  1.88        0.36000001  1.27999997  1.84000003
   1.27999997  1.88        1.          0.75999999]]
expected [[ 3.07999992  2.03999996  3.          1.72000003  2.79999995  2.51999998
   2.03999996  3.55999994  1.88        2.07999992]]
gpu [[ 3.07999992  2.03999996  3.          1.72000003  2.79999995  2.51999998
   2.03999996  3.55999994  1.88        2.07999992]]
diff 0.0
func add float32
a [[ 1.79999995  0.08        1.12        1.36000001  1.51999998  0.68000001
   0.75999999  1.67999995  0.88        1.32000005]]

hughperkins

(env3) (tensorflow-cl|…10△2) ~/git/tensorflow-cl$ git diff
diff --git a/tensorflow/core/common_runtime/gpu/gpu_event_mgr.cc b/tensorflow/core/common_runtime/gpu/gpu_event_mgr.cc
index f18ee5e..ba664ea 100644
--- a/tensorflow/core/common_runtime/gpu/gpu_event_mgr.cc
+++ b/tensorflow/core/common_runtime/gpu/gpu_event_mgr.cc
@@ -18,6 +18,8 @@ limitations under the License.
 #include "tensorflow/core/platform/stream_executor.h"
 #include "tensorflow/core/protobuf/config.pb.h"

+#include <sstream>

hughperkins

#0  0x00007fff86ffb250 in ?? ()
#1  0x00007fffd91d8f27 in std::function<void ()>::function(std::function<void ()> const&) (this=0x7fff9f7fdac0,
    __x=...) at /usr/include/c++/5/functional:2238
#2  0x00007fffdb866067 in tensorflow::EventMgr::FreeMemory (this=0x13753e0, to_free=...)
    at ./tensorflow/core/common_runtime/gpu/gpu_event_mgr.h:127
#3  0x00007fffdbb0c029 in tensorflow::EventMgr::PollLoop (this=0x13753e0)
    at tensorflow/core/common_runtime/gpu/gpu_event_mgr.cc:140
#4  0x00007fffdbb0b9cd in tensorflow::EventMgr::<lambda()>::operator()(void) const (__closure=0x1409c00)
    at tensorflow/core/common_runtime/gpu/gpu_event_mgr.cc:76
#5  0x00007fffdbb0c5bb in std::_Function_handler<void(), tensorflow::EventMgr::StartPollingLoop()::<lambda()> >::_M_invoke(const std::_Any_data &) (__functor=...) at /usr/include/c++/5/functional:1871

hughperkins

#0  0x00007fff867fa2e0 in ?? ()
#1  0x00007fffd91d8c07 in std::function<void ()>::function(std::function<void ()> const&) ()
   from /home/ubuntu/env3/lib/python3.5/site-packages/tensorflow/python/_pywrap_tensorflow.so
#2  0x00007fffdb865d47 in tensorflow::EventMgr::FreeMemory(tensorflow::gtl::InlinedVector<tensorflow::EventMgr::InUse, 4> const&) () from /home/ubuntu/env3/lib/python3.5/site-packages/tensorflow/python/_pywrap_tensorflow.so
#3  0x00007fffdbb0bd09 in tensorflow::EventMgr::PollLoop() ()
   from /home/ubuntu/env3/lib/python3.5/site-packages/tensorflow/python/_pywrap_tensorflow.so
#4  0x00007fffdbb0b6ad in tensorflow::EventMgr::StartPollingLoop()::{lambda()#1}::operator()() const ()
   from /home/ubuntu/env3/lib/python3.5/site-packages/tensorflow/python/_pywrap_tensorflow.so
#5  0x00007fffdbb0c29b in std::_Function_handler<void (), tensorflow::EventMgr::StartPollingLoop()::{lambda()#1}>::_M_invoke(std::_Any_data const&) ()
   from /home/ubuntu/env3/lib/python3.5/site-packages/tensorflow/python/_pyw

hughperkins

// original kernelName: [_ZN10tensorflow12_GLOBAL__N_113SplitOpKernelIfEEvPKT_iiiNS_21CudaDeviceArrayStructIPS2_Li8EEE]
// unique kernelName: [_ZN10tensorflow12_GLOBAL__N_113SplitOpKernelIfEEvPKT_iiiNS_21CudaDeviceArrayStructIPS2_Li8EEE_0_1]
// short kernelname: [_ZN10tensorflow12_GL]

struct tensorflow__CudaDeviceArrayStruct {
    int f0;
    float* f1[8];
    global float** f2;
};

hughperkins

struct class_tensorflow__random__Array {
    int f0[4];
};
struct class_tensorflow__random__Array_0 {
    int f0[2];
};
struct class_tensorflow__random__NormalDistribution {
    char f0;
};
struct class_tensorflow__random__PhiloxRandom {

hughperkins

__internal__ build log: 
<program source>:209:27: warning: & has lower precedence than ==; == will be evaluated first
    v62 = v59 + (v51 == 0 & v50);
                 ~~~~~~~~~^
<program source>:209:27: note: place parentheses around the == expression to silence this warning
    v62 = v59 + (v51 == 0 & v50);
                          ^
                 (       )
<program source>:209:27: note: place parentheses around the & expression to evaluate it first
    v62 = v59 + (v51 == 0 & v50);

hughperkins

/* Minimal declarations for CUDA support.  Testing purposes only. */

#define __constant__ __attribute__((constant))
#define __device__ __attribute__((device))
#define __global__ extern "C" __attribute__((global))
#define __host__ __attribute__((host))
#define __shared__ __attribute__((shared))
#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
#define __forceinline__ __attribute__((always_inline))