Reductions in CUDA use the thrust library for speed and have limited
       functionality.
542356
def reduce_cuda(g, x, axes, dtype):
    """"""
    if axes != 0:
        raise NotImplementedError("'axes' keyword is not implemented for CUDA")

    return g(x, dtype=dtype)
783513
def _run_chunk(self, chunk, kernel_options, device_options, tuning_options):
        """"""

        #detect language and create high-level device interface
        self.dev = DeviceInterface(kernel_options.kernel_string, iterations=tuning_options.iterations, **device_options)

        #move data to the GPU
        gpu_args = self.dev.ready_argument_list(kernel_options.arguments)

        results = []

        for element in chunk:
            params = dict(OrderedDict(zip(tuning_options.tune_params.keys(), element)))

            try:
                time = self.dev.compile_and_benchmark(gpu_args, params, kernel_options, tuning_options)

                params['time'] = time
                results.append(params)
            except Exception:
                params['time'] = None
                results.append(params)

        return results
424920
def main( gpu:Param("GPU to run on", str)=None ):
    """"""
    gpu = setup_distrib(gpu)
    n_gpus = num_distrib()
    path = url2path(URLs.CIFAR)
    ds_tfms = ([*rand_pad(4, 32), flip_lr(p=0.5)], [])
    workers = min(16, num_cpus()//n_gpus)
    data = ImageDataBunch.from_folder(path, valid='test', ds_tfms=ds_tfms, bs=512//n_gpus,
                                      num_workers=workers).normalize(cifar_stats)
    learn = Learner(data, wrn_22(), metrics=accuracy)
    if gpu is None: learn.model = nn.DataParallel(learn.model)
    else: learn.to_distributed(gpu)
    learn.to_fp16()
    learn.fit_one_cycle(35, 3e-3, wd=0.4)
461687
def xstep(self, S, W, lmbda, dimK):
        """"""

        if self.opt['CUDA_CBPDN']:
            Z = cuda.cbpdnmsk(self.D.squeeze(), S[..., 0], W.squeeze(), lmbda,
                              self.opt['CBPDN'])
            Z = Z.reshape(self.cri.Nv + (1, 1, self.cri.M,))
            self.Z[:] = np.asarray(Z, dtype=self.dtype)
            self.Zf = sl.rfftn(self.Z, self.cri.Nv, self.cri.axisN)
            self.Sf = sl.rfftn(S.reshape(self.cri.shpS), self.cri.Nv,
                               self.cri.axisN)
            self.xstep_itstat = None
        else:
            # Create X update object (external representation is expected!)
            xstep = cbpdn.ConvBPDNMaskDcpl(self.D.squeeze(), S, lmbda, W,
                                           self.opt['CBPDN'], dimK=dimK,
                                           dimN=self.cri.dimN)
            xstep.solve()
            self.Sf = sl.rfftn(S.reshape(self.cri.shpS), self.cri.Nv,
                               self.cri.axisN)
            self.setcoef(xstep.getcoef())
            self.xstep_itstat = xstep.itstat[-1] if xstep.itstat else None
1066221
def isemhash_bottleneck(x,
                        bottleneck_bits,
                        bottleneck_noise,
                        discretize_warmup_steps,
                        mode,
                        isemhash_noise_dev=0.5,
                        isemhash_mix_prob=0.5):
  """"""
  with tf.variable_scope("isemhash_bottleneck"):
    x = tf.layers.dense(x, bottleneck_bits, name="dense")
    y = common_layers.saturating_sigmoid(x)
    if isemhash_noise_dev > 0 and mode == tf.estimator.ModeKeys.TRAIN:
      noise = tf.truncated_normal(
          common_layers.shape_list(x), mean=0.0, stddev=isemhash_noise_dev)
      y = common_layers.saturating_sigmoid(x + noise)
    d = tf.to_float(tf.less(0.5, y)) + y - tf.stop_gradient(y)
    d = 2.0 * d - 1.0  # Move from [0, 1] to [-1, 1].
    if mode == tf.estimator.ModeKeys.TRAIN:  # Flip some bits.
      noise = tf.random_uniform(common_layers.shape_list(x))
      noise = 2.0 * tf.to_float(tf.less(bottleneck_noise, noise)) - 1.0
      d *= noise
      d = common_layers.mix(
          d,
          2.0 * y - 1.0,
          discretize_warmup_steps,
          mode == tf.estimator.ModeKeys.TRAIN,
          max_prob=isemhash_mix_prob)
    return d, 0.0
776553
def cublasDsyrk(handle, uplo, trans, n, k, alpha, A, lda, beta, C, ldc):
    
    
    status = _libcublas.cublasDsyrk_v2(handle,
                                       _CUBLAS_FILL_MODE[uplo], 
                                       _CUBLAS_OP[trans], 
                                       n, k, ctypes.byref(cuda.cuFloatComplex(alpha.real,     
                                                                              alpha.imag)),
                                       int(A), lda, 
                                       ctypes.byref(cuda.cuFloatComplex(beta.real,
                                                                        beta.imag)),
                                       int(C), ldc)
    cublasCheckStatus(status)
1158312
private static void compute(Kernel kernel, int[] in, int[] out, int width, int height, boolean alpha, int edge)
    {
        final float[] matrix = kernel.getMatrix();
        final int cols = kernel.getWidth();
        final int cols2 = cols / 2;

        for (int y = 0; y < height; y++)
        {
            int index = y;
            final int ioffset = y * width;
            for (int x = 0; x < width; x++)
            {
                compute(matrix, in, out, x, index, ioffset, cols2, width, alpha, edge);
                index += height;
            }
        }
    }
270034
def get_gpus_in_use(max_devices=None):
    
    from turicreate.util import _get_cuda_gpus
    gpu_indices = get_gpu_ids_in_use(max_devices=max_devices)
    gpus = _get_cuda_gpus()
    return [gpus[index] for index in gpu_indices]
988274
def cublasCher2k(handle, uplo, trans, n, k, alpha, A, lda, B, ldb, beta, C, ldc):
    
    
    status = _libcublas.cublasCher2k_v2(handle, 
                                        _CUBLAS_FILL_MODE[uplo], 
                                        _CUBLAS_OP[trans], 
                                        n, k, ctypes.byref(cuda.cuFloatComplex(alpha.real,                 
                                                                               alpha.imag)),
                                        int(A), lda, int(B), ldb, 
                                        ctypes.byref(cuda.cuFloatComplex(beta.real,
                                                                         beta.imag)),
                                        int(C), ldc)
    cublasCheckStatus(status)
1040019
@Override
    public void memcpySpecial(DataBuffer dstBuffer, Pointer srcPointer, long length, long dstOffset) {
        CudaContext context = getCudaContext();
        AllocationPoint point = ((BaseCudaDataBuffer) dstBuffer).getAllocationPoint();

        Pointer dP = new CudaPointer((point.getPointers().getHostPointer().address()) + dstOffset);

        val profH = PerformanceTracker.getInstance().helperStartTransaction();

        if (nativeOps.memcpyAsync(dP, srcPointer, length, CudaConstants.cudaMemcpyHostToHost, context.getOldStream()) == 0)
            throw new ND4JIllegalStateException("memcpyAsync failed");

        PerformanceTracker.getInstance().helperRegisterTransaction(point.getDeviceId(), profH, point.getNumberOfBytes(),MemcpyDirection.HOST_TO_HOST);

        if (point.getAllocationStatus() == AllocationStatus.DEVICE) {
            Pointer rDP = new CudaPointer(point.getPointers().getDevicePointer().address() + dstOffset);

            val profD = PerformanceTracker.getInstance().helperStartTransaction();

            if (nativeOps.memcpyAsync(rDP, dP, length, CudaConstants.cudaMemcpyHostToDevice, context.getOldStream()) == 0)
                throw new ND4JIllegalStateException("memcpyAsync failed");

            context.syncOldStream();

            PerformanceTracker.getInstance().helperRegisterTransaction(point.getDeviceId(), profD, point.getNumberOfBytes(),MemcpyDirection.HOST_TO_DEVICE);
        }

        context.syncOldStream();


        point.tickDeviceWrite();
    }
285258
public long compute(java.util.Vector<Long> e, java.util.Vector<java.util.List<Edge>> c, java.util.Vector<java.util.List<EdgeFlow>> x) {
        assert e.size() == c.size();
        assert x.size() == c.size();

        numNodes = e.size();
        nodesToQ = new java.util.Vector<Integer>();
        for (int i = 0; i < numNodes; i++) {
            nodesToQ.add(0);
        }

        // init flow
        for (int from = 0; from < numNodes; ++from) {
            for (Edge it : c.get(from)) {
                x.get(from).add(new EdgeFlow(it.getTo(), it.getCost(), 0));
                x.get(it.getTo()).add(new EdgeFlow(from, -it.getCost(), 0));
            }
        }

        // reduced costs for forward edges (c[i,j]-pi[i]+pi[j])
        // Note that for forward edges the residual capacity is infinity
        java.util.Vector<java.util.List<ReducedEdge>> rCostForward = new java.util.Vector<java.util.List<ReducedEdge>>();
        for (int i = 0; i < numNodes; i++) {
            rCostForward.add(new java.util.LinkedList<ReducedEdge>());
        }
        for (int from = 0; from < numNodes; ++from) {
            for (Edge it : c.get(from)) {
                rCostForward.get(from).add(new ReducedEdge(it.getTo(), it.getCost()));
            }
        }

        // reduced costs and capacity for backward edges
        // (c[j,i]-pi[j]+pi[i])
        // Since the flow at the beginning is 0, the residual capacity is
        // also zero
        java.util.Vector<java.util.List<ResidualReducedEdge>> rCostCapBackward = new java.util.Vector<java.util.List<ResidualReducedEdge>>();
        for (int i = 0; i < numNodes; i++) {
            rCostCapBackward.add(new java.util.LinkedList<ResidualReducedEdge>());
        }
        for (int from = 0; from < numNodes; ++from) {
            for (Edge it : c.get(from)) {
                rCostCapBackward.get(it.getTo()).add(
                        new ResidualReducedEdge(from, -it.getCost(), 0));
            }
        }

        // Max supply TODO:demand?, given U?, optimization-> min out of
        // demand,supply
        long u = 0;
        for (int i = 0; i < numNodes; i++) {
            if (e.get(i) > u) {
                u = e.get(i);
            }
        }
        //long delta = (long) (Math.pow(2.0, Math.ceil(Math.log((double) U) / Math.log(2.0))));

        java.util.Vector<Long> d = new java.util.Vector<Long>();
        java.util.Vector<Integer> prev = new java.util.Vector<Integer>();
        for (int i = 0; i < numNodes; i++) {
            d.add(0L);
            prev.add(0);
        }
        long delta = 1;
        long tmp;
        while (true) { // until we break when S or T is empty
            long maxSupply = 0;
            int k = 0;
            for (int i = 0; i < numNodes; i++) {
                if (e.get(i) > 0) {
                    if (maxSupply < e.get(i)) {
                        maxSupply = e.get(i);
                        k = i;
                    }
                }
            }
            if (maxSupply == 0) {
                break;
            }
            delta = maxSupply;

            int[] l = new int[1];
            computeShortestPath(d, prev, k, rCostForward, rCostCapBackward,
                    e, l);

            // find delta (minimum on the path from k to l)
            // delta= e[k];
            // if (-e[l]<delta) delta= e[k];
            int to = l[0];
            do {
                int from = prev.get(to);
                assert from != to;

                // residual
                int itccb = 0;
                while ((itccb < rCostCapBackward.get(from).size())
                        && (rCostCapBackward.get(from).get(itccb).getTo() != to)) {
                    itccb++;
                }
                if (itccb < rCostCapBackward.get(from).size()) {
                    if (rCostCapBackward.get(from).get(itccb).getResidualCapacity() < delta) {
                        delta = rCostCapBackward.get(from).get(itccb).getResidualCapacity();
                    }
                }

                to = from;
            } while (to != k);

            // augment delta flow from k to l (backwards actually...)
            to = l[0];
            do {
                int from = prev.get(to);
                assert from != to;

                // TODO - might do here O(n) can be done in O(1)
                int itx = 0;
                while (x.get(from).get(itx).getTo() != to) {
                    itx++;
                }
                tmp = x.get(from).get(itx).getFlow();
                x.get(from).get(itx).setFlow(tmp + delta);

                // update residual for backward edges
                int itccb = 0;
                while ((itccb < rCostCapBackward.get(to).size())
                        && (rCostCapBackward.get(to).get(itccb).getTo() != from)) {
                    itccb++;
                }
                if (itccb < rCostCapBackward.get(to).size()) {
                    tmp = rCostCapBackward.get(to).get(itccb).getResidualCapacity();
                    rCostCapBackward.get(to).get(itccb).setResidualCapacity(tmp + delta);
                }
                itccb = 0;
                while ((itccb < rCostCapBackward.get(from).size())
                        && (rCostCapBackward.get(from).get(itccb).getTo() != to)) {
                    itccb++;
                }
                if (itccb < rCostCapBackward.get(from).size()) {
                    tmp = rCostCapBackward.get(from).get(itccb).getResidualCapacity();
                    rCostCapBackward.get(from).get(itccb).setResidualCapacity(tmp - delta);
                }

                // update e
                e.set(to, e.get(to) + delta);
                e.set(from, e.get(from) - delta);

                to = from;
            } while (to != k);
        }

        // compute distance from x
        long dist = 0;
        for (int from = 0; from < numNodes; from++) {
            for (EdgeFlow it : x.get(from)) {
                dist += it.getCost() * it.getFlow();
            }
        }
        return dist;
    }
743060
public static KernelFunction autoKernel(Vec dataPoints )
    {
        if(dataPoints.length() < 30)
            return GaussKF.getInstance();
        else if(dataPoints.length() < 1000)
            return EpanechnikovKF.getInstance();
        else//For very large data sets, Uniform is FAST and just as accurate
            return UniformKF.getInstance();
    }
132128
def gpu_with_max_free_mem():
    ""
    mem_all = gpu_mem_get_all()
    if not len(mem_all): return None, 0
    free_all = np.array([x.free for x in mem_all])
    id = np.argmax(free_all)
    return id, free_all[id]
1141569
static void pruneIncorrectShape(FastQueue<Grid> grids , int numRows, int numCols ) {
		// prune clusters which can't be a member calibration target
		for (int i = grids.size()-1; i >= 0; i--) {
			Grid g = grids.get(i);
			if ((g.rows != numRows || g.columns != numCols) && (g.rows != numCols || g.columns != numRows)) {
				grids.remove(i);
			}
		}
	}
107706
def barrier():
    
    if torch.distributed.is_available() and torch.distributed.is_initialized():
        torch.distributed.all_reduce(torch.cuda.FloatTensor(1))
        torch.cuda.synchronize()
1292491
def _reduce(self):
        """"""
        ctx = context.cpu()
        if self._stype == 'default':
            block = self.list_data()
            data = ndarray.add_n(*(w.copyto(ctx) for w in block)) / len(block)
        else:
            # fetch all rows for 'row_sparse' param
            all_row_ids = ndarray.arange(0, self.shape[0], dtype='int64', ctx=ctx)
            data = ndarray.zeros(self.shape, stype='row_sparse', ctx=ctx)
            self._trainer._row_sparse_pull(self, data, all_row_ids, full_idx=True)
        return data
266323
def to_gpu(*args):
    
    if len(args) > 1:
        return (cp.asarray(x) for x in args)
    else:
        return cp.asarray(args[0])