use of org.apache.sysml.runtime.matrix.data.SparseBlockMCSR in project incubator-systemml by apache.
the class MatrixReader method createOutputMatrixBlock.
/**
* NOTE: mallocDense controls if the output matrix blocks is fully allocated, this can be redundant
* if binary block read and single block.
*
* @param rlen number of rows
* @param clen number of columns
* @param bclen number of columns in a block
* @param brlen number of rows in a block
* @param estnnz estimated number of non-zeros
* @param mallocDense if true and not sparse, allocate dense block unsafe
* @param mallocSparse if true and sparse, allocate sparse rows block
* @return matrix block
* @throws IOException if IOException occurs
*/
protected static MatrixBlock createOutputMatrixBlock(long rlen, long clen, int bclen, int brlen, long estnnz, boolean mallocDense, boolean mallocSparse) throws IOException {
// check input dimension
if (!OptimizerUtils.isValidCPDimensions(rlen, clen))
throw new DMLRuntimeException("Matrix dimensions too large for CP runtime: " + rlen + " x " + clen);
// determine target representation (sparse/dense)
boolean sparse = MatrixBlock.evalSparseFormatInMemory(rlen, clen, estnnz);
int numThreads = OptimizerUtils.getParallelBinaryReadParallelism();
long numBlocks = (long) Math.ceil((double) rlen / brlen);
// prepare result matrix block
MatrixBlock ret = new MatrixBlock((int) rlen, (int) clen, sparse, estnnz);
if (!sparse && mallocDense)
ret.allocateDenseBlockUnsafe((int) rlen, (int) clen);
else if (sparse && mallocSparse) {
ret.allocateSparseRowsBlock();
SparseBlock sblock = ret.getSparseBlock();
// create synchronization points for MCSR (start row per block row)
if (// multiple col blocks
sblock instanceof SparseBlockMCSR && clen > bclen && clen >= 0 && bclen > 0 && rlen >= 0 && brlen > 0) {
// adaptive change from scalar to row could cause synchronization issues
if (numThreads <= numBlocks)
for (int i = 0; i < rlen; i += brlen) sblock.allocate(i, Math.max((int) (estnnz / rlen), 2), (int) clen);
else
// allocate all rows to avoid contention
for (int i = 0; i < rlen; i++) sblock.allocate(i, Math.max((int) (estnnz / rlen), 2), (int) clen);
}
}
return ret;
}
use of org.apache.sysml.runtime.matrix.data.SparseBlockMCSR in project incubator-systemml by apache.
the class GPUObject method copyFromHostToDevice.
void copyFromHostToDevice() throws DMLRuntimeException {
LOG.trace("GPU : copyFromHostToDevice, on " + this + ", GPUContext=" + getGPUContext());
long start = 0;
if (DMLScript.STATISTICS)
start = System.nanoTime();
MatrixBlock tmp = mat.acquireRead();
if (tmp.isInSparseFormat()) {
int[] rowPtr = null;
int[] colInd = null;
double[] values = null;
tmp.recomputeNonZeros();
long nnz = tmp.getNonZeros();
mat.getMatrixCharacteristics().setNonZeros(nnz);
SparseBlock block = tmp.getSparseBlock();
boolean copyToDevice = true;
if (block == null && tmp.getNonZeros() == 0) {
// // Allocate empty block --> not necessary
// // To reproduce this, see org.apache.sysml.test.integration.applications.dml.ID3DMLTest
// rowPtr = new int[0];
// colInd = new int[0];
// values = new double[0];
copyToDevice = false;
} else if (block == null && tmp.getNonZeros() != 0) {
throw new DMLRuntimeException("Expected CP sparse block to be not null.");
} else {
// CSR is the preferred format for cuSparse GEMM
// Converts MCSR and COO to CSR
SparseBlockCSR csrBlock = null;
long t0 = 0;
if (block instanceof SparseBlockCSR) {
csrBlock = (SparseBlockCSR) block;
} else if (block instanceof SparseBlockCOO) {
// TODO - should we do this on the GPU using cusparse<t>coo2csr() ?
if (DMLScript.STATISTICS)
t0 = System.nanoTime();
SparseBlockCOO cooBlock = (SparseBlockCOO) block;
csrBlock = new SparseBlockCSR(toIntExact(mat.getNumRows()), cooBlock.rowIndexes(), cooBlock.indexes(), cooBlock.values());
if (DMLScript.STATISTICS)
GPUStatistics.cudaSparseConversionTime.addAndGet(System.nanoTime() - t0);
if (DMLScript.STATISTICS)
GPUStatistics.cudaSparseConversionCount.incrementAndGet();
} else if (block instanceof SparseBlockMCSR) {
if (DMLScript.STATISTICS)
t0 = System.nanoTime();
SparseBlockMCSR mcsrBlock = (SparseBlockMCSR) block;
csrBlock = new SparseBlockCSR(mcsrBlock.getRows(), toIntExact(mcsrBlock.size()));
if (DMLScript.STATISTICS)
GPUStatistics.cudaSparseConversionTime.addAndGet(System.nanoTime() - t0);
if (DMLScript.STATISTICS)
GPUStatistics.cudaSparseConversionCount.incrementAndGet();
} else {
throw new DMLRuntimeException("Unsupported sparse matrix format for CUDA operations");
}
rowPtr = csrBlock.rowPointers();
colInd = csrBlock.indexes();
values = csrBlock.values();
}
allocateSparseMatrixOnDevice();
if (copyToDevice) {
CSRPointer.copyToDevice(getJcudaSparseMatrixPtr(), tmp.getNumRows(), tmp.getNonZeros(), rowPtr, colInd, values);
}
} else {
double[] data = tmp.getDenseBlock();
if (data == null && tmp.getSparseBlock() != null)
throw new DMLRuntimeException("Incorrect sparsity calculation");
else if (data == null && tmp.getNonZeros() != 0)
throw new DMLRuntimeException("MatrixBlock is not allocated");
else if (tmp.getNonZeros() == 0)
data = new double[tmp.getNumRows() * tmp.getNumColumns()];
// Copy dense block
allocateDenseMatrixOnDevice();
cudaMemcpy(getJcudaDenseMatrixPtr(), Pointer.to(data), getDoubleSizeOf(mat.getNumRows() * mat.getNumColumns()), cudaMemcpyHostToDevice);
}
mat.release();
if (DMLScript.STATISTICS)
GPUStatistics.cudaToDevTime.addAndGet(System.nanoTime() - start);
if (DMLScript.STATISTICS)
GPUStatistics.cudaToDevCount.addAndGet(1);
}
Aggregations