hughperkins

$ py.test -sv test/test_inline_refcount.py -k test2

============================= test session starts ==============================
platform linux2 -- Python 2.7.11+, pytest-2.9.2, py-1.4.31, pluggy-0.3.1 -- [deleted]/bin/python2.7
cachedir: .cache
rootdir: [deleted]/pytorch, inifile: pytest.ini
plugins: pep8-1.0.6
collecting ... collected 3 items

test/test_inline_refcount.py::test_inline_refcount_test2 <TestInlineRefCount> initialization

hughperkins

# Note – this is not a bash script (some of the steps require reboot)
# I named it .sh just so Github does correct syntax highlighting.
#
# This is also available as an AMI in us-east-1 (virginia): ami-cf5028a5
#
# The CUDA part is mostly based on this excellent blog post:
# http://tleyden.github.io/blog/2014/10/25/cuda-6-dot-5-on-aws-gpu-instance-running-ubuntu-14-dot-04/

# Install various packages
sudo apt-get update

hughperkins

(env3cl) ubuntu@peach:~/git/tensorflow-blas$ python tensorflow/stream_executor/cl/test/test_gradients.py
CLBlas::initialize_cublas()
Using NVIDIA Corporation , OpenCL platform: NVIDIA CUDA
Using OpenCL device: GeForce 940M
X [[ 1.  0.  1.  0.]
 [ 1.  0.  0.  1.]
 [ 0.  1.  1.  0.]
 [ 0.  1.  0.  1.]]
y [[ 1.  0.]
 [ 0.  1.]

hughperkins

I tensorflow/core/common_runtime/simple_placer.cc:819] Variable_1: /job:localhost/replica:0/task:0/cpu:0
Variable_1/read: /job:localhost/replica:0/task:0/cpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] Variable_1/read: /job:localhost/replica:0/task:0/cpu:0
Variable_1/Assign: /job:localhost/replica:0/task:0/cpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] Variable_1/Assign: /job:localhost/replica:0/task:0/cpu:0
Variable: /job:localhost/replica:0/task:0/cpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] Variable: /job:localhost/replica:0/task:0/cpu:0
Assign: /job:localhost/replica:0/task:0/cpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] Assign: /job:localhost/replica:0/task:0/cpu:0
Variable/read: /job:localhost/replica:0/task:0/cpu:0

hughperkins

kernel void _ZN5Eigen8internal15EigenMetaKernelINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi2ELi0ElEELi0ENS_11MakePointerEEEKNS_17TensorReductionOpINS0_10MaxReducerIfEEKNS_5arrayIlLm2EEEKNS4_INS5_IfLi4ELi0ElEELi0ES7_EEEEEENS_9GpuDeviceEEElEEvT_T0_(global struct Eigen__TensorEvaluator_5_nopointers* eval_nopointers, global float* eval_ptr0, long eval_ptr_offset0, global float* eval_ptr1, long eval_ptr_offset1, global float* eval_ptr2, long eval_ptr_offset2, long size) {
    float accum_0_i_i_i;
    float accum_1_i_i_i;
    float accum_1_i_i_i_unr;
    float accum_2_i_i_i_1_lcssa;
    float accum_2_i_i_i_lcssa;
    float accum_2_i_i_i_lcssa_unr;
    float accum_3_i_i_i;
    float accum_3_i_i_i_lcssa;
    long i_01_i;

hughperkins

inline float __shfl_down_3(float v0, int v1, int v2) {
    local float mem[1024];
    int tid = get_local_id(0);
    int warpid = tid % 32;
    int warpstart = tid - warpid;
    mem[tid] = v0;
    //barrier(CLK_LOCAL_MEM_FENCE);
    int warpsrc = warpid + v1;
    warpsrc = warpsrc >= 32 ? warpid : warpsrc;

hughperkins

inline float __shfl_down_3(float v0, int v1, int v2) {
    local float mem[1024];
    int tid = get_local_id(0);
    int warpid = tid % 32;
    int warpstart = tid - warpid;
    mem[tid] = v0;
    //barrier(CLK_LOCAL_MEM_FENCE);
    int warpsrc = warpid + v1;
    warpsrc = warpsrc >= 32 ? warpid : warpsrc;

hughperkins

W: /job:localhost/replica:0/task:0/gpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] W: /job:localhost/replica
:0/task:0/gpu:0
Assign: /job:localhost/replica:0/task:0/gpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] Assign: /job:localhost/re
plica:0/task:0/gpu:0
W/read: /job:localhost/replica:0/task:0/gpu:0
I tensorflow/core/common_runtime/simple_placer.cc:819] W/read: /job:localhost/re
plica:0/task:0/gpu:0
W/Assign: /job:localhost/replica:0/task:0/gpu:0

hughperkins

$ objdump -r build/test-cocl-cuda_sample.o | grep R_
000000000000000f R_X86_64_32       s.build/test-cocl-cuda_sample-device.cl._Z8setValuePfif
0000000000000014 R_X86_64_32       __opencl_sourcecodebuild/test-cocl-cuda_sample-device.cl
0000000000000019 R_X86_64_PC32     configureKernel-0x0000000000000004
0000000000000021 R_X86_64_PC32     _Z20setKernelArgCharStarPc-0x0000000000000004
0000000000000028 R_X86_64_PC32     _Z17setKernelArgInt32i-0x0000000000000004
0000000000000033 R_X86_64_PC32     _Z17setKernelArgFloatf-0x0000000000000004
000000000000003e R_X86_64_PC32     _Z8kernelGov-0x0000000000000004
000000000000006e R_X86_64_PC32     cudaMalloc-0x0000000000000004
0000000000000089 R_X86_64_PC32     cudaConfigureCall-0x0000000000000004

hughperkins

kernel void _ZN5Eigen8internal15EigenMetaKernelINS_15TensorEvaluatorIKNS_14TensorAssignOpINS_9TensorMapINS_6TensorIfLi1ELi1EiEELi16ENS_11MakePointerEEEKNS_18TensorCwiseUnaryOpINS0_14scalar_sqrt_opIfEEKNS4_INS5_IKfLi1ELi1EiEELi16ES7_EEEEEENS_9GpuDeviceEEEiEEvT_T0_(
        global char* eval_nopointers, global float* eval_ptr0, long eval_ptr0_offset, global float* eval_ptr1, long eval_ptr1_offset, int size, local int *scratch) {
    int i1_02_i;
    int i_03_i;
    int v2;
    int v5;
    int v7;
    int v9;
    int v11;
    int v23;