| 29 | |
| 30 | template<typename T> |
| 31 | Array<T> join(const int jdim, const Array<T> &first, const Array<T> &second) { |
| 32 | // All dimensions except join dimension must be equal |
| 33 | const dim4 &fdims{first.dims()}; |
| 34 | const dim4 &sdims{second.dims()}; |
| 35 | // Compute output dims |
| 36 | dim4 odims(fdims); |
| 37 | odims.dims[jdim] += sdims.dims[jdim]; |
| 38 | Array<T> out{createEmptyArray<T>(odims)}; |
| 39 | const cudaStream_t activeStream{getActiveStream()}; |
| 40 | |
| 41 | // topspeed is achieved when byte size(in+out) ~= L2CacheSize |
| 42 | // |
| 43 | // 1 array: memcpy always copies 1 array. topspeed |
| 44 | // --> size(in) < L2CacheSize/2 |
| 45 | // 2 arrays: topspeeds |
| 46 | // - size(in) < L2CacheSize/2/2 |
| 47 | // --> JIT can copy 2 arrays in // and is fastest |
| 48 | // (condition: array sizes have to be identical) |
| 49 | // - size(in) < L2CacheSize/2 |
| 50 | // --> memcpy will achieve highest speed, although the kernel |
| 51 | // has to be called twice |
| 52 | // - size(in) >= L2CacheSize/2 |
| 53 | // --> memcpy will achieve veryLargeArray speed. The kernel |
| 54 | // will be called twice |
| 55 | if (fdims.dims[jdim] == sdims.dims[jdim]) { |
| 56 | const size_t L2CacheSize{getL2CacheSize(getActiveDeviceId())}; |
| 57 | if (!(first.isReady() || second.isReady()) || |
| 58 | (fdims.elements() * sizeof(T) * 2 * 2 < L2CacheSize)) { |
| 59 | // Both arrays have same size & everything fits into the cache, |
| 60 | // so treat in 1 JIT kernel, iso individual copies which is |
| 61 | // always slower |
| 62 | const dim_t *outStrides{out.strides().dims}; |
| 63 | vector<Param<T>> outputs{ |
| 64 | {out.get(), fdims.dims, outStrides}, |
| 65 | {out.get() + fdims.dims[jdim] * outStrides[jdim], sdims.dims, |
| 66 | outStrides}}; |
| 67 | // Extend the life of the returned node, by saving the |
| 68 | // corresponding shared_ptr |
| 69 | const Node_ptr fNode{first.getNode()}; |
| 70 | const Node_ptr sNode{second.getNode()}; |
| 71 | vector<Node *> nodes{fNode.get(), sNode.get()}; |
| 72 | evalNodes(outputs, nodes); |
| 73 | return out; |
| 74 | } |
| 75 | // continue because individually processing is faster |
| 76 | } |
| 77 | |
| 78 | // Handle each array individually |
| 79 | if (first.isReady()) { |
| 80 | if (1LL + jdim >= first.ndims() && first.isLinear()) { |
| 81 | // first & out are linear |
| 82 | CUDA_CHECK(cudaMemcpyAsync(out.get(), first.get(), |
| 83 | first.elements() * sizeof(T), |
| 84 | cudaMemcpyDeviceToDevice, activeStream)); |
| 85 | } else { |
| 86 | kernel::memcopy<T>(out, first, first.ndims()); |
| 87 | } |
| 88 | } else { |
nothing calls this directly
no test coverage detected