           auto z = reinterpret_cast<Z*>(vz);
 
     __shared__ int xRank, yRank, zRank;
-    __shared__ Nd4jLong zLen, totalThreads;  // xLen == zLen, except when xRank = 1, in this case zLen = 2*xLen
+    __shared__ Nd4jLong zLen, totalThreads, *sharedMem;  // xLen == zLen, except when xRank = 1, in this case zLen = 2*xLen
 
     if (threadIdx.x == 0) {
+        extern __shared__ unsigned char shmem[];
+        sharedMem = reinterpret_cast<Nd4jLong*>(shmem);
+
         xRank = shape::rank(xShapeInfo);
         yRank = shape::rank(yShapeInfo);
         zRank = shape::rank(zShapeInfo);

void TrueBroadcastBoolHelper<X,Y>::exec(const nd4j::broadcast::BoolOps opNum, const NDArray& xArr, const NDArray& yArr, NDArray& zArr) {

    dim3 launchDims;
    launchDims.x = 128; //MAX_NUM_THREADS / 8;   // threadsPerBlock
    launchDims.y = 256; //(zArr.lengthOf() + launchDims.x - 1) / launchDims.x;  // blocksPerGrid
    launchDims.z = 1024; // sharedMem

    PointersManager manager(xArr.getContext(), "TrueBroadcastBoolHelper<X,Y>::exec");


    }
    __syncthreads();

    auto xCoords = sharedMem + threadIdx.x * (xRank + yRank + zRank);
    auto yCoords = xCoords + xRank;
    auto zCoords = yCoords + yRank;

    const auto tid = blockIdx.x * blockDim.x + threadIdx.x;

