You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@systemml.apache.org by na...@apache.org on 2017/06/10 19:07:09 UTC
[1/3] systemml git commit: [FIX] Fixed nested parfor for GPUs
Repository: systemml
Updated Branches:
refs/heads/master 0bcae49ff -> f58717564
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUObject.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUObject.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUObject.java
index 0ed34c5..366eee5 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUObject.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUObject.java
@@ -66,73 +66,91 @@ public class GPUObject {
private static final Log LOG = LogFactory.getLog(GPUObject.class.getName());
- /** GPUContext that owns this GPUObject */
+ /**
+ * GPUContext that owns this GPUObject
+ */
private final GPUContext gpuContext;
- /** Pointer to the underlying dense matrix block on GPU */
+ /**
+ * Pointer to the underlying dense matrix block on GPU
+ */
private Pointer jcudaDenseMatrixPtr = null;
- /** Pointer to the underlying sparse matrix block on GPU */
+ /**
+ * Pointer to the underlying sparse matrix block on GPU
+ */
private CSRPointer jcudaSparseMatrixPtr = null;
- /** An optional tensor descriptor (and shape) that can be set by a tensor instruction such as convolution,
+ /**
+ * An optional tensor descriptor (and shape) that can be set by a tensor instruction such as convolution,
* maxpooling and exploited by a subsequent non-tensor instruction such as relu
*/
private cudnnTensorDescriptor tensorDescriptor = null;
- /** the shape of this tensor, if in fact this is a tensor */
- private int [] tensorShape = null;
+ /**
+ * the shape of this tensor, if in fact this is a tensor
+ */
+ private int[] tensorShape = null;
- /** whether the block attached to this {@link GPUContext} is dirty on the device and needs to be copied back to host */
+ /**
+ * whether the block attached to this {@link GPUContext} is dirty on the device and needs to be copied back to host
+ */
protected boolean dirty = false;
- /** number of read/write locks on this object (this GPUObject is being used in a current instruction) */
+ /**
+ * number of read/write locks on this object (this GPUObject is being used in a current instruction)
+ */
protected AtomicInteger locks = new AtomicInteger(0);
- /** Timestamp, needed by {@link GPUContext#evict(long)} */
+ /**
+ * Timestamp, needed by {@link GPUContext#evict(long)}
+ */
AtomicLong timestamp = new AtomicLong(0);
- /** Whether this block is in sparse format */
+ /**
+ * Whether this block is in sparse format
+ */
protected boolean isSparse = false;
- /** Enclosing {@link MatrixObject} instance */
+ /**
+ * Enclosing {@link MatrixObject} instance
+ */
protected MatrixObject mat = null;
-// private Pointer allocate(String instName, long size) throws DMLRuntimeException {
-// return getGPUContext().allocate(instName, size);
-// }
+ // private Pointer allocate(String instName, long size) throws DMLRuntimeException {
+ // return getGPUContext().allocate(instName, size);
+ // }
@Override
public Object clone() {
GPUObject me = this;
GPUObject that = new GPUObject(me.gpuContext, me.mat);
if (me.tensorShape != null) {
- that.tensorShape = new int[me.tensorShape.length];
- System.arraycopy(me.tensorShape, 0, that.tensorShape, 0, me.tensorShape.length);
- that.allocateTensorDescriptor(me.tensorShape[0], me.tensorShape[1], me.tensorShape[2], me.tensorShape[3]);
- }
+ that.tensorShape = new int[me.tensorShape.length];
+ System.arraycopy(me.tensorShape, 0, that.tensorShape, 0, me.tensorShape.length);
+ that.allocateTensorDescriptor(me.tensorShape[0], me.tensorShape[1], me.tensorShape[2], me.tensorShape[3]);
+ }
that.dirty = me.dirty;
that.locks = new AtomicInteger(me.locks.get());
that.timestamp = new AtomicLong(me.timestamp.get());
that.isSparse = me.isSparse;
try {
- if (me.jcudaDenseMatrixPtr != null) {
- long rows = me.mat.getNumRows();
- long cols = me.mat.getNumColumns();
- long size = rows * cols * Sizeof.DOUBLE;
- me.gpuContext.ensureFreeSpace((int)size);
- that.jcudaDenseMatrixPtr = allocate(size);
- cudaMemcpy(that.jcudaDenseMatrixPtr, me.jcudaDenseMatrixPtr, size, cudaMemcpyDeviceToDevice);
- }
-
- if (me.getJcudaSparseMatrixPtr() != null){
- long rows = mat.getNumRows();
- that.jcudaSparseMatrixPtr = me.jcudaSparseMatrixPtr.clone((int)rows);
- }
+ if (me.jcudaDenseMatrixPtr != null) {
+ long rows = me.mat.getNumRows();
+ long cols = me.mat.getNumColumns();
+ long size = rows * cols * Sizeof.DOUBLE;
+ me.gpuContext.ensureFreeSpace((int) size);
+ that.jcudaDenseMatrixPtr = allocate(size);
+ cudaMemcpy(that.jcudaDenseMatrixPtr, me.jcudaDenseMatrixPtr, size, cudaMemcpyDeviceToDevice);
+ }
+ if (me.getJcudaSparseMatrixPtr() != null) {
+ long rows = mat.getNumRows();
+ that.jcudaSparseMatrixPtr = me.jcudaSparseMatrixPtr.clone((int) rows);
+ }
- } catch (DMLRuntimeException e){
+ } catch (DMLRuntimeException e) {
throw new RuntimeException(e);
}
@@ -147,9 +165,9 @@ public class GPUObject {
getGPUContext().cudaFreeHelper(toFree);
}
-// private void cudaFreeHelper(Pointer toFree, boolean eager) throws DMLRuntimeException {
-// getGPUContext().cudaFreeHelper(toFree, eager);
-// }
+ // private void cudaFreeHelper(Pointer toFree, boolean eager) throws DMLRuntimeException {
+ // getGPUContext().cudaFreeHelper(toFree, eager);
+ // }
private void cudaFreeHelper(String instName, Pointer toFree, boolean eager) throws DMLRuntimeException {
getGPUContext().cudaFreeHelper(instName, toFree, eager);
@@ -161,24 +179,27 @@ public class GPUObject {
/**
* Transposes a dense matrix on the GPU by calling the cublasDgeam operation
- * @param gCtx a valid {@link GPUContext}
- * @param densePtr Pointer to dense matrix on the GPU
- * @param m rows in ouput matrix
- * @param n columns in output matrix
- * @param lda rows in input matrix
- * @param ldc columns in output matrix
- * @return transposed matrix
+ *
+ * @param gCtx a valid {@link GPUContext}
+ * @param densePtr Pointer to dense matrix on the GPU
+ * @param m rows in ouput matrix
+ * @param n columns in output matrix
+ * @param lda rows in input matrix
+ * @param ldc columns in output matrix
* @throws DMLRuntimeException if operation failed
+ * @return transposed matrix
*/
- public static Pointer transpose(GPUContext gCtx, Pointer densePtr, int m, int n, int lda, int ldc) throws DMLRuntimeException {
+ public static Pointer transpose(GPUContext gCtx, Pointer densePtr, int m, int n, int lda, int ldc)
+ throws DMLRuntimeException {
LOG.trace("GPU : transpose of block of size [" + m + "," + n + "]" + ", GPUContext=" + gCtx);
- Pointer alpha = Pointer.to(new double[]{1.0});
- Pointer beta = Pointer.to(new double[]{0.0});
+ Pointer alpha = Pointer.to(new double[] { 1.0 });
+ Pointer beta = Pointer.to(new double[] { 0.0 });
Pointer A = densePtr;
- Pointer C = gCtx.allocate(((long)m)*getDoubleSizeOf(n));
+ Pointer C = gCtx.allocate(((long) m) * getDoubleSizeOf(n));
// Transpose the matrix to get a dense matrix
- JCublas2.cublasDgeam(gCtx.getCublasHandle(), CUBLAS_OP_T, CUBLAS_OP_T, m, n, alpha, A, lda, beta, new Pointer(), lda, C, ldc);
+ JCublas2.cublasDgeam(gCtx.getCublasHandle(), CUBLAS_OP_T, CUBLAS_OP_T, m, n, alpha, A, lda, beta, new Pointer(),
+ lda, C, ldc);
return C;
}
@@ -187,15 +208,17 @@ public class GPUObject {
* Since the allocated matrix is temporary, bookkeeping is not updated.
* Also note that the input dense matrix is expected to be in COLUMN MAJOR FORMAT
* Caller is responsible for deallocating memory on GPU.
- * @param gCtx a valid {@link GPUContext}
+ *
+ * @param gCtx a valid {@link GPUContext}
* @param cusparseHandle handle to cusparse library
- * @param densePtr [in] dense matrix pointer on the GPU in row major
- * @param rows number of rows
- * @param cols number of columns
+ * @param densePtr [in] dense matrix pointer on the GPU in row major
+ * @param rows number of rows
+ * @param cols number of columns
* @return CSR (compressed sparse row) pointer
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
- public static CSRPointer columnMajorDenseToRowMajorSparse(GPUContext gCtx, cusparseHandle cusparseHandle, Pointer densePtr, int rows, int cols) throws DMLRuntimeException {
+ public static CSRPointer columnMajorDenseToRowMajorSparse(GPUContext gCtx, cusparseHandle cusparseHandle,
+ Pointer densePtr, int rows, int cols) throws DMLRuntimeException {
cusparseMatDescr matDescr = CSRPointer.getDefaultCuSparseMatrixDescriptor();
Pointer nnzPerRowPtr = null;
Pointer nnzTotalDevHostPtr = null;
@@ -205,24 +228,31 @@ public class GPUObject {
nnzTotalDevHostPtr = gCtx.allocate(getIntSizeOf(1));
// Output is in dense vector format, convert it to CSR
- cusparseDnnz(cusparseHandle, cusparseDirection.CUSPARSE_DIRECTION_ROW, rows, cols, matDescr, densePtr, rows, nnzPerRowPtr, nnzTotalDevHostPtr);
+ cusparseDnnz(cusparseHandle, cusparseDirection.CUSPARSE_DIRECTION_ROW, rows, cols, matDescr, densePtr, rows,
+ nnzPerRowPtr, nnzTotalDevHostPtr);
//cudaDeviceSynchronize();
- int[] nnzC = {-1};
+ int[] nnzC = { -1 };
- long t2=0;
- if (DMLScript.STATISTICS) t2 = System.nanoTime();
+ long t2 = 0;
+ if (DMLScript.STATISTICS)
+ t2 = System.nanoTime();
cudaMemcpy(Pointer.to(nnzC), nnzTotalDevHostPtr, getIntSizeOf(1), cudaMemcpyDeviceToHost);
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - t2);
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevCount.addAndGet(1);
-
- if (nnzC[0] == -1){
- throw new DMLRuntimeException("cusparseDnnz did not calculate the correct number of nnz from the sparse-matrix vector mulitply on the GPU");
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - t2);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevCount.addAndGet(1);
+
+ if (nnzC[0] == -1) {
+ throw new DMLRuntimeException(
+ "cusparseDnnz did not calculate the correct number of nnz from the sparse-matrix vector mulitply on the GPU");
}
- LOG.trace("GPU : col-major dense size[" + rows + "," + cols + "] to row-major sparse of with nnz = " + nnzC[0] + ", GPUContext=" + gCtx);
+ LOG.trace("GPU : col-major dense size[" + rows + "," + cols + "] to row-major sparse of with nnz = " + nnzC[0]
+ + ", GPUContext=" + gCtx);
CSRPointer C = CSRPointer.allocateEmpty(gCtx, nnzC[0], rows);
- cusparseDdense2csr(cusparseHandle, rows, cols, matDescr, densePtr, rows, nnzPerRowPtr, C.val, C.rowPtr, C.colInd);
+ cusparseDdense2csr(cusparseHandle, rows, cols, matDescr, densePtr, rows, nnzPerRowPtr, C.val, C.rowPtr,
+ C.colInd);
//cudaDeviceSynchronize();
gCtx.cudaFreeHelper(nnzPerRowPtr);
@@ -233,11 +263,12 @@ public class GPUObject {
/**
* Gets the double array from GPU memory onto host memory and returns string.
- * @param A Pointer to memory on device (GPU), assumed to point to a double array
+ *
+ * @param A Pointer to memory on device (GPU), assumed to point to a double array
* @param rows rows in matrix A
* @param cols columns in matrix A
* @return the debug string
- * @throws DMLRuntimeException if DMLRuntimeException occurs
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
*/
public static String debugString(Pointer A, long rows, long cols) throws DMLRuntimeException {
StringBuffer sb = new StringBuffer();
@@ -245,8 +276,8 @@ public class GPUObject {
double[] tmp = new double[len];
cudaMemcpy(Pointer.to(tmp), A, getDoubleSizeOf(len), cudaMemcpyDeviceToHost);
int k = 0;
- for (int i=0; i<rows; i++){
- for (int j=0; j<cols; j++){
+ for (int i = 0; i < rows; i++) {
+ for (int j = 0; j < cols; j++) {
sb.append(tmp[k]).append(' ');
k++;
}
@@ -257,6 +288,7 @@ public class GPUObject {
/**
* Convenience method to directly examine the Sparse matrix on GPU
+ *
* @return CSR (compressed sparse row) pointer
*/
public CSRPointer getSparseMatrixCudaPointer() {
@@ -266,14 +298,14 @@ public class GPUObject {
/**
* Convenience method to directly set the sparse matrix on GPU
* Needed for operations like {@link JCusparse#cusparseDcsrgemm(cusparseHandle, int, int, int, int, int, cusparseMatDescr, int, Pointer, Pointer, Pointer, cusparseMatDescr, int, Pointer, Pointer, Pointer, cusparseMatDescr, Pointer, Pointer, Pointer)}
+ *
* @param sparseMatrixPtr CSR (compressed sparse row) pointer
- *
* @throws DMLRuntimeException ?
*/
public void setSparseMatrixCudaPointer(CSRPointer sparseMatrixPtr) throws DMLRuntimeException {
this.jcudaSparseMatrixPtr = sparseMatrixPtr;
this.isSparse = true;
- if(getJcudaDenseMatrixPtr() != null) {
+ if (getJcudaDenseMatrixPtr() != null) {
cudaFreeHelper(getJcudaDenseMatrixPtr());
jcudaDenseMatrixPtr = null;
}
@@ -286,10 +318,10 @@ public class GPUObject {
* @param densePtr dense pointer
* @throws DMLRuntimeException ?
*/
- public void setDenseMatrixCudaPointer(Pointer densePtr) throws DMLRuntimeException{
+ public void setDenseMatrixCudaPointer(Pointer densePtr) throws DMLRuntimeException {
this.jcudaDenseMatrixPtr = densePtr;
this.isSparse = false;
- if(getJcudaSparseMatrixPtr() != null) {
+ if (getJcudaSparseMatrixPtr() != null) {
getJcudaSparseMatrixPtr().deallocate();
jcudaSparseMatrixPtr = null;
}
@@ -303,26 +335,32 @@ public class GPUObject {
*/
public void denseToSparse() throws DMLRuntimeException {
LOG.trace("GPU : dense -> sparse on " + this + ", GPUContext=" + getGPUContext());
- long t0=0;
- if (DMLScript.STATISTICS) t0 = System.nanoTime();
+ long t0 = 0;
+ if (DMLScript.STATISTICS)
+ t0 = System.nanoTime();
cusparseHandle cusparseHandle = getGPUContext().getCusparseHandle();
- if(cusparseHandle == null)
+ if (cusparseHandle == null)
throw new DMLRuntimeException("Expected cusparse to be initialized");
int rows = toIntExact(mat.getNumRows());
int cols = toIntExact(mat.getNumColumns());
- if(getJcudaDenseMatrixPtr() == null || !isAllocated())
+ if (getJcudaDenseMatrixPtr() == null || !isAllocated())
throw new DMLRuntimeException("Expected allocated dense matrix before denseToSparse() call");
denseRowMajorToColumnMajor();
- setSparseMatrixCudaPointer(columnMajorDenseToRowMajorSparse(getGPUContext(), cusparseHandle, getJcudaDenseMatrixPtr(), rows, cols));
+ setSparseMatrixCudaPointer(
+ columnMajorDenseToRowMajorSparse(getGPUContext(), cusparseHandle, getJcudaDenseMatrixPtr(), rows,
+ cols));
// TODO: What if mat.getNnz() is -1 ?
- if (DMLScript.STATISTICS) GPUStatistics.cudaDenseToSparseTime.addAndGet(System.nanoTime() - t0);
- if (DMLScript.STATISTICS) GPUStatistics.cudaDenseToSparseCount.addAndGet(1);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaDenseToSparseTime.addAndGet(System.nanoTime() - t0);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaDenseToSparseCount.addAndGet(1);
}
/**
* Convenience method. Converts Row Major Dense Matrix to Column Major Dense Matrix
+ *
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
public void denseRowMajorToColumnMajor() throws DMLRuntimeException {
@@ -331,7 +369,7 @@ public class GPUObject {
int n = toIntExact(mat.getNumColumns());
int lda = n;
int ldc = m;
- if(!isAllocated()) {
+ if (!isAllocated()) {
throw new DMLRuntimeException("Error in converting row major to column major : data is not allocated");
}
@@ -342,6 +380,7 @@ public class GPUObject {
/**
* Convenience method. Converts Column Major Dense Matrix to Row Major Dense Matrix
+ *
* @throws DMLRuntimeException if error
*/
public void denseColumnMajorToRowMajor() throws DMLRuntimeException {
@@ -351,7 +390,7 @@ public class GPUObject {
int m = toIntExact(mat.getNumColumns());
int lda = n;
int ldc = m;
- if(!isAllocated()) {
+ if (!isAllocated()) {
throw new DMLRuntimeException("Error in converting column major to row major : data is not allocated");
}
@@ -372,22 +411,28 @@ public class GPUObject {
/**
* Convert sparse to dense (Performs transpose, use sparseToColumnMajorDense if the kernel can deal with column major format)
* Also records per instruction invokation of sparseToDense.
- * @param instructionName Name of the instruction for which statistics are recorded in {@link GPUStatistics}
+ *
+ * @param instructionName Name of the instruction for which statistics are recorded in {@link GPUStatistics}
* @throws DMLRuntimeException ?
*/
public void sparseToDense(String instructionName) throws DMLRuntimeException {
LOG.trace("GPU : sparse -> dense on " + this + ", GPUContext=" + getGPUContext());
- long start=0, end=0;
- if (DMLScript.STATISTICS) start = System.nanoTime();
- if(getJcudaSparseMatrixPtr() == null || !isAllocated())
+ long start = 0, end = 0;
+ if (DMLScript.STATISTICS)
+ start = System.nanoTime();
+ if (getJcudaSparseMatrixPtr() == null || !isAllocated())
throw new DMLRuntimeException("Expected allocated sparse matrix before sparseToDense() call");
sparseToColumnMajorDense();
denseColumnMajorToRowMajor();
- if (DMLScript.STATISTICS) end = System.nanoTime();
- if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS) GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_SPARSE_TO_DENSE, end - start);
- if (DMLScript.STATISTICS) GPUStatistics.cudaSparseToDenseTime.addAndGet(end - start);
- if (DMLScript.STATISTICS) GPUStatistics.cudaSparseToDenseCount.addAndGet(1);
+ if (DMLScript.STATISTICS)
+ end = System.nanoTime();
+ if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
+ GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_SPARSE_TO_DENSE, end - start);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaSparseToDenseTime.addAndGet(end - start);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaSparseToDenseCount.addAndGet(1);
}
/**
@@ -397,11 +442,11 @@ public class GPUObject {
*/
public void sparseToColumnMajorDense() throws DMLRuntimeException {
LOG.trace("GPU : sparse -> col-major dense on " + this + ", GPUContext=" + getGPUContext());
- if(getJcudaSparseMatrixPtr() == null || !isAllocated())
+ if (getJcudaSparseMatrixPtr() == null || !isAllocated())
throw new DMLRuntimeException("Expected allocated sparse matrix before sparseToDense() call");
cusparseHandle cusparseHandle = getGPUContext().getCusparseHandle();
- if(cusparseHandle == null)
+ if (cusparseHandle == null)
throw new DMLRuntimeException("Expected cusparse to be initialized");
int rows = toIntExact(mat.getNumRows());
int cols = toIntExact(mat.getNumColumns());
@@ -410,9 +455,10 @@ public class GPUObject {
/**
* Initializes this GPUObject with a {@link MatrixObject} instance which will contain metadata about the enclosing matrix block
+ *
* @param mat2 the matrix block that owns this {@link GPUObject}
*/
- GPUObject(GPUContext gCtx, MatrixObject mat2) {
+ GPUObject(GPUContext gCtx, MatrixObject mat2) {
gpuContext = gCtx;
this.mat = mat2;
}
@@ -423,14 +469,16 @@ public class GPUObject {
/**
* Returns a previously allocated tensor shape or null
+ *
* @return int array of four elements or null
*/
- public int [] getTensorShape() {
+ public int[] getTensorShape() {
return tensorShape;
}
/**
* Returns a previously allocated tensor descriptor or null
+ *
* @return cudnn tensor descriptor
*/
public cudnnTensorDescriptor getTensorDescriptor() {
@@ -439,6 +487,7 @@ public class GPUObject {
/**
* Returns a previously allocated or allocates and returns a tensor descriptor
+ *
* @param N number of images
* @param C number of channels
* @param H height
@@ -446,8 +495,8 @@ public class GPUObject {
* @return cudnn tensor descriptor
*/
public cudnnTensorDescriptor allocateTensorDescriptor(int N, int C, int H, int W) {
- LOG.trace("GPU : allocateTensorDescriptor with [N="+N+",C="+C+",H="+H+",W="+W+"] on " + this);
- if(tensorDescriptor == null) {
+ LOG.trace("GPU : allocateTensorDescriptor with [N=" + N + ",C=" + C + ",H=" + H + ",W=" + W + "] on " + this);
+ if (tensorDescriptor == null) {
tensorDescriptor = new cudnnTensorDescriptor();
cudnnCreateTensorDescriptor(tensorDescriptor);
cudnnSetTensor4dDescriptor(tensorDescriptor, CUDNN_TENSOR_NCHW, CUDNN_DATA_DOUBLE, N, C, H, W);
@@ -461,11 +510,11 @@ public class GPUObject {
}
private static long getDoubleSizeOf(long numElems) {
- return numElems * ((long)jcuda.Sizeof.DOUBLE);
+ return numElems * ((long) jcuda.Sizeof.DOUBLE);
}
private static long getIntSizeOf(long numElems) {
- return numElems * ((long)jcuda.Sizeof.INT);
+ return numElems * ((long) jcuda.Sizeof.INT);
}
public boolean isAllocated() {
@@ -488,7 +537,7 @@ public class GPUObject {
*
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
- public void allocateSparseAndEmpty() throws DMLRuntimeException{
+ public void allocateSparseAndEmpty() throws DMLRuntimeException {
LOG.trace("GPU : allocate sparse and empty block on " + this + ", GPUContext=" + getGPUContext());
setSparseMatrixCudaPointer(CSRPointer.allocateEmpty(getGPUContext(), 0, mat.getNumRows()));
}
@@ -510,7 +559,9 @@ public class GPUObject {
// The "fill" kernel is called which treats the matrix "jcudaDensePtr" like a vector and fills it with value "v"
// If the fill value is 0, no need to call the special kernel, the allocate memsets the allocated region to 0
if (v != 0)
- getGPUContext().getKernels().launchKernel("fill", ExecutionConfig.getConfigForSimpleVectorOperations(numElems), getJcudaDenseMatrixPtr(), v, numElems);
+ getGPUContext().getKernels()
+ .launchKernel("fill", ExecutionConfig.getConfigForSimpleVectorOperations(numElems),
+ getJcudaDenseMatrixPtr(), v, numElems);
}
/**
@@ -520,8 +571,8 @@ public class GPUObject {
* @return true if sparse and empty
* @throws DMLRuntimeException if error
*/
- public boolean isSparseAndEmpty() throws DMLRuntimeException{
- boolean isSparseAndAllocated = isAllocated()&& LibMatrixCUDA.isInSparseFormat(getGPUContext(), mat);
+ public boolean isSparseAndEmpty() throws DMLRuntimeException {
+ boolean isSparseAndAllocated = isAllocated() && LibMatrixCUDA.isInSparseFormat(getGPUContext(), mat);
boolean isEmptyAndSparseAndAllocated = isSparseAndAllocated && getJcudaSparseMatrixPtr().nnz == 0;
return isEmptyAndSparseAndAllocated;
}
@@ -529,13 +580,15 @@ public class GPUObject {
public boolean acquireDeviceRead() throws DMLRuntimeException {
LOG.trace("GPU : acquireDeviceRead on " + this);
boolean transferred = false;
- if(!isAllocated()) {
- LOG.trace("GPU : in acquireDeviceRead, data is not allocated, copying from host, on " + this + ", GPUContext=" + getGPUContext());
+ if (!isAllocated()) {
+ LOG.trace(
+ "GPU : in acquireDeviceRead, data is not allocated, copying from host, on " + this + ", GPUContext="
+ + getGPUContext());
copyFromHostToDevice();
transferred = true;
}
addLock();
- if(!isAllocated())
+ if (!isAllocated())
throw new DMLRuntimeException("Expected device data to be allocated");
return transferred;
}
@@ -543,7 +596,7 @@ public class GPUObject {
public boolean acquireDeviceModifyDense() throws DMLRuntimeException {
LOG.trace("GPU : acquireDeviceModifyDense on " + this + ", GPUContext=" + getGPUContext());
boolean allocated = false;
- if(!isAllocated()) {
+ if (!isAllocated()) {
mat.setDirty(true);
LOG.trace("GPU : data is not allocated, allocating a dense block, on " + this);
// Dense block, size = numRows * numCols
@@ -551,7 +604,7 @@ public class GPUObject {
allocated = true;
}
dirty = true;
- if(!isAllocated())
+ if (!isAllocated())
throw new DMLRuntimeException("Expected device data to be allocated");
return allocated;
}
@@ -560,14 +613,14 @@ public class GPUObject {
LOG.trace("GPU : acquireDeviceModifySparse on " + this + ", GPUContext=" + getGPUContext());
boolean allocated = false;
isSparse = true;
- if(!isAllocated()) {
+ if (!isAllocated()) {
LOG.trace("GPU : data is not allocated, allocating a sparse block, on " + this);
mat.setDirty(true);
allocateSparseMatrixOnDevice();
allocated = true;
}
dirty = true;
- if(!isAllocated())
+ if (!isAllocated())
throw new DMLRuntimeException("Expected device data to be allocated");
return allocated;
}
@@ -578,6 +631,7 @@ public class GPUObject {
/**
* if the data is allocated on the GPU and is dirty, it is copied back to the host memory
+ *
* @return true if a copy to host happened, false otherwise
* @throws CacheException ?
*/
@@ -586,7 +640,8 @@ public class GPUObject {
try {
LOG.trace("GPU : acquireDeviceModifySparse on " + this + ", GPUContext=" + getGPUContext());
if (isAllocated() && dirty) {
- LOG.trace("GPU : data is dirty on device, copying to host, on " + this + ", GPUContext=" + getGPUContext());
+ LOG.trace("GPU : data is dirty on device, copying to host, on " + this + ", GPUContext="
+ + getGPUContext());
copyFromDeviceToHost();
copied = true;
}
@@ -598,6 +653,7 @@ public class GPUObject {
/**
* Updates the locks depending on the eviction policy selected
+ *
* @throws DMLRuntimeException if there is no locked GPU Object or if could not obtain a {@link GPUContext}
*/
private void updateReleaseLocks(int l) throws DMLRuntimeException {
@@ -606,36 +662,45 @@ public class GPUObject {
throw new CacheException("Internal state error : Invalid number of locks on a GPUObject");
}
- LOG.trace("GPU : updateReleaseLocks, new number of locks is " + locks.get() + ", on " + this + ", GPUContext=" + getGPUContext());
+ LOG.trace("GPU : updateReleaseLocks, new number of locks is " + locks.get() + ", on " + this + ", GPUContext="
+ + getGPUContext());
GPUContext.EvictionPolicy evictionPolicy = getGPUContext().evictionPolicy;
- switch (evictionPolicy){
- case LRU : timestamp.set(System.nanoTime()); break;
- case LFU : timestamp.addAndGet(1); break;
- case MIN_EVICT : /* Do Nothing */ break;
- default : throw new CacheException("The eviction policy is not supported:" + evictionPolicy.name());
+ switch (evictionPolicy) {
+ case LRU:
+ timestamp.set(System.nanoTime());
+ break;
+ case LFU:
+ timestamp.addAndGet(1);
+ break;
+ case MIN_EVICT: /* Do Nothing */
+ break;
+ default:
+ throw new CacheException("The eviction policy is not supported:" + evictionPolicy.name());
}
}
/**
* Releases input allocated on GPU
+ *
* @throws DMLRuntimeException if data is not allocated or if there is no locked GPU Object or if could not obtain a {@link GPUContext}
*/
public void releaseInput() throws DMLRuntimeException {
// A read lock is a positive quantity, therefor when the lock is freed, a negative 1 is added
updateReleaseLocks(-1);
- if(!isAllocated())
+ if (!isAllocated())
throw new CacheException("Attempting to release an input before allocating it");
}
/**
* releases output allocated on GPU
+ *
* @throws DMLRuntimeException if data is not allocated or if there is no locked GPU Object or if could not obtain a {@link GPUContext}
*/
public void releaseOutput() throws DMLRuntimeException {
// A write lock is a negative quantity, therefore when the lock is freed, a positive number is added
updateReleaseLocks(1);
dirty = true;
- if(!isAllocated())
+ if (!isAllocated())
throw new CacheException("Attempting to release an output before allocating it");
}
@@ -664,7 +729,7 @@ public class GPUObject {
void deallocateMemoryOnDevice(boolean eager) throws DMLRuntimeException {
LOG.trace("GPU : deallocateMemoryOnDevice, on " + this + ", GPUContext=" + getGPUContext());
- if(getJcudaDenseMatrixPtr() != null) {
+ if (getJcudaDenseMatrixPtr() != null) {
cudaFreeHelper(null, getJcudaDenseMatrixPtr(), eager);
}
if (getJcudaSparseMatrixPtr() != null) {
@@ -672,7 +737,7 @@ public class GPUObject {
}
jcudaDenseMatrixPtr = null;
jcudaSparseMatrixPtr = null;
- if(tensorDescriptor != null) {
+ if (tensorDescriptor != null) {
cudnnDestroyTensorDescriptor(tensorDescriptor);
tensorDescriptor = null;
}
@@ -686,10 +751,9 @@ public class GPUObject {
long clen = mat.getNumColumns();
long nnz = mat.getNnz();
- if(LibMatrixCUDA.isInSparseFormat(getGPUContext(), mat)) {
+ if (LibMatrixCUDA.isInSparseFormat(getGPUContext(), mat)) {
GPUSize = CSRPointer.estimateSize(nnz, rlen);
- }
- else {
+ } else {
GPUSize = getDoubleSizeOf(rlen * clen);
}
return GPUSize;
@@ -697,11 +761,12 @@ public class GPUObject {
void copyFromHostToDevice() throws DMLRuntimeException {
LOG.trace("GPU : copyFromHostToDevice, on " + this + ", GPUContext=" + getGPUContext());
- long start=0;
- if (DMLScript.STATISTICS) start = System.nanoTime();
+ long start = 0;
+ if (DMLScript.STATISTICS)
+ start = System.nanoTime();
MatrixBlock tmp = mat.acquireRead();
- if(tmp.isInSparseFormat()) {
+ if (tmp.isInSparseFormat()) {
int rowPtr[] = null;
int colInd[] = null;
@@ -713,37 +778,42 @@ public class GPUObject {
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];
+ 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) {
+ } else if (block == null && tmp.getNonZeros() != 0) {
throw new DMLRuntimeException("Expected CP sparse block to be not null.");
- }
- else {
+ } 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;
+ 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();
+ 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;
+ 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();
+ 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");
}
@@ -754,30 +824,33 @@ public class GPUObject {
allocateSparseMatrixOnDevice();
- if(copyToDevice) {
- CSRPointer.copyToDevice(getJcudaSparseMatrixPtr(), tmp.getNumRows(), tmp.getNonZeros(), rowPtr, colInd, values);
+ if (copyToDevice) {
+ CSRPointer.copyToDevice(getJcudaSparseMatrixPtr(), tmp.getNumRows(), tmp.getNonZeros(), rowPtr, colInd,
+ values);
}
- }
- else {
+ } else {
double[] data = tmp.getDenseBlock();
- if( data == null && tmp.getSparseBlock() != null )
+ if (data == null && tmp.getSparseBlock() != null)
throw new DMLRuntimeException("Incorrect sparsity calculation");
- else if( data==null && tmp.getNonZeros() != 0 )
+ 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()];
+ 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);
+ 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);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaToDevTime.addAndGet(System.nanoTime() - start);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaToDevCount.addAndGet(1);
}
public static int toIntExact(long l) throws DMLRuntimeException {
@@ -789,37 +862,42 @@ public class GPUObject {
protected void copyFromDeviceToHost() throws DMLRuntimeException {
LOG.trace("GPU : copyFromDeviceToHost, on " + this + ", GPUContext=" + getGPUContext());
- if (getJcudaDenseMatrixPtr() != null && getJcudaSparseMatrixPtr() != null){
+ if (getJcudaDenseMatrixPtr() != null && getJcudaSparseMatrixPtr() != null) {
throw new DMLRuntimeException("Invalid state : JCuda dense/sparse pointer are both allocated");
}
- if(getJcudaDenseMatrixPtr() != null) {
- long start=0;
- if (DMLScript.STATISTICS) start = System.nanoTime();
+ if (getJcudaDenseMatrixPtr() != null) {
+ long start = 0;
+ if (DMLScript.STATISTICS)
+ start = System.nanoTime();
MatrixBlock tmp = new MatrixBlock(toIntExact(mat.getNumRows()), toIntExact(mat.getNumColumns()), false);
tmp.allocateDenseBlock();
- double [] data = tmp.getDenseBlock();
+ double[] data = tmp.getDenseBlock();
- cudaMemcpy(Pointer.to(data), getJcudaDenseMatrixPtr(), getDoubleSizeOf(data.length), cudaMemcpyDeviceToHost);
+ cudaMemcpy(Pointer.to(data), getJcudaDenseMatrixPtr(), getDoubleSizeOf(data.length),
+ cudaMemcpyDeviceToHost);
tmp.recomputeNonZeros();
mat.acquireModify(tmp);
mat.release();
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime()-start);
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevCount.addAndGet(1);
- }
- else if (getJcudaSparseMatrixPtr() != null){
- if(!LibMatrixCUDA.isInSparseFormat(getGPUContext(), mat))
- throw new DMLRuntimeException("Block not in sparse format on host yet the device sparse matrix pointer is not null");
-
- if(this.isSparseAndEmpty()){
- MatrixBlock tmp = new MatrixBlock(); // Empty Block
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - start);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevCount.addAndGet(1);
+ } else if (getJcudaSparseMatrixPtr() != null) {
+ if (!LibMatrixCUDA.isInSparseFormat(getGPUContext(), mat))
+ throw new DMLRuntimeException(
+ "Block not in sparse format on host yet the device sparse matrix pointer is not null");
+
+ if (this.isSparseAndEmpty()) {
+ MatrixBlock tmp = new MatrixBlock(); // Empty Block
mat.acquireModify(tmp);
mat.release();
} else {
- long start=0;
- if (DMLScript.STATISTICS) start = System.nanoTime();
+ long start = 0;
+ if (DMLScript.STATISTICS)
+ start = System.nanoTime();
int rows = toIntExact(mat.getNumRows());
int cols = toIntExact(mat.getNumColumns());
@@ -833,19 +911,21 @@ public class GPUObject {
MatrixBlock tmp = new MatrixBlock(rows, cols, nnz, sparseBlock);
mat.acquireModify(tmp);
mat.release();
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - start);
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevCount.addAndGet(1);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - start);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevCount.addAndGet(1);
}
- }
- else {
- throw new DMLRuntimeException("Cannot copy from device to host as JCuda dense/sparse pointer is not allocated");
+ } else {
+ throw new DMLRuntimeException(
+ "Cannot copy from device to host as JCuda dense/sparse pointer is not allocated");
}
dirty = false;
}
-
/**
* lazily clears the data associated with this {@link GPUObject} instance
+ *
* @throws CacheException ?
*/
public void clearData() throws DMLRuntimeException {
@@ -854,6 +934,7 @@ public class GPUObject {
/**
* Clears the data associated with this {@link GPUObject} instance
+ *
* @param eager whether to be done synchronously or asynchronously
* @throws CacheException ?
*/
@@ -862,24 +943,27 @@ public class GPUObject {
getGPUContext().removeRecordedUsage(this);
}
- /**
- * Pointer to dense matrix
+ /**
+ * Pointer to dense matrix
+ *
* @return ?
*/
public Pointer getJcudaDenseMatrixPtr() {
return jcudaDenseMatrixPtr;
}
- /**
+ /**
* Pointer to sparse matrix
+ *
* @return ?
*/
public CSRPointer getJcudaSparseMatrixPtr() {
return jcudaSparseMatrixPtr;
}
- /**
- * Whether this block is dirty on the GPU
+ /**
+ * Whether this block is dirty on the GPU
+ *
* @return ?
*/
public boolean isDirty() {
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/JCudaKernels.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/JCudaKernels.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/JCudaKernels.java
index 4d06831..246aecc 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/JCudaKernels.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/JCudaKernels.java
@@ -38,9 +38,9 @@ import jcuda.runtime.JCuda;
/**
* Utility class that allows LibMatrixCUDA as well as GPUObject to invoke custom CUDA kernels.
- *
- * The utility org.apache.sysml.runtime.instructions.gpu.context.JCudaKernels simplifies the launching of the kernels.
- * For example: to launch a kernel
+ * <p>
+ * The utility org.apache.sysml.runtime.instructions.gpu.context.JCudaKernels simplifies the launching of the kernels.
+ * For example: to launch a kernel
* {@code copyUpperToLowerTriangleDense<<1,1,32,32>>(jcudaDenseMatrixPtr, dim, dim*dim) }, the user has to call:
* {@code kernels.launchKernel("copyUpperToLowerTriangleDense", new ExecutionConfig(1,1,32,32), jcudaDenseMatrixPtr, dim, dim*dim) }
*/
@@ -49,81 +49,75 @@ public class JCudaKernels {
private final static String ptxFileName = "/kernels/SystemML.ptx";
private HashMap<String, CUfunction> kernels = new HashMap<String, CUfunction>();
private CUmodule module;
-// private final int deviceNum;
-
+ // private final int deviceNum;
+
/**
* Loads the kernels in the file ptxFileName. Though cubin files are also supported, we will stick with
* ptx file as they are target-independent similar to Java's .class files.
- * @param deviceNum the device number for which to initiate the driver API
+ *
+ * @param deviceNum the device number for which to initiate the driver API
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
JCudaKernels(int deviceNum) throws DMLRuntimeException {
-// this.deviceNum = deviceNum;
+ // this.deviceNum = deviceNum;
module = new CUmodule();
// Load the kernels specified in the ptxFileName file
checkResult(cuModuleLoadDataEx(module, initKernels(ptxFileName), 0, new int[0], Pointer.to(new int[0])));
}
/**
- * Setups the kernel parameters and launches the kernel using cuLaunchKernel API.
+ * Setups the kernel parameters and launches the kernel using cuLaunchKernel API.
* This function currently supports two dimensional grid and blocks.
- *
- * @param name name of the kernel
- * @param config execution configuration
+ *
+ * @param name name of the kernel
+ * @param config execution configuration
* @param arguments can be of type Pointer, long, double, float and int
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
- public void launchKernel(String name, ExecutionConfig config, Object ... arguments) throws DMLRuntimeException {
+ public void launchKernel(String name, ExecutionConfig config, Object... arguments) throws DMLRuntimeException {
CUfunction function = kernels.get(name);
- if(function == null) {
+ if (function == null) {
// caching functions into hashmap reduces the lookup overhead
function = new CUfunction();
checkResult(cuModuleGetFunction(function, module, name));
}
-
+
// Setup parameters
- Pointer [] kernelParams = new Pointer[arguments.length];
+ Pointer[] kernelParams = new Pointer[arguments.length];
for (int i = 0; i < arguments.length; i++) {
- if(arguments[i] == null) {
+ if (arguments[i] == null) {
throw new DMLRuntimeException("The argument to the kernel cannot be null.");
+ } else if (arguments[i] instanceof Pointer) {
+ kernelParams[i] = Pointer.to((Pointer) arguments[i]);
+ } else if (arguments[i] instanceof Integer) {
+ kernelParams[i] = Pointer.to(new int[] { (Integer) arguments[i] });
+ } else if (arguments[i] instanceof Double) {
+ kernelParams[i] = Pointer.to(new double[] { (Double) arguments[i] });
+ } else if (arguments[i] instanceof Long) {
+ kernelParams[i] = Pointer.to(new long[] { (Long) arguments[i] });
+ } else if (arguments[i] instanceof Float) {
+ kernelParams[i] = Pointer.to(new float[] { (Float) arguments[i] });
+ } else {
+ throw new DMLRuntimeException("The argument of type " + arguments[i].getClass() + " is not supported.");
}
- else if(arguments[i] instanceof Pointer) {
- kernelParams[i] = Pointer.to((Pointer) arguments[i]);
- }
- else if(arguments[i] instanceof Integer) {
- kernelParams[i] = Pointer.to(new int[]{(Integer) arguments[i]});
- }
- else if(arguments[i] instanceof Double) {
- kernelParams[i] = Pointer.to(new double[]{(Double) arguments[i]});
- }
- else if(arguments[i] instanceof Long) {
- kernelParams[i] = Pointer.to(new long[]{(Long) arguments[i]});
- }
- else if(arguments[i] instanceof Float) {
- kernelParams[i] = Pointer.to(new float[]{(Float) arguments[i]});
- }
- else {
- throw new DMLRuntimeException("The argument of type " + arguments[i].getClass() + " is not supported.");
- }
- }
-
+ }
+
// Launches the kernel using CUDA's driver API.
- checkResult(cuLaunchKernel(function,
- config.gridDimX, config.gridDimY, config.gridDimZ,
- config.blockDimX, config.blockDimY, config.blockDimZ,
- config.sharedMemBytes, config.stream, Pointer.to(kernelParams), null));
- JCuda.cudaDeviceSynchronize();
+ checkResult(cuLaunchKernel(function, config.gridDimX, config.gridDimY, config.gridDimZ, config.blockDimX,
+ config.blockDimY, config.blockDimZ, config.sharedMemBytes, config.stream, Pointer.to(kernelParams),
+ null));
+ JCuda.cudaDeviceSynchronize();
+ }
+
+ public static void checkResult(int cuResult) throws DMLRuntimeException {
+ if (cuResult != CUresult.CUDA_SUCCESS) {
+ throw new DMLRuntimeException(CUresult.stringFor(cuResult));
+ }
}
-
- public static void checkResult(int cuResult) throws DMLRuntimeException {
- if (cuResult != CUresult.CUDA_SUCCESS) {
- throw new DMLRuntimeException(CUresult.stringFor(cuResult));
- }
- }
-
- /**
+
+ /**
* Reads the ptx file from resource
- *
+ *
* @param ptxFileName
* @return
* @throws DMLRuntimeException
@@ -131,30 +125,29 @@ public class JCudaKernels {
private Pointer initKernels(String ptxFileName) throws DMLRuntimeException {
InputStream in = null;
ByteArrayOutputStream out = null;
- try {
- in = JCudaKernels.class.getResourceAsStream(ptxFileName);
- if (in != null) {
- out = new ByteArrayOutputStream();
- byte buffer[] = new byte[8192];
- while (true) {
- int read = in.read(buffer);
- if (read == -1)
- break;
- out.write(buffer, 0, read);
- }
- out.write('\0');
- out.flush();
- return Pointer.to(out.toByteArray());
- }
- else {
- throw new DMLRuntimeException( "The input file " + ptxFileName + " not found. (Hint: Please compile SystemML using -DenableGPU=true flag. Example: mvn package -DenableGPU=true).");
- }
- } catch (IOException e) {
- throw new DMLRuntimeException("Could not initialize the kernels", e);
- }
- finally {
- IOUtilFunctions.closeSilently(out);
- IOUtilFunctions.closeSilently(in);
+ try {
+ in = JCudaKernels.class.getResourceAsStream(ptxFileName);
+ if (in != null) {
+ out = new ByteArrayOutputStream();
+ byte buffer[] = new byte[8192];
+ while (true) {
+ int read = in.read(buffer);
+ if (read == -1)
+ break;
+ out.write(buffer, 0, read);
+ }
+ out.write('\0');
+ out.flush();
+ return Pointer.to(out.toByteArray());
+ } else {
+ throw new DMLRuntimeException("The input file " + ptxFileName
+ + " not found. (Hint: Please compile SystemML using -DenableGPU=true flag. Example: mvn package -DenableGPU=true).");
+ }
+ } catch (IOException e) {
+ throw new DMLRuntimeException("Could not initialize the kernels", e);
+ } finally {
+ IOUtilFunctions.closeSilently(out);
+ IOUtilFunctions.closeSilently(in);
}
}
}
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixCUDA.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixCUDA.java b/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixCUDA.java
index 2453b41..48dd391 100644
--- a/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixCUDA.java
+++ b/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixCUDA.java
@@ -422,17 +422,13 @@ public class LibMatrixCUDA {
if (CONVOLUTION_PREFERENCE == cudnnConvolutionFwdPreference.CUDNN_CONVOLUTION_FWD_NO_WORKSPACE) {
algo = jcuda.jcudnn.cudnnConvolutionFwdAlgo.CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
} else if (CONVOLUTION_PREFERENCE == cudnnConvolutionFwdPreference.CUDNN_CONVOLUTION_FWD_PREFER_FASTEST) {
- int[] algos = {
- jcuda.jcudnn.cudnnConvolutionFwdAlgo.CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM,
- jcuda.jcudnn.cudnnConvolutionFwdAlgo.CUDNN_CONVOLUTION_FWD_ALGO_GEMM,
- jcuda.jcudnn.cudnnConvolutionFwdAlgo.CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
- };
+ int[] algos = {-1};
// TODO: Look into FFt, Winograd, etc
// Also ensure that GPU has enough memory to allocate memory
long sizeInBytesArray[] = {0};
- algo = jcuda.jcudnn.JCudnn.cudnnGetConvolutionForwardAlgorithm(getCudnnHandle(gCtx), srcTensorDesc, filterDesc, convDesc, dstTensorDesc,
+ jcuda.jcudnn.JCudnn.cudnnGetConvolutionForwardAlgorithm(getCudnnHandle(gCtx), srcTensorDesc, filterDesc, convDesc, dstTensorDesc,
CONVOLUTION_PREFERENCE, sizeInBytesArray[0], algos);
- cudnnGetConvolutionForwardWorkspaceSize(getCudnnHandle(gCtx), srcTensorDesc, filterDesc, convDesc, dstTensorDesc, algo, sizeInBytesArray);
+ cudnnGetConvolutionForwardWorkspaceSize(getCudnnHandle(gCtx), srcTensorDesc, filterDesc, convDesc, dstTensorDesc, algos[0], sizeInBytesArray);
if (sizeInBytesArray[0] != 0)
workSpace = gCtx.allocate(sizeInBytesArray[0]);
sizeInBytes = sizeInBytesArray[0];
@@ -1187,7 +1183,7 @@ public class LibMatrixCUDA {
* @throws DMLRuntimeException if an error occurs
*/
public static void relu(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in, String outputName) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
long N = in.getNumRows();
long CHW = in.getNumColumns();
@@ -1237,7 +1233,7 @@ public class LibMatrixCUDA {
public static void matmultTSMM(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject left, String outputName,
boolean isLeftTransposed) throws DMLRuntimeException {
LOG.trace("GPU : matmultTSMM" + ", GPUContext=" + gCtx);
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
if(isInSparseFormat(gCtx, left)) {
// For sparse TSMM, invoke matmult (TODO: possible performance improvement)
@@ -1332,7 +1328,7 @@ public class LibMatrixCUDA {
*/
public static MatrixObject matmult(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject left, MatrixObject right, String outputName,
boolean isLeftTransposed, boolean isRightTransposed) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
LOG.trace("GPU : matmult" + ", GPUContext=" + gCtx);
if(!left.getGPUObject(gCtx).isAllocated() || !right.getGPUObject(gCtx).isAllocated())
@@ -1823,7 +1819,7 @@ public class LibMatrixCUDA {
*/
public static void unaryAggregate(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in1, String output, AggregateUnaryOperator op)
throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
LOG.trace("GPU : unaryAggregate" + ", GPUContext=" + gCtx);
final int REDUCTION_ALL = 1;
@@ -2329,7 +2325,7 @@ public class LibMatrixCUDA {
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
public static void matrixScalarArithmetic(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in, String outputName, boolean isInputTransposed, ScalarOperator op) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
double constant = op.getConstant();
LOG.trace("GPU : matrixScalarArithmetic, scalar: " + constant + ", GPUContext=" + gCtx);
@@ -2409,7 +2405,7 @@ public class LibMatrixCUDA {
*/
public static void matrixScalarArithmetic(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in1, MatrixObject in2,
String outputName, boolean isLeftTransposed, boolean isRightTransposed, BinaryOperator op) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
boolean isCUDALibAvailable = (op.fn instanceof Plus || op.fn instanceof Minus) && !isSparseAndEmpty(gCtx, in1) && !isSparseAndEmpty(gCtx, in2) && !isVector(in1) && !isVector(in2);
if(!isCUDALibAvailable) {
@@ -2447,7 +2443,7 @@ public class LibMatrixCUDA {
*/
private static void matrixScalarOp(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in, String outputName, boolean isInputTransposed,
ScalarOperator op) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
if(isInputTransposed)
throw new DMLRuntimeException("Transposing the input is not supported");
@@ -2503,7 +2499,7 @@ public class LibMatrixCUDA {
*/
private static void matrixMatrixOp(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in1, MatrixObject in2,
String outputName, boolean isLeftTransposed, boolean isRightTransposed, BinaryOperator op) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
boolean isEmpty1 = isSparseAndEmpty(gCtx, in1);
boolean isEmpty2 = isSparseAndEmpty(gCtx, in2);
@@ -2618,7 +2614,7 @@ public class LibMatrixCUDA {
}
private static void deviceCopy(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject src, String outputName, boolean isInputTransposed) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
if(!isInputTransposed)
deviceCopy(ec, gCtx, instName, src, outputName);
@@ -2646,7 +2642,7 @@ public class LibMatrixCUDA {
@SuppressWarnings("unused")
private static void compareAndSet(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in, String outputName, double compareVal, double tolerance,
double ifEqualsVal, double ifLessThanVal, double ifGreaterThanVal) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
Pointer A = getDensePointer(gCtx, in, instName); // TODO: FIXME: Implement sparse kernel
MatrixObject out = ec.getMatrixObject(outputName);
@@ -2674,7 +2670,7 @@ public class LibMatrixCUDA {
* @throws DMLRuntimeException if error
*/
private static void setOutputToConstant(ExecutionContext ec, GPUContext gCtx, String instName, double constant, String outputName) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
if(constant == 0) {
// TODO: Create sparse empty block instead
@@ -2756,7 +2752,7 @@ public class LibMatrixCUDA {
*/
private static void dgeam(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in1, MatrixObject in2, String outputName,
boolean isLeftTransposed, boolean isRightTransposed, double alpha, double beta) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
LOG.trace("GPU : dgeam" + ", GPUContext=" + gCtx);
@@ -2883,7 +2879,7 @@ public class LibMatrixCUDA {
public static void transpose(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in, String outputName) throws DMLRuntimeException {
// C = alpha* op( A ) + beta* op ( B )
// = 1.0 * A^T + 0.0 * A^T
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
dgeam(ec, gCtx, instName, in, in, outputName, true, true, 1.0, 0.0);
}
@@ -3122,7 +3118,7 @@ public class LibMatrixCUDA {
* @throws DMLRuntimeException
*/
private static void unaryOp(ExecutionContext ec, GPUContext gCtx, MatrixObject in1, String kernel, double sparseAndEmptyFillValue, String outputName, String instName, String kernelTimer) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
GPUObject in = in1.getGPUObject(gCtx);
boolean isSparseAndEmpty = in.isSparseAndEmpty();
@@ -3158,7 +3154,7 @@ public class LibMatrixCUDA {
*/
public static void axpy(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in1, MatrixObject in2,
String outputName, double constant) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
Pointer A = getDensePointer(gCtx, in1, instName);
Pointer B = getDensePointer(gCtx, in2, instName);
@@ -3213,7 +3209,7 @@ public class LibMatrixCUDA {
* @throws DMLRuntimeException if an error occurs
*/
public static void solve(ExecutionContext ec, GPUContext gCtx, String instName, MatrixObject in1, MatrixObject in2, String outputName) throws DMLRuntimeException {
- if (ec.getGPUContext() != gCtx)
+ if (ec.getGPUContext(0) != gCtx)
throw new DMLRuntimeException("GPU : Invalid internal state, the GPUContext set with the ExecutionContext is not the same used to run this LibMatrixCUDA function");
// x = solve(A, b)
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixDNNHelper.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixDNNHelper.java b/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixDNNHelper.java
index 40a39f0..5a137cc 100644
--- a/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixDNNHelper.java
+++ b/src/main/java/org/apache/sysml/runtime/matrix/data/LibMatrixDNNHelper.java
@@ -60,6 +60,7 @@ public class LibMatrixDNNHelper {
* Factory method that returns list of callable tasks for performing maxpooling backward operation
*
* @param params convolution parameters
+ * @param performReluBackward whether to perform ReLU backward
* @return list of callable tasks for performing maxpooling backward operation
* @throws DMLRuntimeException if error occurs
*/
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/test/java/org/apache/sysml/test/gpu/GPUTests.java
----------------------------------------------------------------------
diff --git a/src/test/java/org/apache/sysml/test/gpu/GPUTests.java b/src/test/java/org/apache/sysml/test/gpu/GPUTests.java
index 8af9104..ba61dc0 100644
--- a/src/test/java/org/apache/sysml/test/gpu/GPUTests.java
+++ b/src/test/java/org/apache/sysml/test/gpu/GPUTests.java
@@ -20,6 +20,7 @@
package org.apache.sysml.test.gpu;
import java.util.ArrayList;
+import java.util.Formatter;
import java.util.List;
import java.util.Map;
import java.util.Set;
@@ -82,16 +83,14 @@ public abstract class GPUTests extends AutomatedTestBase {
int count = GPUContextPool.getDeviceCount();
int freeCount = GPUContextPool.getAvailableCount();
Assert.assertTrue("All GPUContexts have not been returned to the GPUContextPool", count == freeCount);
- ArrayList<GPUContext> gpuContexts = new ArrayList<>();
- for (int i = 0; i < count; i++) {
- GPUContext gCtx = GPUContextPool.getFromPool();
+
+ List<GPUContext> gCtxs = GPUContextPool.reserveAllGPUContexts();
+ for (GPUContext gCtx : gCtxs) {
gCtx.initializeThread();
gCtx.clearMemory();
- gpuContexts.add(gCtx);
- }
- for (GPUContext gCtx : gpuContexts) {
- GPUContextPool.returnToPool(gCtx);
}
+ GPUContextPool.freeAllGPUContexts();
+
} catch (DMLRuntimeException e) {
// Ignore
@@ -124,6 +123,33 @@ public abstract class GPUTests extends AutomatedTestBase {
}
/**
+ * Generates a random input matrix with a given size and sparsity
+ *
+ * @param spark valid instance of {@link SparkSession}
+ * @param m number of rows
+ * @param n number of columns
+ * @param min min for RNG
+ * @param max max for RNG
+ * @param sparsity sparsity (1 = completely dense, 0 = completely sparse)
+ * @return a random matrix with given size and sparsity
+ */
+ protected Matrix generateInputMatrix(SparkSession spark, int m, int n, double min, double max, double sparsity, int seed) {
+ // Generate a random matrix of size m * n
+ MLContext genMLC = new MLContext(spark);
+ String scriptStr;
+ if (sparsity == 0.0) {
+ scriptStr = "in1 = matrix(0, rows=" + m + ", cols=" + n + ")";
+ } else {
+ scriptStr = "in1 = rand(rows=" + m + ", cols=" + n + ", sparsity = " + sparsity + ", seed= " + seed
+ + ", min=" + min + ", max=" + max + ")";
+ }
+ Script generateScript = ScriptFactory.dmlFromString(scriptStr).out("in1");
+ Matrix in1 = genMLC.execute(generateScript).getMatrix("in1");
+ genMLC.close();
+ return in1;
+ }
+
+ /**
* Asserts that the values in two matrices are in {@link UnaryOpTests#THRESHOLD} of each other
*
* @param expected expected matrix
@@ -145,8 +171,11 @@ public abstract class GPUTests extends AutomatedTestBase {
double actualDouble = actualMB.quickGetValue(i, j);
if (expectedDouble != 0.0 && !Double.isNaN(expectedDouble) && Double.isFinite(expectedDouble)) {
double relativeError = Math.abs((expectedDouble - actualDouble) / expectedDouble);
- Assert.assertTrue("Comparing floating point numbers, relative error(" + relativeError
- + ") is more than threshold (" + getTHRESHOLD() + ")", relativeError < getTHRESHOLD());
+ Formatter format = new Formatter();
+ format.format(
+ "Relative error(%f) is more than threshold (%f). Expected = %f, Actual = %f, differed at [%d, %d]",
+ relativeError, getTHRESHOLD(), expectedDouble, actualDouble, i, j);
+ Assert.assertTrue(format.toString(), relativeError < getTHRESHOLD());
} else {
Assert.assertEquals(expectedDouble, actualDouble, getTHRESHOLD());
}
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/test/java/org/apache/sysml/test/gpu/NeuralNetworkOpTests.java
----------------------------------------------------------------------
diff --git a/src/test/java/org/apache/sysml/test/gpu/NeuralNetworkOpTests.java b/src/test/java/org/apache/sysml/test/gpu/NeuralNetworkOpTests.java
index f1f1ea5..f4e931b 100644
--- a/src/test/java/org/apache/sysml/test/gpu/NeuralNetworkOpTests.java
+++ b/src/test/java/org/apache/sysml/test/gpu/NeuralNetworkOpTests.java
@@ -29,6 +29,8 @@ import org.apache.sysml.runtime.instructions.gpu.context.GPUContext;
import org.apache.sysml.runtime.instructions.gpu.context.GPUContextPool;
import org.apache.sysml.runtime.util.ConvolutionUtils;
import org.apache.sysml.test.utils.TestUtils;
+import org.junit.Assert;
+import org.junit.Ignore;
import org.junit.Test;
/**
@@ -37,8 +39,6 @@ import org.junit.Test;
public class NeuralNetworkOpTests extends GPUTests {
private final static String TEST_NAME = "NeuralNetworkOpTests";
- private final int seed = 42;
-
// The MAX_OP_SIZE is to take into consideration the memory available on the GPU as well as
// limits set by cudnn (operands need to be less than 2GB)
private static final double MAX_OP_SIZE;
@@ -47,17 +47,19 @@ public class NeuralNetworkOpTests extends GPUTests {
double MAX = 0.5 * 1024 * 1024 * 1024; // 0.5 GB (this HAS to be less than 2GB)
try {
// Cap the maximum allowed operand size to 1/3rd of the usable GPU memory or MAX, whichever is lesser
- GPUContext gCtx = GPUContextPool.getFromPool();
- long availableMemory = gCtx.getAvailableMemory();
+ List<GPUContext> gCtxs = GPUContextPool.reserveAllGPUContexts();
+ long availableMemory = gCtxs.get(0).getAvailableMemory();
double averageMemoryPerOperand = availableMemory / 3.0;
MAX_OP_SIZE = Math.min(averageMemoryPerOperand, MAX);
- GPUContextPool.returnToPool(gCtx);
+ GPUContextPool.freeAllGPUContexts();
} catch (DMLRuntimeException e) {
throw new RuntimeException(e);
}
}
+ private final int seed = 42;
+
// More comprehensive but time consuming tests
/*
private final List<Integer> Nlst = Arrays.asList(128, 64, 32);
@@ -73,14 +75,13 @@ public class NeuralNetworkOpTests extends GPUTests {
private final List<Integer> padWidthLst = Arrays.asList(3, 1);
private final List<Double> sparsitylst = Arrays.asList(1.0); // Only test for dense
*/
-
private final List<Integer> Nlst = Arrays.asList(128, 64);
private final List<Integer> Clst = Arrays.asList(30, 3);
private final List<Integer> Hlst = Arrays.asList(256, 64);
private final List<Integer> Wlst = Arrays.asList(256, 64);
private final List<Integer> Klst = Arrays.asList(30, 20);
- private final List<Integer> Rlst = Arrays.asList(128, 3, 1);
- private final List<Integer> Slst = Arrays.asList(128, 3, 1);
+ private final List<Integer> Rlst = Arrays.asList(128, 3);
+ private final List<Integer> Slst = Arrays.asList(128, 3);
private final List<Integer> strideHeightLst = Arrays.asList(9, 1);
private final List<Integer> strideWidthLst = Arrays.asList(9, 1);
private final List<Integer> padHeightLst = Arrays.asList(3, 1);
@@ -143,9 +144,9 @@ public class NeuralNetworkOpTests extends GPUTests {
filterSizeInMB, N, K, P, Q, doutSizeInMB,
strideH, strideW, padH, padW);
Matrix image = generateInputMatrix(spark, (int) N,
- (int) (C * H * W), sparsity, seed);
+ (int) (C * H * W), -127, 127, sparsity, seed);
Matrix filter = generateInputMatrix(spark, (int) K,
- (int) (C * R * S), sparsity, seed);
+ (int) (C * R * S), -127, 127, sparsity, seed);
HashMap<String, Object> inputs = new HashMap<>();
inputs.put("N", N);
inputs.put("C", C);
@@ -181,6 +182,77 @@ public class NeuralNetworkOpTests extends GPUTests {
}
}
+ @Ignore
+ @Test
+ /**
+ * Ignored test to iron out issues
+ */ public void testConv2dOneCase() {
+ String scriptStr = "O = conv2d(image, filter, padding=[padH, padW], stride=[strideH, strideW], input_shape=[N,C,H,W], filter_shape=[K,C,R,S]); print(toString(O, decimal=10, rows=1000, cols=1000))";
+
+ long N = 32;
+ long C = 3;
+ long H = 128;
+ long W = 128;
+
+ long K = 30;
+ long R = 64;
+ long S = 64;
+
+ long padH = 9;
+ long padW = 9;
+ long strideH = 3;
+ long strideW = 3;
+
+ double sparsity = 1.0;
+
+ // Make sure ops fit in GPU memory and within constraints of cudnn
+ long imageSize = N * C * H * W * 8l;
+ if (imageSize > MAX_OP_SIZE) // image size
+ Assert.fail();
+ long filterSize = K * C * R * S * 8l;
+ if (filterSize > MAX_OP_SIZE) // filter size
+ Assert.fail();
+ // filter is smaller than image + padding
+ if (R > (H + padH) || S > (W + padW))
+ Assert.fail();
+
+ int P = (int) ConvolutionUtils.getP(H, R, strideH, padH);
+ int Q = (int) ConvolutionUtils.getQ(W, S, strideW, padW);
+
+ long doutSize = N * K * P * Q * 8l;
+ if (doutSize > MAX_OP_SIZE) // dout/output size
+ Assert.fail();
+
+ double imageSizeInMB = imageSize / (1024.0 * 1024.0);
+ double filterSizeInMB = filterSize / (1024.0 * 1024.0);
+ double doutSizeInMB = doutSize / (1024.0 * 1024.0);
+ System.out
+ .format("conv2d, image[%d,%d,%d,%d](%.1fMB), filter[%d,%d,%d,%d](%.1f), dout[%d,%d,%d,%d](%.1fMB), stride[%d,%d], padding[%d,%d]",
+ N, C, H, W, imageSizeInMB, N, C, R, S, filterSizeInMB, N, K, P, Q, doutSizeInMB, strideH,
+ strideW, padH, padW);
+ Matrix image = generateInputMatrix(spark, (int) N, (int) (C * H * W), -1, 1, sparsity, seed);
+ Matrix filter = generateInputMatrix(spark, (int) K, (int) (C * R * S), -1, 1.0, sparsity, seed);
+ HashMap<String, Object> inputs = new HashMap<>();
+ inputs.put("N", N);
+ inputs.put("C", C);
+ inputs.put("H", H);
+ inputs.put("W", W);
+ inputs.put("K", K);
+ inputs.put("R", R);
+ inputs.put("S", S);
+ inputs.put("strideH", strideH);
+ inputs.put("strideW", strideW);
+ inputs.put("padH", padH);
+ inputs.put("padW", padW);
+ inputs.put("image", image);
+ inputs.put("filter", filter);
+ List<Object> outCPU = runOnCPU(spark, scriptStr, inputs, Arrays.asList("O"));
+ List<Object> outGPU = runOnGPU(spark, scriptStr, inputs, Arrays.asList("O"));
+ assertHeavyHitterPresent("gpu_conv2d");
+ assertEqualObjects(outCPU.get(0), outGPU.get(0));
+ clearGPUMemory();
+ }
+
@Test
public void testConv2dBackwardFilter() {
String scriptStr = "O = conv2d_backward_filter(image, dout, padding=[padH, padW], stride=[strideH, strideW], input_shape=[N,C,H,W], filter_shape=[K,C,R,S])";
@@ -226,9 +298,9 @@ public class NeuralNetworkOpTests extends GPUTests {
filterSizeInMB, N, K, P, Q, doutSizeInMB,
strideH, strideW, padH, padW);
Matrix image = generateInputMatrix(spark, (int) N,
- (int) (C * H * W), sparsity, seed);
+ (int) (C * H * W), 0.-127.0, 127, sparsity, seed);
Matrix dout = generateInputMatrix(spark, (int) N,
- (int) (K * P * Q), sparsity, seed);
+ (int) (K * P * Q), 0.-127.0, 127, sparsity, seed);
HashMap<String, Object> inputs = new HashMap<>();
inputs.put("N", N);
inputs.put("C", C);
@@ -310,9 +382,9 @@ public class NeuralNetworkOpTests extends GPUTests {
strideH, strideW, padH, padW);
Matrix filter = generateInputMatrix(spark, (int) K,
- (int) (C * R * S), sparsity, seed);
+ (int) (C * R * S), 0.-127.0, 127, sparsity, seed);
Matrix dout = generateInputMatrix(spark, (int) N,
- (int) (K * P * Q), sparsity, seed);
+ (int) (K * P * Q), 0.-127.0, 127, sparsity, seed);
HashMap<String, Object> inputs = new HashMap<>();
inputs.put("N", N);
inputs.put("C", C);
@@ -392,7 +464,7 @@ public class NeuralNetworkOpTests extends GPUTests {
P, Q, doutSizeInMB, strideH, strideW, padH, padW);
Matrix image = generateInputMatrix(spark, (int) N,
- (int) (C * H * W), sparsity, seed);
+ (int) (C * H * W), 0.-127.0, 127, sparsity, seed);
HashMap<String, Object> inputs = new HashMap<>();
inputs.put("N", N);
inputs.put("C", C);
@@ -469,9 +541,9 @@ public class NeuralNetworkOpTests extends GPUTests {
P, Q, doutSizeInMB, strideH, strideW, padH, padW);
Matrix image = generateInputMatrix(spark, (int) N,
- (int) (C * H * W), sparsity, seed);
+ (int) (C * H * W), 0.-127.0, 127, sparsity, seed);
Matrix dout = generateInputMatrix(spark, (int) N, (int) (C * P * Q),
- sparsity, seed);
+ 0.-127.0, 127, sparsity, seed);
HashMap<String, Object> inputs = new HashMap<>();
inputs.put("N", N);
inputs.put("C", C);
[3/3] systemml git commit: [FIX] Fixed nested parfor for GPUs
Posted by na...@apache.org.
[FIX] Fixed nested parfor for GPUs
Additionally
- Fixed intellij codestyle accordingly
- Fixed formatting of some GPU related source files
Closes #532
Project: http://git-wip-us.apache.org/repos/asf/systemml/repo
Commit: http://git-wip-us.apache.org/repos/asf/systemml/commit/f5871756
Tree: http://git-wip-us.apache.org/repos/asf/systemml/tree/f5871756
Diff: http://git-wip-us.apache.org/repos/asf/systemml/diff/f5871756
Branch: refs/heads/master
Commit: f587175647a84a3825b174b4d29c0398be17331f
Parents: 0bcae49
Author: Nakul Jindal <na...@gmail.com>
Authored: Sat Jun 10 12:06:47 2017 -0700
Committer: Nakul Jindal <na...@gmail.com>
Committed: Sat Jun 10 12:06:47 2017 -0700
----------------------------------------------------------------------
dev/code-style/systemml-style-intellij.xml | 37 +-
.../apache/sysml/api/ScriptExecutorUtils.java | 17 +-
.../controlprogram/ParForProgramBlock.java | 19 +-
.../context/ExecutionContext.java | 64 +-
.../controlprogram/parfor/LocalParWorker.java | 2 +-
.../cp/FunctionCallCPInstruction.java | 12 +-
.../gpu/AggregateBinaryGPUInstruction.java | 4 +-
.../gpu/AggregateUnaryGPUInstruction.java | 2 +-
.../gpu/ConvolutionGPUInstruction.java | 18 +-
.../instructions/gpu/MMTSJGPUInstruction.java | 2 +-
.../gpu/MatrixBuiltinGPUInstruction.java | 30 +-
.../MatrixMatrixArithmeticGPUInstruction.java | 2 +-
.../gpu/MatrixMatrixAxpyGPUInstruction.java | 2 +-
.../gpu/MatrixMatrixBuiltinGPUInstruction.java | 2 +-
.../instructions/gpu/ReorgGPUInstruction.java | 2 +-
.../ScalarMatrixArithmeticGPUInstruction.java | 2 +-
.../instructions/gpu/context/CSRPointer.java | 922 ++++++-------
.../gpu/context/ExecutionConfig.java | 85 +-
.../instructions/gpu/context/GPUContext.java | 1257 +++++++++---------
.../gpu/context/GPUContextPool.java | 266 ++--
.../instructions/gpu/context/GPUObject.java | 454 ++++---
.../instructions/gpu/context/JCudaKernels.java | 141 +-
.../runtime/matrix/data/LibMatrixCUDA.java | 42 +-
.../runtime/matrix/data/LibMatrixDNNHelper.java | 1 +
.../org/apache/sysml/test/gpu/GPUTests.java | 47 +-
.../sysml/test/gpu/NeuralNetworkOpTests.java | 106 +-
26 files changed, 1917 insertions(+), 1621 deletions(-)
----------------------------------------------------------------------
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/dev/code-style/systemml-style-intellij.xml
----------------------------------------------------------------------
diff --git a/dev/code-style/systemml-style-intellij.xml b/dev/code-style/systemml-style-intellij.xml
index 1ad3209..b4a53b4 100644
--- a/dev/code-style/systemml-style-intellij.xml
+++ b/dev/code-style/systemml-style-intellij.xml
@@ -1,28 +1,27 @@
<!--
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
+* Licensed to the Apache Software Foundation (ASF) under one
+* or more contributor license agreements. See the NOTICE file
+* distributed with this work for additional information
+* regarding copyright ownership. The ASF licenses this file
+* to you under the Apache License, Version 2.0 (the
+* "License"); you may not use this file except in compliance
+* with the License. You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an
+* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+* KIND, either express or implied. See the License for the
+* specific language governing permissions and limitations
+* under the License.
-->
-
<code_scheme name="SystemML Format">
<option name="CLASS_COUNT_TO_USE_IMPORT_ON_DEMAND" value="999" />
<option name="NAMES_COUNT_TO_USE_IMPORT_ON_DEMAND" value="999" />
<option name="IMPORT_LAYOUT_TABLE">
<value>
- <package name="" withSubpackages="true" static="false" />
+ <package name="" withSubpackages="true" static="true" />
<emptyLine />
<package name="java" withSubpackages="true" static="false" />
<emptyLine />
@@ -32,7 +31,7 @@
<emptyLine />
<package name="com" withSubpackages="true" static="false" />
<emptyLine />
- <package name="" withSubpackages="true" static="true" />
+ <package name="" withSubpackages="true" static="false" />
</value>
</option>
<codeStyleSettings language="JAVA">
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/api/ScriptExecutorUtils.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/api/ScriptExecutorUtils.java b/src/main/java/org/apache/sysml/api/ScriptExecutorUtils.java
index 674a011..2895aa4 100644
--- a/src/main/java/org/apache/sysml/api/ScriptExecutorUtils.java
+++ b/src/main/java/org/apache/sysml/api/ScriptExecutorUtils.java
@@ -19,6 +19,8 @@
package org.apache.sysml.api;
+import java.util.List;
+
import org.apache.sysml.api.mlcontext.ScriptExecutor;
import org.apache.sysml.conf.DMLConfig;
import org.apache.sysml.hops.codegen.SpoofCompiler;
@@ -79,23 +81,22 @@ public class ScriptExecutorUtils {
// GPUs
GPUContextPool.PER_PROCESS_MAX_GPUS = dmlconf.getIntValue(DMLConfig.MAX_GPUS_PER_PROCESS);
Statistics.startRunTimer();
- GPUContext gCtx = null;
try {
// run execute (w/ exception handling to ensure proper shutdown)
if (DMLScript.USE_ACCELERATOR && ec != null) {
- gCtx = GPUContextPool.getFromPool();
- if (gCtx == null) {
+ List<GPUContext> gCtxs = GPUContextPool.reserveAllGPUContexts();
+ if (gCtxs == null) {
throw new DMLRuntimeException(
"GPU : Could not create GPUContext, either no GPU or all GPUs currently in use");
}
- gCtx.initializeThread();
- ec.setGPUContext(gCtx);
+ gCtxs.get(0).initializeThread();
+ ec.setGPUContexts(gCtxs);
}
rtprog.execute(ec);
} finally { // ensure cleanup/shutdown
- if (DMLScript.USE_ACCELERATOR && ec.getGPUContext() != null) {
- ec.getGPUContext().clearTemporaryMemory();
- GPUContextPool.returnToPool(ec.getGPUContext());
+ if (DMLScript.USE_ACCELERATOR && !ec.getGPUContexts().isEmpty()) {
+ ec.getGPUContexts().forEach(gCtx -> gCtx.clearTemporaryMemory());
+ GPUContextPool.freeAllGPUContexts();
}
if (dmlconf.getBooleanValue(DMLConfig.CODEGEN))
SpoofCompiler.cleanupCodeGenerator();
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/controlprogram/ParForProgramBlock.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/controlprogram/ParForProgramBlock.java b/src/main/java/org/apache/sysml/runtime/controlprogram/ParForProgramBlock.java
index c9dcc22..95e28e7 100644
--- a/src/main/java/org/apache/sysml/runtime/controlprogram/ParForProgramBlock.java
+++ b/src/main/java/org/apache/sysml/runtime/controlprogram/ParForProgramBlock.java
@@ -632,9 +632,7 @@ public class ParForProgramBlock extends ForProgramBlock
{
case LOCAL: //create parworkers as local threads
if (DMLScript.USE_ACCELERATOR) {
- GPUContextPool.returnToPool(ec.getGPUContext());
- ec.setGPUContext(null);
- setDegreeOfParallelism(GPUContextPool.getDeviceCount());
+ setDegreeOfParallelism(ec.getNumGPUContexts());
}
executeLocalParFor(ec, iterVar, from, to, incr);
break;
@@ -757,7 +755,7 @@ public class ParForProgramBlock extends ForProgramBlock
{
//create parallel workers as (lazy) deep copies
//including preparation of update-in-place variables
- workers[i] = createParallelWorker( _pwIDs[i], queue, ec );
+ workers[i] = createParallelWorker( _pwIDs[i], queue, ec, i);
threads[i] = new Thread( workers[i] );
threads[i].setPriority(Thread.MAX_PRIORITY);
}
@@ -833,11 +831,9 @@ public class ParForProgramBlock extends ForProgramBlock
// the main thread to use the GPUContext
if (DMLScript.USE_ACCELERATOR) {
for (int i = 0; i < _numThreads; i++) {
- GPUContext gCtx = workers[i].getExecutionContext().getGPUContext();
- GPUContextPool.returnToPool(gCtx);
+ workers[i].getExecutionContext().setGPUContexts(null);
}
- ec.setGPUContext(GPUContextPool.getFromPool());
- ec.getGPUContext().initializeThread();
+ ec.getGPUContext(0).initializeThread();
}
}
finally
@@ -1386,10 +1382,11 @@ public class ParForProgramBlock extends ForProgramBlock
* @param pwID parworker id
* @param queue task queue
* @param ec execution context
+ * @param index the index of the worker
* @return local parworker
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
- private LocalParWorker createParallelWorker(long pwID, LocalTaskQueue<Task> queue, ExecutionContext ec)
+ private LocalParWorker createParallelWorker(long pwID, LocalTaskQueue<Task> queue, ExecutionContext ec, int index)
throws DMLRuntimeException
{
LocalParWorker pw = null;
@@ -1420,9 +1417,9 @@ public class ParForProgramBlock extends ForProgramBlock
ExecutionContext cpEc = ProgramConverter.createDeepCopyExecutionContext(ec);
// If GPU mode is enabled, gets a GPUContext from the pool of GPUContexts
- // and sets it in the ExecutionContext
+ // and sets it in the ExecutionContext of the parfor
if (DMLScript.USE_ACCELERATOR){
- cpEc.setGPUContext(GPUContextPool.getFromPool());
+ cpEc.setGPUContexts(Arrays.asList(ec.getGPUContext(index)));
}
//prepare basic update-in-place variables (vars dropped on result merge)
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/controlprogram/context/ExecutionContext.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/controlprogram/context/ExecutionContext.java b/src/main/java/org/apache/sysml/runtime/controlprogram/context/ExecutionContext.java
index 735f394..fb179f5 100644
--- a/src/main/java/org/apache/sysml/runtime/controlprogram/context/ExecutionContext.java
+++ b/src/main/java/org/apache/sysml/runtime/controlprogram/context/ExecutionContext.java
@@ -20,7 +20,9 @@
package org.apache.sysml.runtime.controlprogram.context;
import java.util.ArrayList;
+import java.util.Collection;
import java.util.HashMap;
+import java.util.List;
import org.apache.commons.logging.Log;
import org.apache.commons.logging.LogFactory;
@@ -66,7 +68,10 @@ public class ExecutionContext {
//debugging (optional)
protected DebugState _dbState = null;
- protected GPUContext _gpuContext = null;
+ /**
+ * List of {@link GPUContext}s owned by this {@link ExecutionContext}
+ */
+ protected List<GPUContext> _gpuContexts = new ArrayList<>();
protected ExecutionContext()
{
@@ -99,13 +104,42 @@ public class ExecutionContext {
_variables = vars;
}
- public GPUContext getGPUContext() {
- return _gpuContext;
+ /**
+ * Get the i-th GPUContext
+ * @param index index of the GPUContext
+ * @return a valid GPUContext or null if the indexed GPUContext does not exist.
+ */
+ public GPUContext getGPUContext(int index) {
+ try {
+ return _gpuContexts.get(index);
+ } catch (IndexOutOfBoundsException e){
+ return null;
+ }
}
- public void setGPUContext(GPUContext _gpuContext) {
- this._gpuContext = _gpuContext;
- }
+ /**
+ * Sets the list of GPUContexts
+ * @param gpuContexts a collection of GPUContexts
+ */
+ public void setGPUContexts(List<GPUContext> gpuContexts){
+ _gpuContexts = gpuContexts;
+ }
+
+ /**
+ * Gets the list of GPUContexts
+ * @return a list of GPUContexts
+ */
+ public List<GPUContext> getGPUContexts() {
+ return _gpuContexts;
+ }
+
+ /**
+ * Gets the number of GPUContexts
+ * @return number of GPUContexts
+ */
+ public int getNumGPUContexts() {
+ return _gpuContexts.size();
+ }
/* -------------------------------------------------------
* Methods to handle variables and associated data objects
@@ -238,7 +272,7 @@ public class ExecutionContext {
throws DMLRuntimeException
{
MatrixObject mo = allocateGPUMatrixObject(varName);
- boolean allocated = mo.getGPUObject(getGPUContext()).acquireDeviceModifyDense();
+ boolean allocated = mo.getGPUObject(getGPUContext(0)).acquireDeviceModifyDense();
mo.getMatrixCharacteristics().setNonZeros(-1);
return new Pair<MatrixObject, Boolean>(mo, allocated);
}
@@ -257,7 +291,7 @@ public class ExecutionContext {
{
MatrixObject mo = allocateGPUMatrixObject(varName);
mo.getMatrixCharacteristics().setNonZeros(nnz);
- boolean allocated = mo.getGPUObject(getGPUContext()).acquireDeviceModifySparse();
+ boolean allocated = mo.getGPUObject(getGPUContext(0)).acquireDeviceModifySparse();
return new Pair<MatrixObject, Boolean>(mo, allocated);
}
@@ -269,12 +303,12 @@ public class ExecutionContext {
*/
public MatrixObject allocateGPUMatrixObject(String varName) throws DMLRuntimeException {
MatrixObject mo = getMatrixObject(varName);
- if( mo.getGPUObject(getGPUContext()) == null ) {
- GPUObject newGObj = getGPUContext().createGPUObject(mo);
+ if( mo.getGPUObject(getGPUContext(0)) == null ) {
+ GPUObject newGObj = getGPUContext(0).createGPUObject(mo);
// The lock is added here for an output block
// so that any block currently in use is not deallocated by eviction on the GPU
newGObj.addLock();
- mo.setGPUObject(getGPUContext(), newGObj);
+ mo.setGPUObject(getGPUContext(0), newGObj);
}
return mo;
}
@@ -282,7 +316,7 @@ public class ExecutionContext {
public Pair<MatrixObject, Boolean> getMatrixInputForGPUInstruction(String varName)
throws DMLRuntimeException
{
- GPUContext gCtx = getGPUContext();
+ GPUContext gCtx = getGPUContext(0);
boolean copied = false;
MatrixObject mo = getMatrixObject(varName);
if(mo == null) {
@@ -322,7 +356,7 @@ public class ExecutionContext {
throws DMLRuntimeException
{
MatrixObject mo = getMatrixObject(varName);
- mo.getGPUObject(getGPUContext()).releaseInput();
+ mo.getGPUObject(getGPUContext(0)).releaseInput();
}
/**
@@ -374,10 +408,10 @@ public class ExecutionContext {
public void releaseMatrixOutputForGPUInstruction(String varName) throws DMLRuntimeException {
MatrixObject mo = getMatrixObject(varName);
- if(mo.getGPUObject(getGPUContext()) == null || !mo.getGPUObject(getGPUContext()).isAllocated()) {
+ if(mo.getGPUObject(getGPUContext(0)) == null || !mo.getGPUObject(getGPUContext(0)).isAllocated()) {
throw new DMLRuntimeException("No output is allocated on GPU");
}
- mo.getGPUObject(getGPUContext()).releaseOutput();
+ mo.getGPUObject(getGPUContext(0)).releaseOutput();
}
public void setMatrixOutput(String varName, MatrixBlock outputData)
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/controlprogram/parfor/LocalParWorker.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/controlprogram/parfor/LocalParWorker.java b/src/main/java/org/apache/sysml/runtime/controlprogram/parfor/LocalParWorker.java
index c4684ec..636b1f8 100644
--- a/src/main/java/org/apache/sysml/runtime/controlprogram/parfor/LocalParWorker.java
+++ b/src/main/java/org/apache/sysml/runtime/controlprogram/parfor/LocalParWorker.java
@@ -83,7 +83,7 @@ public class LocalParWorker extends ParWorker implements Runnable
// Initialize this GPUContext to this thread
if (DMLScript.USE_ACCELERATOR)
- _ec.getGPUContext().initializeThread();
+ _ec.getGPUContext(0).initializeThread();
//setup compiler config for worker thread
ConfigurationManager.setLocalConfig(_cconf);
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/cp/FunctionCallCPInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/cp/FunctionCallCPInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/cp/FunctionCallCPInstruction.java
index 9cc6bb2..3cd2633 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/cp/FunctionCallCPInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/cp/FunctionCallCPInstruction.java
@@ -168,9 +168,9 @@ public class FunctionCallCPInstruction extends CPInstruction
// and copy the function arguments into the created table.
ExecutionContext fn_ec = ExecutionContextFactory.createContext(false, ec.getProgram());
if (DMLScript.USE_ACCELERATOR) {
- fn_ec.setGPUContext(ec.getGPUContext());
- ec.setGPUContext(null);
- fn_ec.getGPUContext().initializeThread();
+ fn_ec.setGPUContexts(ec.getGPUContexts());
+ ec.setGPUContexts(null);
+ fn_ec.getGPUContext(0).initializeThread();
}
fn_ec.setVariables(functionVariables);
// execute the function block
@@ -206,9 +206,9 @@ public class FunctionCallCPInstruction extends CPInstruction
ec.unpinVariables(_boundInputParamNames, pinStatus);
if (DMLScript.USE_ACCELERATOR) {
- ec.setGPUContext(fn_ec.getGPUContext());
- fn_ec.setGPUContext(null);
- ec.getGPUContext().initializeThread();
+ ec.setGPUContexts(fn_ec.getGPUContexts());
+ fn_ec.setGPUContexts(null);
+ ec.getGPUContext(0).initializeThread();
}
// add the updated binding for each return variable to the variables in original symbol table
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateBinaryGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateBinaryGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateBinaryGPUInstruction.java
index 2531c17..0c0a4b2 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateBinaryGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateBinaryGPUInstruction.java
@@ -96,7 +96,7 @@ public class AggregateBinaryGPUInstruction extends GPUInstruction
int clen = (int) (_isRightTransposed ? m2.getNumRows() : m2.getNumColumns());
ec.setMetaData(_output.getName(), rlen, clen);
- LibMatrixCUDA.matmult(ec, ec.getGPUContext(), getExtendedOpcode(), m1, m2, _output.getName(), _isLeftTransposed, _isRightTransposed);
+ LibMatrixCUDA.matmult(ec, ec.getGPUContext(0), getExtendedOpcode(), m1, m2, _output.getName(), _isLeftTransposed, _isRightTransposed);
//release inputs/outputs
ec.releaseMatrixInputForGPUInstruction(_input1.getName());
@@ -113,6 +113,6 @@ public class AggregateBinaryGPUInstruction extends GPUInstruction
@SuppressWarnings("unused")
private boolean isSparse(ExecutionContext ec, String var) throws DMLRuntimeException {
MatrixObject mo = ec.getMatrixObject(var);
- return LibMatrixCUDA.isInSparseFormat(ec.getGPUContext(), mo);
+ return LibMatrixCUDA.isInSparseFormat(ec.getGPUContext(0), mo);
}
}
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateUnaryGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateUnaryGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateUnaryGPUInstruction.java
index 8bdd397..5d01820 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateUnaryGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/AggregateUnaryGPUInstruction.java
@@ -93,7 +93,7 @@ public class AggregateUnaryGPUInstruction extends GPUInstruction {
ec.setMetaData(_output.getName(), rlen, 1);
}
- LibMatrixCUDA.unaryAggregate(ec, ec.getGPUContext(), getExtendedOpcode(), in1, _output.getName(), (AggregateUnaryOperator)_optr);
+ LibMatrixCUDA.unaryAggregate(ec, ec.getGPUContext(0), getExtendedOpcode(), in1, _output.getName(), (AggregateUnaryOperator)_optr);
//release inputs/outputs
ec.releaseMatrixInputForGPUInstruction(_input1.getName());
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/ConvolutionGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/ConvolutionGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/ConvolutionGPUInstruction.java
index 9d4cd1f..e5ea097 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/ConvolutionGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/ConvolutionGPUInstruction.java
@@ -186,9 +186,9 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), input.getNumRows(), input.getNumColumns());
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
if(instOpcode.equalsIgnoreCase("bias_add"))
- LibMatrixCUDA.biasAdd(ec.getGPUContext(), getExtendedOpcode(), input, bias, out);
+ LibMatrixCUDA.biasAdd(ec.getGPUContext(0), getExtendedOpcode(), input, bias, out);
else if(instOpcode.equalsIgnoreCase("bias_multiply"))
- LibMatrixCUDA.biasMultiply(ec.getGPUContext(), getExtendedOpcode(), input, bias, out);
+ LibMatrixCUDA.biasMultiply(ec.getGPUContext(0), getExtendedOpcode(), input, bias, out);
// release inputs/outputs
ec.releaseMatrixInputForGPUInstruction(_input1.getName());
ec.releaseMatrixInputForGPUInstruction(_input2.getName());
@@ -202,7 +202,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
ec.setMetaData(_output.getName(), input.getNumRows(), input.getNumColumns());
- LibMatrixCUDA.reluBackward(ec.getGPUContext(), getExtendedOpcode(), input, dout, out);
+ LibMatrixCUDA.reluBackward(ec.getGPUContext(0), getExtendedOpcode(), input, dout, out);
// release inputs/outputs
ec.releaseMatrixInputForGPUInstruction(_input1.getName());
ec.releaseMatrixInputForGPUInstruction(_input2.getName());
@@ -253,7 +253,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), N, K * P * Q);
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
- LibMatrixCUDA.conv2d(ec.getGPUContext(), getExtendedOpcode(), image, filter, out, N, C, H, W,
+ LibMatrixCUDA.conv2d(ec.getGPUContext(0), getExtendedOpcode(), image, filter, out, N, C, H, W,
K, R, S, pad_h, pad_w, stride_h, stride_w, P, Q);
}
else if (instOpcode.equalsIgnoreCase("conv2d_bias_add")) {
@@ -268,7 +268,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), N, K * P * Q);
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
- LibMatrixCUDA.conv2dBiasAdd(ec.getGPUContext(), getExtendedOpcode(), image, bias, filter, out, N, C, H, W,
+ LibMatrixCUDA.conv2dBiasAdd(ec.getGPUContext(0), getExtendedOpcode(), image, bias, filter, out, N, C, H, W,
K, R, S, pad_h, pad_w, stride_h, stride_w, P, Q);
}
else if (instOpcode.equalsIgnoreCase("conv2d_backward_filter")) {
@@ -283,7 +283,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), K, C * R * S);
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
- LibMatrixCUDA.conv2dBackwardFilter(ec.getGPUContext(), getExtendedOpcode(), image, dout, out, N, C, H, W,
+ LibMatrixCUDA.conv2dBackwardFilter(ec.getGPUContext(0), getExtendedOpcode(), image, dout, out, N, C, H, W,
K, R, S, pad_h, pad_w, stride_h, stride_w, P, Q);
// TODO: For now always copy the device data to host
// ec.gpuCtx.copyDeviceToHost(outputBlock);
@@ -300,7 +300,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), N, C * H * W);
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
- LibMatrixCUDA.conv2dBackwardData(ec.getGPUContext(), getExtendedOpcode(), filter, dout, out, N, C, H, W,
+ LibMatrixCUDA.conv2dBackwardData(ec.getGPUContext(0), getExtendedOpcode(), filter, dout, out, N, C, H, W,
K, R, S, pad_h, pad_w, stride_h, stride_w, P, Q);
}
else if (instOpcode.equalsIgnoreCase("maxpooling")) {
@@ -313,7 +313,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), N, C * P * Q);
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
if(instOpcode.equalsIgnoreCase("maxpooling"))
- LibMatrixCUDA.maxpooling(ec.getGPUContext(), getExtendedOpcode(), image, out, N, C, H, W,
+ LibMatrixCUDA.maxpooling(ec.getGPUContext(0), getExtendedOpcode(), image, out, N, C, H, W,
K, R, S, pad_h, pad_w, stride_h, stride_w, P, Q);
}
else if (instOpcode.equalsIgnoreCase("maxpooling_backward")) {
@@ -328,7 +328,7 @@ public class ConvolutionGPUInstruction extends GPUInstruction
ec.setMetaData(_output.getName(), N, C * H * W);
MatrixObject out = getDenseMatrixOutputForGPUInstruction(ec, _output.getName());
- LibMatrixCUDA.maxpoolingBackward(ec.getGPUContext(), getExtendedOpcode(), image, dout, out, N, C, H, W,
+ LibMatrixCUDA.maxpoolingBackward(ec.getGPUContext(0), getExtendedOpcode(), image, dout, out, N, C, H, W,
K, R, S, pad_h, pad_w, stride_h, stride_w, P, Q);
}
else {
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/MMTSJGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MMTSJGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MMTSJGPUInstruction.java
index c147a6f..55656f0 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MMTSJGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MMTSJGPUInstruction.java
@@ -107,7 +107,7 @@ public class MMTSJGPUInstruction extends GPUInstruction
//execute operations
ec.setMetaData(_output.getName(), rlen, clen);
- LibMatrixCUDA.matmultTSMM(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName(), isLeftTransposed);
+ LibMatrixCUDA.matmultTSMM(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName(), isLeftTransposed);
ec.releaseMatrixInputForGPUInstruction(_input.getName());
ec.releaseMatrixOutputForGPUInstruction(_output.getName());
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixBuiltinGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixBuiltinGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixBuiltinGPUInstruction.java
index 7b50285..beeacee 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixBuiltinGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixBuiltinGPUInstruction.java
@@ -44,35 +44,35 @@ public class MatrixBuiltinGPUInstruction extends BuiltinUnaryGPUInstruction {
switch(opcode) {
case "sel+":
- LibMatrixCUDA.relu(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.relu(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "exp":
- LibMatrixCUDA.exp(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.exp(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "sqrt":
- LibMatrixCUDA.sqrt(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.sqrt(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "log":
- LibMatrixCUDA.log(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.log(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "round":
- LibMatrixCUDA.round(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.round(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "floor":
- LibMatrixCUDA.floor(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.floor(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "ceil":
- LibMatrixCUDA.ceil(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.ceil(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "abs":
- LibMatrixCUDA.abs(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.abs(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "sin":
- LibMatrixCUDA.sin(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.sin(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "cos":
- LibMatrixCUDA.cos(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.cos(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "tan":
- LibMatrixCUDA.tan(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.tan(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "asin":
- LibMatrixCUDA.asin(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.asin(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "acos":
- LibMatrixCUDA.acos(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.acos(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "atan":
- LibMatrixCUDA.atan(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.atan(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
case "sign":
- LibMatrixCUDA.sign(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName()); break;
+ LibMatrixCUDA.sign(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName()); break;
default:
throw new DMLRuntimeException("Unsupported GPU operator:" + opcode);
}
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixArithmeticGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixArithmeticGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixArithmeticGPUInstruction.java
index 9573a1a..a03f9b1 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixArithmeticGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixArithmeticGPUInstruction.java
@@ -71,7 +71,7 @@ public class MatrixMatrixArithmeticGPUInstruction extends ArithmeticBinaryGPUIns
ec.setMetaData(_output.getName(), (int)rlen, (int)clen);
BinaryOperator bop = (BinaryOperator) _optr;
- LibMatrixCUDA.matrixScalarArithmetic(ec, ec.getGPUContext(), getExtendedOpcode(), in1, in2, _output.getName(), isLeftTransposed, isRightTransposed, bop);
+ LibMatrixCUDA.matrixScalarArithmetic(ec, ec.getGPUContext(0), getExtendedOpcode(), in1, in2, _output.getName(), isLeftTransposed, isRightTransposed, bop);
ec.releaseMatrixInputForGPUInstruction(_input1.getName());
ec.releaseMatrixInputForGPUInstruction(_input2.getName());
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixAxpyGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixAxpyGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixAxpyGPUInstruction.java
index 58905d6..e430e29 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixAxpyGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixAxpyGPUInstruction.java
@@ -103,7 +103,7 @@ public class MatrixMatrixAxpyGPUInstruction extends ArithmeticBinaryGPUInstructi
" and input2:" + rlen2 + " X " + clen2);
}
- LibMatrixCUDA.axpy(ec, ec.getGPUContext(), getExtendedOpcode(), in1, in2, _output.getName(), multiplier*scalar.getDoubleValue());
+ LibMatrixCUDA.axpy(ec, ec.getGPUContext(0), getExtendedOpcode(), in1, in2, _output.getName(), multiplier*scalar.getDoubleValue());
ec.releaseMatrixInputForGPUInstruction(_input1.getName());
ec.releaseMatrixInputForGPUInstruction(_input2.getName());
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixBuiltinGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixBuiltinGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixBuiltinGPUInstruction.java
index 8936735..e60a3d7 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixBuiltinGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/MatrixMatrixBuiltinGPUInstruction.java
@@ -46,7 +46,7 @@ public class MatrixMatrixBuiltinGPUInstruction extends BuiltinBinaryGPUInstructi
if(opcode.equals("solve")) {
ec.setMetaData(output.getName(), mat1.getNumColumns(), 1);
- LibMatrixCUDA.solve(ec, ec.getGPUContext(), getExtendedOpcode(), mat1, mat2, output.getName());
+ LibMatrixCUDA.solve(ec, ec.getGPUContext(0), getExtendedOpcode(), mat1, mat2, output.getName());
} else {
throw new DMLRuntimeException("Unsupported GPU operator:" + opcode);
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/ReorgGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/ReorgGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/ReorgGPUInstruction.java
index 53d56a3..bc63d12 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/ReorgGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/ReorgGPUInstruction.java
@@ -81,7 +81,7 @@ public class ReorgGPUInstruction extends GPUInstruction
//execute operation
ec.setMetaData(_output.getName(), rlen, clen);
- LibMatrixCUDA.transpose(ec, ec.getGPUContext(), getExtendedOpcode(), mat, _output.getName());
+ LibMatrixCUDA.transpose(ec, ec.getGPUContext(0), getExtendedOpcode(), mat, _output.getName());
//release inputs/outputs
ec.releaseMatrixInputForGPUInstruction(_input.getName());
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/ScalarMatrixArithmeticGPUInstruction.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/ScalarMatrixArithmeticGPUInstruction.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/ScalarMatrixArithmeticGPUInstruction.java
index 64cb6c4..ea4665a 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/ScalarMatrixArithmeticGPUInstruction.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/ScalarMatrixArithmeticGPUInstruction.java
@@ -60,7 +60,7 @@ public class ScalarMatrixArithmeticGPUInstruction extends ArithmeticBinaryGPUIns
ScalarOperator sc_op = (ScalarOperator) _optr;
sc_op.setConstant(constant.getDoubleValue());
- LibMatrixCUDA.matrixScalarArithmetic(ec, ec.getGPUContext(), getExtendedOpcode(), in1, _output.getName(), isTransposed, sc_op);
+ LibMatrixCUDA.matrixScalarArithmetic(ec, ec.getGPUContext(0), getExtendedOpcode(), in1, _output.getName(), isTransposed, sc_op);
ec.releaseMatrixInputForGPUInstruction(mat.getName());
ec.releaseMatrixOutputForGPUInstruction(_output.getName());
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/CSRPointer.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/CSRPointer.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/CSRPointer.java
index 0ff9d14..b15dd69 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/CSRPointer.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/CSRPointer.java
@@ -52,453 +52,477 @@ import jcuda.jcusparse.cusparsePointerMode;
*/
public class CSRPointer {
- private static final Log LOG = LogFactory.getLog(CSRPointer.class.getName());
-
- private static final double ULTRA_SPARSITY_TURN_POINT = 0.0004;
-
- /** {@link GPUContext} instance to track the GPU to do work on */
- private final GPUContext gpuContext;
-
- public static cusparseMatDescr matrixDescriptor;
-
- /** Number of non zeroes */
- public long nnz;
-
- /** double array of non zero values */
- public Pointer val;
-
- /** integer array of start of all rows and end of last row + 1 */
- public Pointer rowPtr;
-
- /** integer array of nnz values' column indices */
- public Pointer colInd;
-
- /** descriptor of matrix, only CUSPARSE_MATRIX_TYPE_GENERAL supported */
- public cusparseMatDescr descr;
-
-
- public CSRPointer clone(int rows) throws DMLRuntimeException {
- CSRPointer me = this;
- CSRPointer that = new CSRPointer(me.getGPUContext());
-
- that.allocateMatDescrPointer();
- long totalSize = estimateSize(me.nnz, rows);
- that.gpuContext.ensureFreeSpace(totalSize);
-
- that.nnz = me.nnz;
- that.val = allocate(that.nnz * Sizeof.DOUBLE);
- that.rowPtr = allocate(rows * Sizeof.DOUBLE);
- that.colInd = allocate(that.nnz * Sizeof.DOUBLE);
-
- cudaMemcpy(that.val, me.val, that.nnz * Sizeof.DOUBLE, cudaMemcpyDeviceToDevice);
- cudaMemcpy(that.rowPtr, me.rowPtr, rows * Sizeof.DOUBLE, cudaMemcpyDeviceToDevice);
- cudaMemcpy(that.colInd, me.colInd, that.nnz * Sizeof.DOUBLE, cudaMemcpyDeviceToDevice);
-
- return that;
- }
-
- /**
- * Default constructor to help with Factory method {@link #allocateEmpty(GPUContext, long, long)}
- * @param gCtx a valid {@link GPUContext}
- */
- private CSRPointer(GPUContext gCtx) {
- gpuContext = gCtx;
- val = new Pointer();
- rowPtr = new Pointer();
- colInd = new Pointer();
- allocateMatDescrPointer();
- }
-
-// private Pointer allocate(String instName, long size) throws DMLRuntimeException {
-// return getGPUContext().allocate(instName, size);
-// }
-
- private Pointer allocate(long size) throws DMLRuntimeException {
- return getGPUContext().allocate(size);
- }
-
-// private void cudaFreeHelper(Pointer toFree) throws DMLRuntimeException {
-// getGPUContext().cudaFreeHelper(toFree);
-// }
-
- private void cudaFreeHelper(Pointer toFree, boolean eager) throws DMLRuntimeException {
- getGPUContext().cudaFreeHelper(toFree, eager);
- }
-
-// private void cudaFreeHelper(String instName, Pointer toFree, boolean eager) throws DMLRuntimeException {
-// getGPUContext().cudaFreeHelper(instName, toFree, eager);
-// }
-
- private static long getDoubleSizeOf(long numElems) {
- return numElems * ((long)jcuda.Sizeof.DOUBLE);
- }
-
- private static long getIntSizeOf(long numElems) {
- return numElems * ((long)jcuda.Sizeof.INT);
- }
-
- private GPUContext getGPUContext() {
- return gpuContext;
- }
-
- public static int toIntExact(long l) throws DMLRuntimeException {
- if (l < Integer.MIN_VALUE || l > Integer.MAX_VALUE) {
- throw new DMLRuntimeException("Cannot be cast to int:" + l);
- }
- return (int) l;
- }
-
- /**
- * @return Singleton default matrix descriptor object
- * (set with CUSPARSE_MATRIX_TYPE_GENERAL, CUSPARSE_INDEX_BASE_ZERO)
- */
- public static cusparseMatDescr getDefaultCuSparseMatrixDescriptor() {
- if (matrixDescriptor == null) {
- // Code from JCuda Samples - http://www.jcuda.org/samples/JCusparseSample.java
- matrixDescriptor = new cusparseMatDescr();
- cusparseCreateMatDescr(matrixDescriptor);
- cusparseSetMatType(matrixDescriptor, CUSPARSE_MATRIX_TYPE_GENERAL);
- cusparseSetMatIndexBase(matrixDescriptor, CUSPARSE_INDEX_BASE_ZERO);
- }
- return matrixDescriptor;
- }
-
- /**
- * Estimate the size of a CSR matrix in GPU memory
- * Size of pointers is not needed and is not added in
- *
- * @param nnz2 number of non zeroes
- * @param rows number of rows
- * @return size estimate
- */
- public static long estimateSize(long nnz2, long rows) {
- long sizeofValArray = getDoubleSizeOf(nnz2);
- long sizeofRowPtrArray = getIntSizeOf(rows + 1);
- long sizeofColIndArray = getIntSizeOf(nnz2);
- long sizeofDescr = getIntSizeOf(4);
- // From the CUSPARSE documentation, the cusparseMatDescr in native code is represented as:
- // typedef struct {
- // cusparseMatrixType_t MatrixType;
- // cusparseFillMode_t FillMode;
- // cusparseDiagType_t DiagType;
- // cusparseIndexBase_t IndexBase;
- // } cusparseMatDescr_t;
- long tot = sizeofValArray + sizeofRowPtrArray + sizeofColIndArray + sizeofDescr;
- return tot;
- }
-
- /**
- * Static method to copy a CSR sparse matrix from Host to Device
- *
- * @param dest [input] destination location (on GPU)
- * @param rows number of rows
- * @param nnz number of non-zeroes
- * @param rowPtr integer array of row pointers
- * @param colInd integer array of column indices
- * @param values double array of non zero values
- */
- public static void copyToDevice(CSRPointer dest, int rows, long nnz, int[] rowPtr, int[] colInd, double[] values) {
- CSRPointer r = dest;
- long t0 = 0;
- if (DMLScript.STATISTICS) t0 = System.nanoTime();
- r.nnz = nnz;
- cudaMemcpy(r.rowPtr, Pointer.to(rowPtr), getIntSizeOf(rows + 1), cudaMemcpyHostToDevice);
- cudaMemcpy(r.colInd, Pointer.to(colInd), getIntSizeOf(nnz), cudaMemcpyHostToDevice);
- cudaMemcpy(r.val, Pointer.to(values), getDoubleSizeOf(nnz), cudaMemcpyHostToDevice);
- if (DMLScript.STATISTICS) GPUStatistics.cudaToDevTime.addAndGet(System.nanoTime() - t0);
- if (DMLScript.STATISTICS) GPUStatistics.cudaToDevCount.addAndGet(3);
- }
-
- /**
- * Static method to copy a CSR sparse matrix from Device to host
- *
- * @param src [input] source location (on GPU)
- * @param rows [input] number of rows
- * @param nnz [input] number of non-zeroes
- * @param rowPtr [output] pre-allocated integer array of row pointers of size (rows+1)
- * @param colInd [output] pre-allocated integer array of column indices of size nnz
- * @param values [output] pre-allocated double array of values of size nnz
- */
- public static void copyToHost(CSRPointer src, int rows, long nnz, int[] rowPtr, int[] colInd, double[] values) {
- CSRPointer r = src;
- long t0 = 0;
- if (DMLScript.STATISTICS) t0 = System.nanoTime();
- cudaMemcpy(Pointer.to(rowPtr), r.rowPtr, getIntSizeOf(rows + 1), cudaMemcpyDeviceToHost);
- cudaMemcpy(Pointer.to(colInd), r.colInd, getIntSizeOf(nnz), cudaMemcpyDeviceToHost);
- cudaMemcpy(Pointer.to(values), r.val, getDoubleSizeOf(nnz), cudaMemcpyDeviceToHost);
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - t0);
- if (DMLScript.STATISTICS) GPUStatistics.cudaFromDevCount.addAndGet(3);
- }
-
- /**
- * Estimates the number of non zero elements from the results of a sparse cusparseDgeam operation
- * C = a op(A) + b op(B)
- * @param gCtx a valid {@link GPUContext}
- * @param handle a valid {@link cusparseHandle}
- * @param A Sparse Matrix A on GPU
- * @param B Sparse Matrix B on GPU
- * @param m Rows in A
- * @param n Columns in Bs
- * @return CSR (compressed sparse row) pointer
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- public static CSRPointer allocateForDgeam(GPUContext gCtx, cusparseHandle handle, CSRPointer A, CSRPointer B, int m, int n)
- throws DMLRuntimeException {
- if (A.nnz >= Integer.MAX_VALUE || B.nnz >= Integer.MAX_VALUE) {
- throw new DMLRuntimeException("Number of non zeroes is larger than supported by cuSparse");
- }
- CSRPointer C = new CSRPointer(gCtx);
- step1AllocateRowPointers(gCtx, handle, C, m);
- step2GatherNNZGeam(gCtx, handle, A, B, C, m, n);
- step3AllocateValNInd(gCtx, handle, C);
- return C;
- }
-
- /**
- * Estimates the number of non-zero elements from the result of a sparse matrix multiplication C = A * B
- * and returns the {@link CSRPointer} to C with the appropriate GPU memory.
- *
- * @param gCtx ?
- * @param handle a valid {@link cusparseHandle}
- * @param A Sparse Matrix A on GPU
- * @param transA 'T' if A is to be transposed, 'N' otherwise
- * @param B Sparse Matrix B on GPU
- * @param transB 'T' if B is to be transposed, 'N' otherwise
- * @param m Rows in A
- * @param n Columns in B
- * @param k Columns in A / Rows in B
- * @return a {@link CSRPointer} instance that encapsulates the CSR matrix on GPU
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- public static CSRPointer allocateForMatrixMultiply(GPUContext gCtx, cusparseHandle handle, CSRPointer A, int transA, CSRPointer B, int transB, int m, int n, int k)
- throws DMLRuntimeException {
- // Following the code example at http://docs.nvidia.com/cuda/cusparse/#cusparse-lt-t-gt-csrgemm and at
- // https://github.com/jcuda/jcuda-matrix-utils/blob/master/JCudaMatrixUtils/src/test/java/org/jcuda/matrix/samples/JCusparseSampleDgemm.java
- CSRPointer C = new CSRPointer(gCtx);
- step1AllocateRowPointers(gCtx, handle, C, m);
- step2GatherNNZGemm(gCtx, handle, A, transA, B, transB, C, m, n, k);
- step3AllocateValNInd(gCtx, handle, C);
- return C;
- }
-
- /**
- * Check for ultra sparsity
- *
- * @param rows number of rows
- * @param cols number of columns
- * @return true if ultra sparse
- */
- public boolean isUltraSparse(int rows, int cols) {
- double sp = ((double) nnz / rows / cols);
- return sp < ULTRA_SPARSITY_TURN_POINT;
- }
-
-// ==============================================================================================
-
-// The following methods estimate the memory needed for sparse matrices that are
-// results of operations on other sparse matrices using the cuSparse Library.
-// The operation is C = op(A) binaryOperation op(B), C is the output and A & B are the inputs
-// op = whether to transpose or not
-// binaryOperation = For cuSparse, +, - are *(matmul) are supported
-
-// From CuSparse Manual,
-// Since A and B have different sparsity patterns, cuSPARSE adopts a two-step approach
-// to complete sparse matrix C. In the first step, the user allocates csrRowPtrC of m+1
-// elements and uses function cusparseXcsrgeamNnz() to determine csrRowPtrC
-// and the total number of nonzero elements. In the second step, the user gathers nnzC
-//(number of nonzero elements of matrix C) from either (nnzC=*nnzTotalDevHostPtr)
-// or (nnzC=csrRowPtrC(m)-csrRowPtrC(0)) and allocates csrValC, csrColIndC of
-// nnzC elements respectively, then finally calls function cusparse[S|D|C|Z]csrgeam()
-// to complete matrix C.
-
- /**
- * Initializes {@link #descr} to CUSPARSE_MATRIX_TYPE_GENERAL,
- * the default that works for DGEMM.
- */
- private void allocateMatDescrPointer() {
- this.descr = getDefaultCuSparseMatrixDescriptor();
- }
-
- /**
- * Factory method to allocate an empty CSR Sparse matrix on the GPU
- *
- * @param gCtx ?
- * @param nnz2 number of non-zeroes
- * @param rows number of rows
- * @return a {@link CSRPointer} instance that encapsulates the CSR matrix on GPU
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- public static CSRPointer allocateEmpty(GPUContext gCtx, long nnz2, long rows) throws DMLRuntimeException {
- LOG.trace("GPU : allocateEmpty from CSRPointer with nnz=" + nnz2 + " and rows=" + rows + ", GPUContext=" + gCtx);
- assert nnz2 > -1 : "Incorrect usage of internal API, number of non zeroes is less than 0 when trying to allocate sparse data on GPU";
- CSRPointer r = new CSRPointer(gCtx);
- r.nnz = nnz2;
- if (nnz2 == 0) {
- // The convention for an empty sparse matrix is to just have an instance of the CSRPointer object
- // with no memory allocated on the GPU.
- return r;
- }
- gCtx.ensureFreeSpace(getDoubleSizeOf(nnz2) + getIntSizeOf(rows + 1) + getIntSizeOf(nnz2));
- // increment the cudaCount by 1 for the allocation of all 3 arrays
- r.val = gCtx.allocate(null, getDoubleSizeOf(nnz2));
- r.rowPtr = gCtx.allocate(null, getIntSizeOf(rows + 1));
- r.colInd = gCtx.allocate(null, getIntSizeOf(nnz2));
- return r;
- }
-
- /**
- * Allocate row pointers of m+1 elements
- * @param gCtx a valid {@link GPUContext}
- * @param handle a valid {@link cusparseHandle}
- * @param C Output matrix
- * @param rowsC number of rows in C
- * @throws DMLRuntimeException ?
- */
- private static void step1AllocateRowPointers(GPUContext gCtx, cusparseHandle handle, CSRPointer C, int rowsC) throws DMLRuntimeException {
- LOG.trace("GPU : step1AllocateRowPointers" + ", GPUContext=" + gCtx);
- cusparseSetPointerMode(handle, cusparsePointerMode.CUSPARSE_POINTER_MODE_HOST);
- //cudaDeviceSynchronize;
- // Do not increment the cudaCount of allocations on GPU
- C.rowPtr = gCtx.allocate(getIntSizeOf((long) rowsC + 1));
- }
-
- /**
- * Determine total number of nonzero element for the cusparseDgeam operation.
- * This is done from either (nnzC=*nnzTotalDevHostPtr) or (nnzC=csrRowPtrC(m)-csrRowPtrC(0))
- * @param gCtx a valid {@link GPUContext}
- * @param handle a valid {@link cusparseHandle}
- * @param A Sparse Matrix A on GPU
- * @param B Sparse Matrix B on GPU
- * @param C Output Sparse Matrix C on GPU
- * @param m Rows in C
- * @param n Columns in C
- * @throws DMLRuntimeException ?
- */
- private static void step2GatherNNZGeam(GPUContext gCtx, cusparseHandle handle, CSRPointer A, CSRPointer B, CSRPointer C, int m, int n) throws DMLRuntimeException {
- LOG.trace("GPU : step2GatherNNZGeam for DGEAM" + ", GPUContext=" + gCtx);
- int[] CnnzArray = {-1};
- cusparseXcsrgeamNnz(handle, m, n,
- A.descr, toIntExact(A.nnz), A.rowPtr, A.colInd,
- B.descr, toIntExact(B.nnz), B.rowPtr, B.colInd,
- C.descr, C.rowPtr, Pointer.to(CnnzArray));
- //cudaDeviceSynchronize;
- if (CnnzArray[0] != -1) {
- C.nnz = CnnzArray[0];
- } else {
- int baseArray[] = {0};
- cudaMemcpy(Pointer.to(CnnzArray), C.rowPtr.withByteOffset(getIntSizeOf(m)), getIntSizeOf(1), cudaMemcpyDeviceToHost);
- cudaMemcpy(Pointer.to(baseArray), C.rowPtr, getIntSizeOf(1), cudaMemcpyDeviceToHost);
- C.nnz = CnnzArray[0] - baseArray[0];
- }
- }
-
-// ==============================================================================================
-
- /**
- * Determine total number of nonzero element for the cusparseDgemm operation.
- * @param gCtx a valid {@link GPUContext}
- * @param handle a valid {@link cusparseHandle}
- * @param A Sparse Matrix A on GPU
- * @param transA op - whether A is transposed
- * @param B Sparse Matrix B on GPU
- * @param transB op - whether B is transposed
- * @param C Output Sparse Matrix C on GPU
- * @param m Number of rows of sparse matrix op ( A ) and C
- * @param n Number of columns of sparse matrix op ( B ) and C
- * @param k Number of columns/rows of sparse matrix op ( A ) / op ( B )
- * @throws DMLRuntimeException ?
- */
- private static void step2GatherNNZGemm(GPUContext gCtx, cusparseHandle handle, CSRPointer A, int transA, CSRPointer B, int transB, CSRPointer C, int m, int n, int k) throws DMLRuntimeException {
- LOG.trace("GPU : step2GatherNNZGemm for DGEMM" + ", GPUContext=" + gCtx);
- int[] CnnzArray = {-1};
- if (A.nnz >= Integer.MAX_VALUE || B.nnz >= Integer.MAX_VALUE) {
- throw new DMLRuntimeException("Number of non zeroes is larger than supported by cuSparse");
- }
- cusparseXcsrgemmNnz(handle, transA, transB, m, n, k,
- A.descr, toIntExact(A.nnz), A.rowPtr, A.colInd,
- B.descr, toIntExact(B.nnz), B.rowPtr, B.colInd,
- C.descr, C.rowPtr, Pointer.to(CnnzArray));
- //cudaDeviceSynchronize;
- if (CnnzArray[0] != -1) {
- C.nnz = CnnzArray[0];
- } else {
- int baseArray[] = {0};
- cudaMemcpy(Pointer.to(CnnzArray), C.rowPtr.withByteOffset(getIntSizeOf(m)), getIntSizeOf(1), cudaMemcpyDeviceToHost);
- cudaMemcpy(Pointer.to(baseArray), C.rowPtr, getIntSizeOf(1), cudaMemcpyDeviceToHost);
- C.nnz = CnnzArray[0] - baseArray[0];
- }
- }
-
- /**
- * Allocate val and index pointers.
- * @param gCtx a valid {@link GPUContext}
- * @param handle a valid {@link cusparseHandle}
- * @param C Output sparse matrix on GPU
- * @throws DMLRuntimeException ?
- */
- private static void step3AllocateValNInd(GPUContext gCtx, cusparseHandle handle, CSRPointer C) throws DMLRuntimeException {
- LOG.trace("GPU : step3AllocateValNInd" + ", GPUContext=" + gCtx);
- // Increment cudaCount by one when all three arrays of CSR sparse array are allocated
- C.val = gCtx.allocate(null, getDoubleSizeOf(C.nnz));
- C.colInd = gCtx.allocate(null, getIntSizeOf(C.nnz));
- }
-
- /**
- * Copies this CSR matrix on the GPU to a dense column-major matrix
- * on the GPU. This is a temporary matrix for operations such as
- * cusparseDcsrmv.
- * Since the allocated matrix is temporary, bookkeeping is not updated.
- * The caller is responsible for calling "free" on the returned Pointer object
- *
- * @param cusparseHandle a valid {@link cusparseHandle}
- * @param cublasHandle a valid {@link cublasHandle}
- * @param rows number of rows in this CSR matrix
- * @param cols number of columns in this CSR matrix
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- * @return A {@link Pointer} to the allocated dense matrix (in column-major format)
- */
- public Pointer toColumnMajorDenseMatrix(cusparseHandle cusparseHandle, cublasHandle cublasHandle, int rows, int cols) throws DMLRuntimeException {
- LOG.trace("GPU : sparse -> column major dense (inside CSRPointer) on " + this + ", GPUContext=" + getGPUContext());
- long size = ((long) rows) * getDoubleSizeOf((long) cols);
- Pointer A = allocate(size);
- // If this sparse block is empty, the allocated dense matrix, initialized to zeroes, will be returned.
- if (val != null && rowPtr != null && colInd != null && nnz > 0) {
- // Note: cusparseDcsr2dense method cannot handle empty blocks
- cusparseDcsr2dense(cusparseHandle, rows, cols, descr, val, rowPtr, colInd, A, rows);
- //cudaDeviceSynchronize;
- } else {
- LOG.debug("in CSRPointer, the values array, row pointers array or column indices array was null");
- }
- return A;
- }
-
- /**
- * Calls cudaFree lazily on the allocated {@link Pointer} instances
- *
- * @throws DMLRuntimeException ?
- */
- public void deallocate() throws DMLRuntimeException {
- deallocate(false);
- }
-
- /**
- * Calls cudaFree lazily or eagerly on the allocated {@link Pointer} instances
- *
- * @param eager whether to do eager or lazy cudaFrees
- * @throws DMLRuntimeException ?
- */
- public void deallocate(boolean eager) throws DMLRuntimeException {
- if (nnz > 0) {
- cudaFreeHelper(val, eager);
- cudaFreeHelper(rowPtr, eager);
- cudaFreeHelper(colInd, eager);
- }
- }
-
- @Override
- public String toString() {
- return "CSRPointer{" +
- "nnz=" + nnz +
- '}';
- }
+ private static final Log LOG = LogFactory.getLog(CSRPointer.class.getName());
+
+ private static final double ULTRA_SPARSITY_TURN_POINT = 0.0004;
+ public static cusparseMatDescr matrixDescriptor;
+ /**
+ * {@link GPUContext} instance to track the GPU to do work on
+ */
+ private final GPUContext gpuContext;
+ /**
+ * Number of non zeroes
+ */
+ public long nnz;
+
+ /**
+ * double array of non zero values
+ */
+ public Pointer val;
+
+ /**
+ * integer array of start of all rows and end of last row + 1
+ */
+ public Pointer rowPtr;
+
+ /**
+ * integer array of nnz values' column indices
+ */
+ public Pointer colInd;
+
+ /**
+ * descriptor of matrix, only CUSPARSE_MATRIX_TYPE_GENERAL supported
+ */
+ public cusparseMatDescr descr;
+
+ /**
+ * Default constructor to help with Factory method {@link #allocateEmpty(GPUContext, long, long)}
+ *
+ * @param gCtx a valid {@link GPUContext}
+ */
+ private CSRPointer(GPUContext gCtx) {
+ gpuContext = gCtx;
+ val = new Pointer();
+ rowPtr = new Pointer();
+ colInd = new Pointer();
+ allocateMatDescrPointer();
+ }
+
+ private static long getDoubleSizeOf(long numElems) {
+ return numElems * ((long) jcuda.Sizeof.DOUBLE);
+ }
+
+ // private Pointer allocate(String instName, long size) throws DMLRuntimeException {
+ // return getGPUContext().allocate(instName, size);
+ // }
+
+ private static long getIntSizeOf(long numElems) {
+ return numElems * ((long) jcuda.Sizeof.INT);
+ }
+
+ // private void cudaFreeHelper(Pointer toFree) throws DMLRuntimeException {
+ // getGPUContext().cudaFreeHelper(toFree);
+ // }
+
+ public static int toIntExact(long l) throws DMLRuntimeException {
+ if (l < Integer.MIN_VALUE || l > Integer.MAX_VALUE) {
+ throw new DMLRuntimeException("Cannot be cast to int:" + l);
+ }
+ return (int) l;
+ }
+
+ // private void cudaFreeHelper(String instName, Pointer toFree, boolean eager) throws DMLRuntimeException {
+ // getGPUContext().cudaFreeHelper(instName, toFree, eager);
+ // }
+
+ /**
+ * @return Singleton default matrix descriptor object
+ * (set with CUSPARSE_MATRIX_TYPE_GENERAL, CUSPARSE_INDEX_BASE_ZERO)
+ */
+ public static cusparseMatDescr getDefaultCuSparseMatrixDescriptor() {
+ if (matrixDescriptor == null) {
+ // Code from JCuda Samples - http://www.jcuda.org/samples/JCusparseSample.java
+ matrixDescriptor = new cusparseMatDescr();
+ cusparseCreateMatDescr(matrixDescriptor);
+ cusparseSetMatType(matrixDescriptor, CUSPARSE_MATRIX_TYPE_GENERAL);
+ cusparseSetMatIndexBase(matrixDescriptor, CUSPARSE_INDEX_BASE_ZERO);
+ }
+ return matrixDescriptor;
+ }
+
+ /**
+ * Estimate the size of a CSR matrix in GPU memory
+ * Size of pointers is not needed and is not added in
+ *
+ * @param nnz2 number of non zeroes
+ * @param rows number of rows
+ * @return size estimate
+ */
+ public static long estimateSize(long nnz2, long rows) {
+ long sizeofValArray = getDoubleSizeOf(nnz2);
+ long sizeofRowPtrArray = getIntSizeOf(rows + 1);
+ long sizeofColIndArray = getIntSizeOf(nnz2);
+ long sizeofDescr = getIntSizeOf(4);
+ // From the CUSPARSE documentation, the cusparseMatDescr in native code is represented as:
+ // typedef struct {
+ // cusparseMatrixType_t MatrixType;
+ // cusparseFillMode_t FillMode;
+ // cusparseDiagType_t DiagType;
+ // cusparseIndexBase_t IndexBase;
+ // } cusparseMatDescr_t;
+ long tot = sizeofValArray + sizeofRowPtrArray + sizeofColIndArray + sizeofDescr;
+ return tot;
+ }
+
+ /**
+ * Static method to copy a CSR sparse matrix from Host to Device
+ *
+ * @param dest [input] destination location (on GPU)
+ * @param rows number of rows
+ * @param nnz number of non-zeroes
+ * @param rowPtr integer array of row pointers
+ * @param colInd integer array of column indices
+ * @param values double array of non zero values
+ */
+ public static void copyToDevice(CSRPointer dest, int rows, long nnz, int[] rowPtr, int[] colInd, double[] values) {
+ CSRPointer r = dest;
+ long t0 = 0;
+ if (DMLScript.STATISTICS)
+ t0 = System.nanoTime();
+ r.nnz = nnz;
+ cudaMemcpy(r.rowPtr, Pointer.to(rowPtr), getIntSizeOf(rows + 1), cudaMemcpyHostToDevice);
+ cudaMemcpy(r.colInd, Pointer.to(colInd), getIntSizeOf(nnz), cudaMemcpyHostToDevice);
+ cudaMemcpy(r.val, Pointer.to(values), getDoubleSizeOf(nnz), cudaMemcpyHostToDevice);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaToDevTime.addAndGet(System.nanoTime() - t0);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaToDevCount.addAndGet(3);
+ }
+
+ /**
+ * Static method to copy a CSR sparse matrix from Device to host
+ *
+ * @param src [input] source location (on GPU)
+ * @param rows [input] number of rows
+ * @param nnz [input] number of non-zeroes
+ * @param rowPtr [output] pre-allocated integer array of row pointers of size (rows+1)
+ * @param colInd [output] pre-allocated integer array of column indices of size nnz
+ * @param values [output] pre-allocated double array of values of size nnz
+ */
+ public static void copyToHost(CSRPointer src, int rows, long nnz, int[] rowPtr, int[] colInd, double[] values) {
+ CSRPointer r = src;
+ long t0 = 0;
+ if (DMLScript.STATISTICS)
+ t0 = System.nanoTime();
+ cudaMemcpy(Pointer.to(rowPtr), r.rowPtr, getIntSizeOf(rows + 1), cudaMemcpyDeviceToHost);
+ cudaMemcpy(Pointer.to(colInd), r.colInd, getIntSizeOf(nnz), cudaMemcpyDeviceToHost);
+ cudaMemcpy(Pointer.to(values), r.val, getDoubleSizeOf(nnz), cudaMemcpyDeviceToHost);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevTime.addAndGet(System.nanoTime() - t0);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaFromDevCount.addAndGet(3);
+ }
+
+ /**
+ * Estimates the number of non zero elements from the results of a sparse cusparseDgeam operation
+ * C = a op(A) + b op(B)
+ *
+ * @param gCtx a valid {@link GPUContext}
+ * @param handle a valid {@link cusparseHandle}
+ * @param A Sparse Matrix A on GPU
+ * @param B Sparse Matrix B on GPU
+ * @param m Rows in A
+ * @param n Columns in Bs
+ * @return CSR (compressed sparse row) pointer
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public static CSRPointer allocateForDgeam(GPUContext gCtx, cusparseHandle handle, CSRPointer A, CSRPointer B, int m,
+ int n) throws DMLRuntimeException {
+ if (A.nnz >= Integer.MAX_VALUE || B.nnz >= Integer.MAX_VALUE) {
+ throw new DMLRuntimeException("Number of non zeroes is larger than supported by cuSparse");
+ }
+ CSRPointer C = new CSRPointer(gCtx);
+ step1AllocateRowPointers(gCtx, handle, C, m);
+ step2GatherNNZGeam(gCtx, handle, A, B, C, m, n);
+ step3AllocateValNInd(gCtx, handle, C);
+ return C;
+ }
+
+ /**
+ * Estimates the number of non-zero elements from the result of a sparse matrix multiplication C = A * B
+ * and returns the {@link CSRPointer} to C with the appropriate GPU memory.
+ *
+ * @param gCtx ?
+ * @param handle a valid {@link cusparseHandle}
+ * @param A Sparse Matrix A on GPU
+ * @param transA 'T' if A is to be transposed, 'N' otherwise
+ * @param B Sparse Matrix B on GPU
+ * @param transB 'T' if B is to be transposed, 'N' otherwise
+ * @param m Rows in A
+ * @param n Columns in B
+ * @param k Columns in A / Rows in B
+ * @return a {@link CSRPointer} instance that encapsulates the CSR matrix on GPU
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public static CSRPointer allocateForMatrixMultiply(GPUContext gCtx, cusparseHandle handle, CSRPointer A, int transA,
+ CSRPointer B, int transB, int m, int n, int k) throws DMLRuntimeException {
+ // Following the code example at http://docs.nvidia.com/cuda/cusparse/#cusparse-lt-t-gt-csrgemm and at
+ // https://github.com/jcuda/jcuda-matrix-utils/blob/master/JCudaMatrixUtils/src/test/java/org/jcuda/matrix/samples/JCusparseSampleDgemm.java
+ CSRPointer C = new CSRPointer(gCtx);
+ step1AllocateRowPointers(gCtx, handle, C, m);
+ step2GatherNNZGemm(gCtx, handle, A, transA, B, transB, C, m, n, k);
+ step3AllocateValNInd(gCtx, handle, C);
+ return C;
+ }
+
+ /**
+ * Factory method to allocate an empty CSR Sparse matrix on the GPU
+ *
+ * @param gCtx ?
+ * @param nnz2 number of non-zeroes
+ * @param rows number of rows
+ * @return a {@link CSRPointer} instance that encapsulates the CSR matrix on GPU
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public static CSRPointer allocateEmpty(GPUContext gCtx, long nnz2, long rows) throws DMLRuntimeException {
+ LOG.trace(
+ "GPU : allocateEmpty from CSRPointer with nnz=" + nnz2 + " and rows=" + rows + ", GPUContext=" + gCtx);
+ assert nnz2
+ > -1 : "Incorrect usage of internal API, number of non zeroes is less than 0 when trying to allocate sparse data on GPU";
+ CSRPointer r = new CSRPointer(gCtx);
+ r.nnz = nnz2;
+ if (nnz2 == 0) {
+ // The convention for an empty sparse matrix is to just have an instance of the CSRPointer object
+ // with no memory allocated on the GPU.
+ return r;
+ }
+ gCtx.ensureFreeSpace(getDoubleSizeOf(nnz2) + getIntSizeOf(rows + 1) + getIntSizeOf(nnz2));
+ // increment the cudaCount by 1 for the allocation of all 3 arrays
+ r.val = gCtx.allocate(null, getDoubleSizeOf(nnz2));
+ r.rowPtr = gCtx.allocate(null, getIntSizeOf(rows + 1));
+ r.colInd = gCtx.allocate(null, getIntSizeOf(nnz2));
+ return r;
+ }
+
+ /**
+ * Allocate row pointers of m+1 elements
+ *
+ * @param gCtx a valid {@link GPUContext}
+ * @param handle a valid {@link cusparseHandle}
+ * @param C Output matrix
+ * @param rowsC number of rows in C
+ * @throws DMLRuntimeException ?
+ */
+ private static void step1AllocateRowPointers(GPUContext gCtx, cusparseHandle handle, CSRPointer C, int rowsC)
+ throws DMLRuntimeException {
+ LOG.trace("GPU : step1AllocateRowPointers" + ", GPUContext=" + gCtx);
+ cusparseSetPointerMode(handle, cusparsePointerMode.CUSPARSE_POINTER_MODE_HOST);
+ //cudaDeviceSynchronize;
+ // Do not increment the cudaCount of allocations on GPU
+ C.rowPtr = gCtx.allocate(getIntSizeOf((long) rowsC + 1));
+ }
+
+ /**
+ * Determine total number of nonzero element for the cusparseDgeam operation.
+ * This is done from either (nnzC=*nnzTotalDevHostPtr) or (nnzC=csrRowPtrC(m)-csrRowPtrC(0))
+ *
+ * @param gCtx a valid {@link GPUContext}
+ * @param handle a valid {@link cusparseHandle}
+ * @param A Sparse Matrix A on GPU
+ * @param B Sparse Matrix B on GPU
+ * @param C Output Sparse Matrix C on GPU
+ * @param m Rows in C
+ * @param n Columns in C
+ * @throws DMLRuntimeException ?
+ */
+ private static void step2GatherNNZGeam(GPUContext gCtx, cusparseHandle handle, CSRPointer A, CSRPointer B,
+ CSRPointer C, int m, int n) throws DMLRuntimeException {
+ LOG.trace("GPU : step2GatherNNZGeam for DGEAM" + ", GPUContext=" + gCtx);
+ int[] CnnzArray = { -1 };
+ cusparseXcsrgeamNnz(handle, m, n, A.descr, toIntExact(A.nnz), A.rowPtr, A.colInd, B.descr, toIntExact(B.nnz),
+ B.rowPtr, B.colInd, C.descr, C.rowPtr, Pointer.to(CnnzArray));
+ //cudaDeviceSynchronize;
+ if (CnnzArray[0] != -1) {
+ C.nnz = CnnzArray[0];
+ } else {
+ int baseArray[] = { 0 };
+ cudaMemcpy(Pointer.to(CnnzArray), C.rowPtr.withByteOffset(getIntSizeOf(m)), getIntSizeOf(1),
+ cudaMemcpyDeviceToHost);
+ cudaMemcpy(Pointer.to(baseArray), C.rowPtr, getIntSizeOf(1), cudaMemcpyDeviceToHost);
+ C.nnz = CnnzArray[0] - baseArray[0];
+ }
+ }
+
+ /**
+ * Determine total number of nonzero element for the cusparseDgemm operation.
+ *
+ * @param gCtx a valid {@link GPUContext}
+ * @param handle a valid {@link cusparseHandle}
+ * @param A Sparse Matrix A on GPU
+ * @param transA op - whether A is transposed
+ * @param B Sparse Matrix B on GPU
+ * @param transB op - whether B is transposed
+ * @param C Output Sparse Matrix C on GPU
+ * @param m Number of rows of sparse matrix op ( A ) and C
+ * @param n Number of columns of sparse matrix op ( B ) and C
+ * @param k Number of columns/rows of sparse matrix op ( A ) / op ( B )
+ * @throws DMLRuntimeException ?
+ */
+ private static void step2GatherNNZGemm(GPUContext gCtx, cusparseHandle handle, CSRPointer A, int transA,
+ CSRPointer B, int transB, CSRPointer C, int m, int n, int k) throws DMLRuntimeException {
+ LOG.trace("GPU : step2GatherNNZGemm for DGEMM" + ", GPUContext=" + gCtx);
+ int[] CnnzArray = { -1 };
+ if (A.nnz >= Integer.MAX_VALUE || B.nnz >= Integer.MAX_VALUE) {
+ throw new DMLRuntimeException("Number of non zeroes is larger than supported by cuSparse");
+ }
+ cusparseXcsrgemmNnz(handle, transA, transB, m, n, k, A.descr, toIntExact(A.nnz), A.rowPtr, A.colInd, B.descr,
+ toIntExact(B.nnz), B.rowPtr, B.colInd, C.descr, C.rowPtr, Pointer.to(CnnzArray));
+ //cudaDeviceSynchronize;
+ if (CnnzArray[0] != -1) {
+ C.nnz = CnnzArray[0];
+ } else {
+ int baseArray[] = { 0 };
+ cudaMemcpy(Pointer.to(CnnzArray), C.rowPtr.withByteOffset(getIntSizeOf(m)), getIntSizeOf(1),
+ cudaMemcpyDeviceToHost);
+ cudaMemcpy(Pointer.to(baseArray), C.rowPtr, getIntSizeOf(1), cudaMemcpyDeviceToHost);
+ C.nnz = CnnzArray[0] - baseArray[0];
+ }
+ }
+
+ /**
+ * Allocate val and index pointers.
+ *
+ * @param gCtx a valid {@link GPUContext}
+ * @param handle a valid {@link cusparseHandle}
+ * @param C Output sparse matrix on GPU
+ * @throws DMLRuntimeException ?
+ */
+ private static void step3AllocateValNInd(GPUContext gCtx, cusparseHandle handle, CSRPointer C)
+ throws DMLRuntimeException {
+ LOG.trace("GPU : step3AllocateValNInd" + ", GPUContext=" + gCtx);
+ // Increment cudaCount by one when all three arrays of CSR sparse array are allocated
+ C.val = gCtx.allocate(null, getDoubleSizeOf(C.nnz));
+ C.colInd = gCtx.allocate(null, getIntSizeOf(C.nnz));
+ }
+
+ // ==============================================================================================
+
+ // The following methods estimate the memory needed for sparse matrices that are
+ // results of operations on other sparse matrices using the cuSparse Library.
+ // The operation is C = op(A) binaryOperation op(B), C is the output and A & B are the inputs
+ // op = whether to transpose or not
+ // binaryOperation = For cuSparse, +, - are *(matmul) are supported
+
+ // From CuSparse Manual,
+ // Since A and B have different sparsity patterns, cuSPARSE adopts a two-step approach
+ // to complete sparse matrix C. In the first step, the user allocates csrRowPtrC of m+1
+ // elements and uses function cusparseXcsrgeamNnz() to determine csrRowPtrC
+ // and the total number of nonzero elements. In the second step, the user gathers nnzC
+ //(number of nonzero elements of matrix C) from either (nnzC=*nnzTotalDevHostPtr)
+ // or (nnzC=csrRowPtrC(m)-csrRowPtrC(0)) and allocates csrValC, csrColIndC of
+ // nnzC elements respectively, then finally calls function cusparse[S|D|C|Z]csrgeam()
+ // to complete matrix C.
+
+ public CSRPointer clone(int rows) throws DMLRuntimeException {
+ CSRPointer me = this;
+ CSRPointer that = new CSRPointer(me.getGPUContext());
+
+ that.allocateMatDescrPointer();
+ long totalSize = estimateSize(me.nnz, rows);
+ that.gpuContext.ensureFreeSpace(totalSize);
+
+ that.nnz = me.nnz;
+ that.val = allocate(that.nnz * Sizeof.DOUBLE);
+ that.rowPtr = allocate(rows * Sizeof.DOUBLE);
+ that.colInd = allocate(that.nnz * Sizeof.DOUBLE);
+
+ cudaMemcpy(that.val, me.val, that.nnz * Sizeof.DOUBLE, cudaMemcpyDeviceToDevice);
+ cudaMemcpy(that.rowPtr, me.rowPtr, rows * Sizeof.DOUBLE, cudaMemcpyDeviceToDevice);
+ cudaMemcpy(that.colInd, me.colInd, that.nnz * Sizeof.DOUBLE, cudaMemcpyDeviceToDevice);
+
+ return that;
+ }
+
+ private Pointer allocate(long size) throws DMLRuntimeException {
+ return getGPUContext().allocate(size);
+ }
+
+ private void cudaFreeHelper(Pointer toFree, boolean eager) throws DMLRuntimeException {
+ getGPUContext().cudaFreeHelper(toFree, eager);
+ }
+
+ private GPUContext getGPUContext() {
+ return gpuContext;
+ }
+
+ // ==============================================================================================
+
+ /**
+ * Check for ultra sparsity
+ *
+ * @param rows number of rows
+ * @param cols number of columns
+ * @return true if ultra sparse
+ */
+ public boolean isUltraSparse(int rows, int cols) {
+ double sp = ((double) nnz / rows / cols);
+ return sp < ULTRA_SPARSITY_TURN_POINT;
+ }
+
+ /**
+ * Initializes {@link #descr} to CUSPARSE_MATRIX_TYPE_GENERAL,
+ * the default that works for DGEMM.
+ */
+ private void allocateMatDescrPointer() {
+ this.descr = getDefaultCuSparseMatrixDescriptor();
+ }
+
+ /**
+ * Copies this CSR matrix on the GPU to a dense column-major matrix
+ * on the GPU. This is a temporary matrix for operations such as
+ * cusparseDcsrmv.
+ * Since the allocated matrix is temporary, bookkeeping is not updated.
+ * The caller is responsible for calling "free" on the returned Pointer object
+ *
+ * @param cusparseHandle a valid {@link cusparseHandle}
+ * @param cublasHandle a valid {@link cublasHandle}
+ * @param rows number of rows in this CSR matrix
+ * @param cols number of columns in this CSR matrix
+ * @return A {@link Pointer} to the allocated dense matrix (in column-major format)
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public Pointer toColumnMajorDenseMatrix(cusparseHandle cusparseHandle, cublasHandle cublasHandle, int rows,
+ int cols) throws DMLRuntimeException {
+ LOG.trace("GPU : sparse -> column major dense (inside CSRPointer) on " + this + ", GPUContext="
+ + getGPUContext());
+ long size = ((long) rows) * getDoubleSizeOf((long) cols);
+ Pointer A = allocate(size);
+ // If this sparse block is empty, the allocated dense matrix, initialized to zeroes, will be returned.
+ if (val != null && rowPtr != null && colInd != null && nnz > 0) {
+ // Note: cusparseDcsr2dense method cannot handle empty blocks
+ cusparseDcsr2dense(cusparseHandle, rows, cols, descr, val, rowPtr, colInd, A, rows);
+ //cudaDeviceSynchronize;
+ } else {
+ LOG.debug("in CSRPointer, the values array, row pointers array or column indices array was null");
+ }
+ return A;
+ }
+
+ /**
+ * Calls cudaFree lazily on the allocated {@link Pointer} instances
+ *
+ * @throws DMLRuntimeException ?
+ */
+ public void deallocate() throws DMLRuntimeException {
+ deallocate(false);
+ }
+
+ /**
+ * Calls cudaFree lazily or eagerly on the allocated {@link Pointer} instances
+ *
+ * @param eager whether to do eager or lazy cudaFrees
+ * @throws DMLRuntimeException ?
+ */
+ public void deallocate(boolean eager) throws DMLRuntimeException {
+ if (nnz > 0) {
+ cudaFreeHelper(val, eager);
+ cudaFreeHelper(rowPtr, eager);
+ cudaFreeHelper(colInd, eager);
+ }
+ }
+
+ @Override
+ public String toString() {
+ return "CSRPointer{" + "nnz=" + nnz + '}';
+ }
}
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/ExecutionConfig.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/ExecutionConfig.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/ExecutionConfig.java
index ce5c5ff..ef000c2 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/ExecutionConfig.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/ExecutionConfig.java
@@ -31,25 +31,24 @@ import jcuda.driver.CUstream;
* Java Wrapper to specify CUDA execution configuration for launching custom kernels
*/
public class ExecutionConfig {
- public int gridDimX;
+ public int gridDimX;
public int gridDimY = 1;
public int gridDimZ = 1;
- public int blockDimX;
+ public int blockDimX;
public int blockDimY = 1;
public int blockDimZ = 1;
public int sharedMemBytes = 0;
public CUstream stream = null;
-
-
private static HashMap<Integer, Integer> maxBlockDimForDevice = new HashMap<Integer, Integer>();
/**
* Convenience constructor for setting the number of blocks, number of threads and the
* shared memory size
- * @param gridDimX Number of blocks (for CUDA Kernel)
- * @param blockDimX Number of threads per block (for CUDA Kernel)
- * @param sharedMemBytes Amount of Shared memory (for CUDA Kernel)
+ *
+ * @param gridDimX Number of blocks (for CUDA Kernel)
+ * @param blockDimX Number of threads per block (for CUDA Kernel)
+ * @param sharedMemBytes Amount of Shared memory (for CUDA Kernel)
*/
public ExecutionConfig(int gridDimX, int blockDimX, int sharedMemBytes) {
this.gridDimX = gridDimX;
@@ -58,13 +57,13 @@ public class ExecutionConfig {
}
/**
- * Use this for simple vector operations and use following in the kernel
- * <code>
- * int index = blockIdx.x * blockDim.x + threadIdx.x
+ * Use this for simple vector operations and use following in the kernel
+ * <code>
+ * int index = blockIdx.x * blockDim.x + threadIdx.x
* </code>
- *
+ * <p>
* This tries to schedule as minimum grids as possible.
- *
+ *
* @param numCells number of cells
* @return execution configuration
* @throws DMLRuntimeException if DMLRuntimeException occurs
@@ -72,19 +71,19 @@ public class ExecutionConfig {
public static ExecutionConfig getConfigForSimpleVectorOperations(int numCells) throws DMLRuntimeException {
int deviceNumber = 0;
int blockDimX = getMaxBlockDim(deviceNumber);
- int gridDimX = (int)Math.ceil((double)numCells / blockDimX);
+ int gridDimX = (int) Math.ceil((double) numCells / blockDimX);
return new ExecutionConfig(gridDimX, blockDimX);
}
-
+
/**
- * Use this for simple matrix operations and use following in the kernel
- * <code>
+ * Use this for simple matrix operations and use following in the kernel
+ * <code>
* int ix = blockIdx.x * blockDim.x + threadIdx.x;
* int iy = blockIdx.y * blockDim.y + threadIdx.y;
* </code>
- *
+ * <p>
* This tries to schedule as minimum grids as possible.
- *
+ *
* @param rlen number of rows
* @param clen number of columns
* @return execution configuration
@@ -94,45 +93,45 @@ public class ExecutionConfig {
int deviceNumber = 0;
int maxBlockDim = getMaxBlockDim(deviceNumber);
int blockDimX = (int) Math.min(maxBlockDim, rlen);
- int gridDimX = (int)Math.ceil((double)rlen / blockDimX);
- int blockDimY = (int)Math.min(Math.floor(((double)maxBlockDim)/blockDimX), clen);
- int gridDimY = (int)Math.ceil((double)clen / blockDimY);
+ int gridDimX = (int) Math.ceil((double) rlen / blockDimX);
+ int blockDimY = (int) Math.min(Math.floor(((double) maxBlockDim) / blockDimX), clen);
+ int gridDimY = (int) Math.ceil((double) clen / blockDimY);
return new ExecutionConfig(gridDimX, gridDimY, blockDimX, blockDimY);
}
-
+
public ExecutionConfig(int gridDimX, int blockDimX) {
this.gridDimX = gridDimX;
this.blockDimX = blockDimX;
}
-
+
public ExecutionConfig(int gridDimX, int gridDimY, int blockDimX, int blockDimY) {
this.gridDimX = gridDimX;
this.gridDimY = gridDimY;
this.blockDimX = blockDimX;
this.blockDimY = blockDimY;
}
-
-
+
/**
- * Get the CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X of the given device
- *
+ * Get the CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X of the given device
+ *
* @param deviceNumber device number of the given device
* @return The maximum block dimension, in x-direction
* @throws DMLRuntimeException if DMLRuntimeException occurs
*/
- private static int getMaxBlockDim(int deviceNumber) throws DMLRuntimeException {
-// return 32;
- // TODO: Use JCudaDriver.cuOccupancyMaxPotentialBlockSize to chose the block size that maximizes occupancy
- Integer ret = maxBlockDimForDevice.get(deviceNumber);
- if(ret == null) {
- CUdevice device = new CUdevice();
- JCudaKernels.checkResult(jcuda.driver.JCudaDriver.cuDeviceGet(device, deviceNumber));
- int maxBlockDimX[] = {0};
- jcuda.driver.JCudaDriver.cuDeviceGetAttribute(maxBlockDimX, CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, device);
- maxBlockDimForDevice.put(deviceNumber, maxBlockDimX[0]);
- return maxBlockDimX[0];
- }
- return ret;
- }
-
- }
+ private static int getMaxBlockDim(int deviceNumber) throws DMLRuntimeException {
+ // return 32;
+ // TODO: Use JCudaDriver.cuOccupancyMaxPotentialBlockSize to chose the block size that maximizes occupancy
+ Integer ret = maxBlockDimForDevice.get(deviceNumber);
+ if (ret == null) {
+ CUdevice device = new CUdevice();
+ JCudaKernels.checkResult(jcuda.driver.JCudaDriver.cuDeviceGet(device, deviceNumber));
+ int maxBlockDimX[] = { 0 };
+ jcuda.driver.JCudaDriver
+ .cuDeviceGetAttribute(maxBlockDimX, CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, device);
+ maxBlockDimForDevice.put(deviceNumber, maxBlockDimX[0]);
+ return maxBlockDimX[0];
+ }
+ return ret;
+ }
+
+}
[2/3] systemml git commit: [FIX] Fixed nested parfor for GPUs
Posted by na...@apache.org.
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContext.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContext.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContext.java
index 8da67ea..b3c19ef 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContext.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContext.java
@@ -18,14 +18,24 @@
*/
package org.apache.sysml.runtime.instructions.gpu.context;
-import jcuda.Pointer;
-import jcuda.jcublas.cublasHandle;
-import jcuda.jcudnn.cudnnHandle;
-import jcuda.jcusolver.cusolverDnHandle;
-import jcuda.jcusolver.cusolverSpHandle;
-import jcuda.jcusparse.cusparseHandle;
-import jcuda.runtime.JCuda;
-import jcuda.runtime.cudaDeviceProp;
+import static jcuda.jcublas.JCublas2.cublasCreate;
+import static jcuda.jcublas.JCublas2.cublasDestroy;
+import static jcuda.jcudnn.JCudnn.cudnnCreate;
+import static jcuda.jcudnn.JCudnn.cudnnDestroy;
+import static jcuda.jcusolver.JCusolverDn.cusolverDnCreate;
+import static jcuda.jcusolver.JCusolverDn.cusolverDnDestroy;
+import static jcuda.jcusolver.JCusolverSp.cusolverSpCreate;
+import static jcuda.jcusolver.JCusolverSp.cusolverSpDestroy;
+import static jcuda.jcusparse.JCusparse.cusparseCreate;
+import static jcuda.jcusparse.JCusparse.cusparseDestroy;
+import static jcuda.runtime.JCuda.cudaDeviceScheduleBlockingSync;
+import static jcuda.runtime.JCuda.cudaFree;
+import static jcuda.runtime.JCuda.cudaGetDeviceCount;
+import static jcuda.runtime.JCuda.cudaMalloc;
+import static jcuda.runtime.JCuda.cudaMemGetInfo;
+import static jcuda.runtime.JCuda.cudaMemset;
+import static jcuda.runtime.JCuda.cudaSetDevice;
+import static jcuda.runtime.JCuda.cudaSetDeviceFlags;
import java.util.ArrayList;
import java.util.Collections;
@@ -45,24 +55,14 @@ import org.apache.sysml.runtime.instructions.gpu.GPUInstruction;
import org.apache.sysml.utils.GPUStatistics;
import org.apache.sysml.utils.LRUCacheMap;
-import static jcuda.jcublas.JCublas2.cublasCreate;
-import static jcuda.jcublas.JCublas2.cublasDestroy;
-import static jcuda.jcudnn.JCudnn.cudnnCreate;
-import static jcuda.jcudnn.JCudnn.cudnnDestroy;
-import static jcuda.jcusolver.JCusolverDn.cusolverDnCreate;
-import static jcuda.jcusolver.JCusolverDn.cusolverDnDestroy;
-import static jcuda.jcusolver.JCusolverSp.cusolverSpCreate;
-import static jcuda.jcusolver.JCusolverSp.cusolverSpDestroy;
-import static jcuda.jcusparse.JCusparse.cusparseCreate;
-import static jcuda.jcusparse.JCusparse.cusparseDestroy;
-import static jcuda.runtime.JCuda.cudaDeviceScheduleBlockingSync;
-import static jcuda.runtime.JCuda.cudaFree;
-import static jcuda.runtime.JCuda.cudaGetDeviceCount;
-import static jcuda.runtime.JCuda.cudaMalloc;
-import static jcuda.runtime.JCuda.cudaMemGetInfo;
-import static jcuda.runtime.JCuda.cudaMemset;
-import static jcuda.runtime.JCuda.cudaSetDevice;
-import static jcuda.runtime.JCuda.cudaSetDeviceFlags;
+import jcuda.Pointer;
+import jcuda.jcublas.cublasHandle;
+import jcuda.jcudnn.cudnnHandle;
+import jcuda.jcusolver.cusolverDnHandle;
+import jcuda.jcusolver.cusolverSpHandle;
+import jcuda.jcusparse.cusparseHandle;
+import jcuda.runtime.JCuda;
+import jcuda.runtime.cudaDeviceProp;
/**
* Represents a context per GPU accessible through the same JVM
@@ -71,606 +71,643 @@ import static jcuda.runtime.JCuda.cudaSetDeviceFlags;
public class GPUContext {
protected static final Log LOG = LogFactory.getLog(GPUContext.class.getName());
+ /**
+ * currently employed eviction policy
+ */
+ public final EvictionPolicy evictionPolicy = EvictionPolicy.LRU;
+ /**
+ * The minimum CUDA Compute capability needed for SystemML.
+ * After compute capability 3.0, 2^31 - 1 blocks and 1024 threads per block are supported.
+ * If SystemML needs to run on an older card, this logic can be revisited.
+ */
+ final int MAJOR_REQUIRED = 3;
+ final int MINOR_REQUIRED = 0;
+ /**
+ * active device assigned to this GPUContext instance
+ */
+ private final int deviceNum;
+ // Invoke cudaMemGetInfo to get available memory information. Useful if GPU is shared among multiple application.
+ public double GPU_MEMORY_UTILIZATION_FACTOR = ConfigurationManager.getDMLConfig()
+ .getDoubleValue(DMLConfig.GPU_MEMORY_UTILIZATION_FACTOR);
+ /**
+ * Map of free blocks allocate on GPU. maps size_of_block -> pointer on GPU
+ */
+ private LRUCacheMap<Long, LinkedList<Pointer>> freeCUDASpaceMap = new LRUCacheMap<>();
+ /**
+ * To record size of allocated blocks
+ */
+ private HashMap<Pointer, Long> cudaBlockSizeMap = new HashMap<>();
+ /**
+ * list of allocated {@link GPUObject} instances allocated on {@link GPUContext#deviceNum} GPU
+ * These are matrices allocated on the GPU on which rmvar hasn't been called yet.
+ * If a {@link GPUObject} has more than one lock on it, it cannot be freed
+ * If it has zero locks on it, it can be freed, but it is preferrable to keep it around
+ * so that an extraneous host to dev transfer can be avoided
+ */
+ private ArrayList<GPUObject> allocatedGPUObjects = new ArrayList<>();
+ /**
+ * cudnnHandle for Deep Neural Network operations on the GPU
+ */
+ private cudnnHandle cudnnHandle;
+ /**
+ * cublasHandle for BLAS operations on the GPU
+ */
+ private cublasHandle cublasHandle;
+ /**
+ * cusparseHandle for certain sparse BLAS operations on the GPU
+ */
+ private cusparseHandle cusparseHandle;
+ /**
+ * cusolverDnHandle for invoking solve() function on dense matrices on the GPU
+ */
+ private cusolverDnHandle cusolverDnHandle;
+ /**
+ * cusolverSpHandle for invoking solve() function on sparse matrices on the GPU
+ */
+ private cusolverSpHandle cusolverSpHandle;
+ /**
+ * to launch custom CUDA kernel, specific to the active GPU for this GPUContext
+ */
+ private JCudaKernels kernels;
+
+ protected GPUContext(int deviceNum) throws DMLRuntimeException {
+ this.deviceNum = deviceNum;
+ cudaSetDevice(deviceNum);
+
+ cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+
+ long free[] = { 0 };
+ long total[] = { 0 };
+ cudaMemGetInfo(free, total);
+
+ long start = System.nanoTime();
+ cudnnHandle = new cudnnHandle();
+ cudnnCreate(cudnnHandle);
+ cublasHandle = new cublasHandle();
+ cublasCreate(cublasHandle);
+ // For cublas v2, cublasSetPointerMode tells Cublas whether to expect scalar arguments on device or on host
+ // This applies to arguments like "alpha" in Dgemm, and "y" in Ddot.
+ // cublasSetPointerMode(LibMatrixCUDA.cublasHandle, cublasPointerMode.CUBLAS_POINTER_MODE_DEVICE);
+ cusparseHandle = new cusparseHandle();
+ cusparseCreate(cusparseHandle);
+
+ cusolverDnHandle = new cusolverDnHandle();
+ cusolverDnCreate(cusolverDnHandle);
+ cusolverSpHandle = new cusolverSpHandle();
+ cusolverSpCreate(cusolverSpHandle);
+
+ kernels = new JCudaKernels(deviceNum);
+
+ GPUStatistics.cudaLibrariesInitTime = System.nanoTime() - start;
+ LOG.info(" GPU memory - Total: " + (total[0] * (1e-6)) + " MB, Available: " + (free[0] * (1e-6)) + " MB on "
+ + this);
- /** Eviction policies for {@link GPUContext#evict(long)} */
- public enum EvictionPolicy {
- LRU, LFU, MIN_EVICT
}
- /** currently employed eviction policy */
- public final EvictionPolicy evictionPolicy = EvictionPolicy.LRU;
+ public static int cudaGetDevice() {
+ int[] device = new int[1];
+ JCuda.cudaGetDevice(device);
+ return device[0];
+ }
- /** Map of free blocks allocate on GPU. maps size_of_block -> pointer on GPU */
- private LRUCacheMap<Long, LinkedList<Pointer>> freeCUDASpaceMap = new LRUCacheMap<>();
+ public int getDeviceNum() {
+ return deviceNum;
+ }
- /** To record size of allocated blocks */
- private HashMap<Pointer, Long> cudaBlockSizeMap = new HashMap<>();
+ /**
+ * Sets the device for the calling thread.
+ * This method must be called after
+ * {@link org.apache.sysml.runtime.controlprogram.context.ExecutionContext#getGPUContext(int)}
+ * If in a multi-threaded env like parfor, this method must be called when in the
+ * appropriate thread
+ */
+ public void initializeThread() {
+ cudaSetDevice(deviceNum);
+ }
+
+ /**
+ * Convenience method for {@link #allocate(String, long, int)}, defaults statsCount to 1.
+ *
+ * @param size size of data (in bytes) to allocate
+ * @return jcuda pointer
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public Pointer allocate(long size) throws DMLRuntimeException {
+ return allocate(null, size, 1);
+ }
+
+ /**
+ * Convenience method for {@link #allocate(String, long, int)}, defaults statsCount to 1.
+ *
+ * @param instructionName name of instruction for which to record per instruction performance statistics, null if don't want to record
+ * @param size size of data (in bytes) to allocate
+ * @return jcuda pointer
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public Pointer allocate(String instructionName, long size) throws DMLRuntimeException {
+ return allocate(instructionName, size, 1);
+ }
- /** active device assigned to this GPUContext instance */
- private final int deviceNum;
+ /**
+ * Allocates temporary space on the device.
+ * Does not update bookkeeping.
+ * The caller is responsible for freeing up after usage.
+ *
+ * @param instructionName name of instruction for which to record per instruction performance statistics, null if don't want to record
+ * @param size Size of data (in bytes) to allocate
+ * @param statsCount amount to increment the cudaAllocCount by
+ * @return jcuda Pointer
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ public Pointer allocate(String instructionName, long size, int statsCount) throws DMLRuntimeException {
+ long t0 = 0, t1 = 0, end = 0;
+ Pointer A;
+ if (freeCUDASpaceMap.containsKey(size)) {
+ LOG.trace("GPU : in allocate from instruction " + instructionName + ", found free block of size " + (size
+ / 1024.0) + " Kbytes from previously allocated block on " + this);
+ if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
+ t0 = System.nanoTime();
+ LinkedList<Pointer> freeList = freeCUDASpaceMap.get(size);
+ A = freeList.pop();
+ if (freeList.isEmpty())
+ freeCUDASpaceMap.remove(size);
+ if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
+ GPUStatistics
+ .maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_REUSE, System.nanoTime() - t0);
+ } else {
+ LOG.trace(
+ "GPU : in allocate from instruction " + instructionName + ", allocating new block of size " + (size
+ / 1024.0) + " Kbytes on " + this);
+ if (DMLScript.STATISTICS)
+ t0 = System.nanoTime();
+ ensureFreeSpace(instructionName, size);
+ A = new Pointer();
+ cudaMalloc(A, size);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaAllocTime.getAndAdd(System.nanoTime() - t0);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaAllocCount.getAndAdd(statsCount);
+ if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
+ GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_ALLOCATE,
+ System.nanoTime() - t0);
+ }
+ // Set all elements to 0 since newly allocated space will contain garbage
+ if (DMLScript.STATISTICS)
+ t1 = System.nanoTime();
+ LOG.trace("GPU : in allocate from instruction " + instructionName + ", setting block of size " + (size / 1024.0)
+ + " Kbytes to zero on " + this);
+ cudaMemset(A, 0, size);
+ if (DMLScript.STATISTICS)
+ end = System.nanoTime();
+ if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
+ GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_SET_ZERO, end - t1);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaMemSet0Time.getAndAdd(end - t1);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaMemSet0Count.getAndAdd(1);
+ cudaBlockSizeMap.put(A, size);
+ return A;
- /** list of allocated {@link GPUObject} instances allocated on {@link GPUContext#deviceNum} GPU
- * These are matrices allocated on the GPU on which rmvar hasn't been called yet.
- * If a {@link GPUObject} has more than one lock on it, it cannot be freed
- * If it has zero locks on it, it can be freed, but it is preferrable to keep it around
- * so that an extraneous host to dev transfer can be avoided */
- private ArrayList<GPUObject> allocatedGPUObjects = new ArrayList<>();
+ }
- /** cudnnHandle for Deep Neural Network operations on the GPU */
- private cudnnHandle cudnnHandle;
+ /**
+ * Does lazy cudaFree calls
+ *
+ * @param toFree {@link Pointer} instance to be freed
+ */
+ public void cudaFreeHelper(final Pointer toFree) {
+ cudaFreeHelper(null, toFree, false);
+ }
- /** cublasHandle for BLAS operations on the GPU */
- private cublasHandle cublasHandle;
+ /**
+ * does lazy/eager cudaFree calls
+ *
+ * @param toFree {@link Pointer} instance to be freed
+ * @param eager true if to be done eagerly
+ */
+ public void cudaFreeHelper(final Pointer toFree, boolean eager) {
+ cudaFreeHelper(null, toFree, eager);
+ }
- /** cusparseHandle for certain sparse BLAS operations on the GPU */
- private cusparseHandle cusparseHandle;
+ /**
+ * Does lazy cudaFree calls
+ *
+ * @param instructionName name of the instruction for which to record per instruction free time, null if do not want to record
+ * @param toFree {@link Pointer} instance to be freed
+ */
+ public void cudaFreeHelper(String instructionName, final Pointer toFree) {
+ cudaFreeHelper(instructionName, toFree, false);
+ }
- /** cusolverDnHandle for invoking solve() function on dense matrices on the GPU */
- private cusolverDnHandle cusolverDnHandle;
+ /**
+ * Does cudaFree calls, lazily
+ *
+ * @param instructionName name of the instruction for which to record per instruction free time, null if do not want to record
+ * @param toFree {@link Pointer} instance to be freed
+ * @param eager true if to be done eagerly
+ */
+ public void cudaFreeHelper(String instructionName, final Pointer toFree, boolean eager) {
+ Pointer dummy = new Pointer();
+ if (toFree == dummy) // trying to free a null pointer
+ return;
+ long t0 = 0;
+ assert cudaBlockSizeMap.containsKey(
+ toFree) : "ERROR : Internal state corrupted, cache block size map is not aware of a block it trying to free up";
+ long size = cudaBlockSizeMap.get(toFree);
+ if (eager) {
+ LOG.trace("GPU : eagerly freeing cuda memory [ " + toFree + " ] for instruction " + instructionName + " on "
+ + this);
+ if (DMLScript.STATISTICS)
+ t0 = System.nanoTime();
+ cudaFree(toFree);
+ cudaBlockSizeMap.remove(toFree);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaDeAllocTime.addAndGet(System.nanoTime() - t0);
+ if (DMLScript.STATISTICS)
+ GPUStatistics.cudaDeAllocCount.addAndGet(1);
+ if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
+ GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_CUDA_FREE,
+ System.nanoTime() - t0);
+ } else {
+ LOG.trace("GPU : lazily freeing cuda memory for instruction " + instructionName + " on " + this);
+ LinkedList<Pointer> freeList = freeCUDASpaceMap.get(size);
+ if (freeList == null) {
+ freeList = new LinkedList<Pointer>();
+ freeCUDASpaceMap.put(size, freeList);
+ }
+ if (freeList.contains(toFree))
+ throw new RuntimeException("GPU : Internal state corrupted, double free");
+ freeList.add(toFree);
+ }
+ }
- /** cusolverSpHandle for invoking solve() function on sparse matrices on the GPU */
- private cusolverSpHandle cusolverSpHandle;
+ /**
+ * Thin wrapper over {@link GPUContext#evict(long)}
+ *
+ * @param size size to check
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ void ensureFreeSpace(long size) throws DMLRuntimeException {
+ ensureFreeSpace(null, size);
+ }
- /** to launch custom CUDA kernel, specific to the active GPU for this GPUContext */
- private JCudaKernels kernels;
+ /**
+ * Thin wrapper over {@link GPUContext#evict(long)}
+ *
+ * @param instructionName instructionName name of the instruction for which performance measurements are made
+ * @param size size to check
+ * @throws DMLRuntimeException if DMLRuntimeException occurs
+ */
+ void ensureFreeSpace(String instructionName, long size) throws DMLRuntimeException {
+ if (size >= getAvailableMemory()) {
+ evict(instructionName, size);
+ }
+ }
/**
- * The minimum CUDA Compute capability needed for SystemML.
- * After compute capability 3.0, 2^31 - 1 blocks and 1024 threads per block are supported.
- * If SystemML needs to run on an older card, this logic can be revisited.
+ * Convenience wrapper over {@link GPUContext#evict(String, long)}
+ *
+ * @param GPUSize Desired size to be freed up on the GPU
+ * @throws DMLRuntimeException If no blocks to free up or if not enough blocks with zero locks on them.
*/
- final int MAJOR_REQUIRED = 3;
- final int MINOR_REQUIRED = 0;
+ protected void evict(final long GPUSize) throws DMLRuntimeException {
+ evict(null, GPUSize);
+ }
+
+ /**
+ * Memory on the GPU is tried to be freed up until either a chunk of needed size is freed up
+ * or it fails.
+ * First the set of reusable blocks is freed up. If that isn't enough, the set of allocated matrix
+ * blocks with zero locks on them is freed up.
+ * The process cycles through the sorted list of allocated {@link GPUObject} instances. Sorting is based on
+ * number of (read) locks that have been obtained on it (reverse order). It repeatedly frees up
+ * blocks on which there are zero locks until the required size has been freed up.
+ * // TODO: update it with hybrid policy
+ *
+ * @param instructionName name of the instruction for which performance measurements are made
+ * @param neededSize desired size to be freed up on the GPU
+ * @throws DMLRuntimeException If no reusable memory blocks to free up or if not enough matrix blocks with zero locks on them.
+ */
+ protected void evict(String instructionName, final long neededSize) throws DMLRuntimeException {
+ LOG.trace("GPU : evict called from " + instructionName + " for size " + neededSize + " on " + this);
+ GPUStatistics.cudaEvictionCount.addAndGet(1);
+ // Release the set of free blocks maintained in a GPUObject.freeCUDASpaceMap
+ // to free up space
+ LRUCacheMap<Long, LinkedList<Pointer>> lruCacheMap = freeCUDASpaceMap;
+ while (lruCacheMap.size() > 0) {
+ if (neededSize <= getAvailableMemory())
+ break;
+ Map.Entry<Long, LinkedList<Pointer>> toFreeListPair = lruCacheMap.removeAndGetLRUEntry();
+ LinkedList<Pointer> toFreeList = toFreeListPair.getValue();
+ Long size = toFreeListPair.getKey();
+ Pointer toFree = toFreeList.pop();
+ if (toFreeList.isEmpty())
+ lruCacheMap.remove(size);
+ cudaFreeHelper(instructionName, toFree, true);
+ }
+
+ if (neededSize <= getAvailableMemory())
+ return;
+
+ if (allocatedGPUObjects.size() == 0) {
+ throw new DMLRuntimeException(
+ "There is not enough memory on device for this matrix, request (" + neededSize + ")");
+ }
+
+ Collections.sort(allocatedGPUObjects, new Comparator<GPUObject>() {
+ @Override
+ public int compare(GPUObject p1, GPUObject p2) {
+ long p1Val = p1.locks.get();
+ long p2Val = p2.locks.get();
+
+ if (p1Val > 0 && p2Val > 0) {
+ // Both are locked, so don't sort
+ return 0;
+ } else if (p1Val > 0 || p2Val > 0) {
+ // Put the unlocked one to RHS
+ return Long.compare(p2Val, p1Val);
+ } else {
+ // Both are unlocked
+
+ if (evictionPolicy == EvictionPolicy.MIN_EVICT) {
+ long p1Size = 0;
+ long p2Size = 0;
+ try {
+ p1Size = p1.getSizeOnDevice() - neededSize;
+ p2Size = p2.getSizeOnDevice() - neededSize;
+ } catch (DMLRuntimeException e) {
+ throw new RuntimeException(e);
+ }
+
+ if (p1Size >= 0 && p2Size >= 0) {
+ return Long.compare(p2Size, p1Size);
+ } else {
+ return Long.compare(p1Size, p2Size);
+ }
+ } else if (evictionPolicy == EvictionPolicy.LRU || evictionPolicy == EvictionPolicy.LFU) {
+ return Long.compare(p2.timestamp.get(), p1.timestamp.get());
+ } else {
+ throw new RuntimeException("Unsupported eviction policy:" + evictionPolicy.name());
+ }
+ }
+ }
+ });
+
+ while (neededSize > getAvailableMemory() && allocatedGPUObjects.size() > 0) {
+ GPUObject toBeRemoved = allocatedGPUObjects.get(allocatedGPUObjects.size() - 1);
+ if (toBeRemoved.locks.get() > 0) {
+ throw new DMLRuntimeException(
+ "There is not enough memory on device for this matrix, request (" + neededSize + ")");
+ }
+ if (toBeRemoved.dirty) {
+ toBeRemoved.copyFromDeviceToHost();
+ }
+
+ toBeRemoved.clearData(true);
+ }
+ }
+
+ /**
+ * Whether the GPU associated with this {@link GPUContext} has recorded the usage of a certain block
+ *
+ * @param o the block
+ * @return true if present, false otherwise
+ */
+ public boolean isBlockRecorded(GPUObject o) {
+ return allocatedGPUObjects.contains(o);
+ }
+
+ /**
+ * @param o {@link GPUObject} instance to record
+ * @see GPUContext#allocatedGPUObjects
+ * Records the usage of a matrix block
+ */
+ public void recordBlockUsage(GPUObject o) {
+ allocatedGPUObjects.add(o);
+ }
- // Invoke cudaMemGetInfo to get available memory information. Useful if GPU is shared among multiple application.
- public double GPU_MEMORY_UTILIZATION_FACTOR = ConfigurationManager.getDMLConfig().getDoubleValue(DMLConfig.GPU_MEMORY_UTILIZATION_FACTOR);
-
- protected GPUContext(int deviceNum) throws DMLRuntimeException {
- this.deviceNum = deviceNum;
- cudaSetDevice(deviceNum);
-
- cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
-
- long free[] = {0};
- long total[] = {0};
- cudaMemGetInfo(free, total);
-
- long start = System.nanoTime();
- cudnnHandle = new cudnnHandle();
- cudnnCreate(cudnnHandle);
- cublasHandle = new cublasHandle();
- cublasCreate(cublasHandle);
- // For cublas v2, cublasSetPointerMode tells Cublas whether to expect scalar arguments on device or on host
- // This applies to arguments like "alpha" in Dgemm, and "y" in Ddot.
- // cublasSetPointerMode(LibMatrixCUDA.cublasHandle, cublasPointerMode.CUBLAS_POINTER_MODE_DEVICE);
- cusparseHandle = new cusparseHandle();
- cusparseCreate(cusparseHandle);
-
- cusolverDnHandle = new cusolverDnHandle();
- cusolverDnCreate(cusolverDnHandle);
- cusolverSpHandle = new cusolverSpHandle();
- cusolverSpCreate(cusolverSpHandle);
-
- kernels = new JCudaKernels(deviceNum);
-
- GPUStatistics.cudaLibrariesInitTime = System.nanoTime() - start;
- LOG.info(" GPU memory - Total: " + (total[0] * (1e-6)) + " MB, Available: " + (free[0] * (1e-6)) + " MB on " + this);
-
- }
-
- public int getDeviceNum() {
- return deviceNum;
- }
-
- /**
- * Sets the device for the calling thread.
- * This method must be called after {@link GPUContextPool#getFromPool()}
- * is called.
- * If in a multi-threaded env like parfor, this method must be called when in the
- * appropriate thread
- */
- public void initializeThread() {
- cudaSetDevice(deviceNum);
- }
-
- public static int cudaGetDevice() {
- int[] device = new int[1];
- JCuda.cudaGetDevice(device);
- return device[0];
- }
-
- /**
- * Convenience method for {@link #allocate(String, long, int)}, defaults statsCount to 1.
- *
- * @param size size of data (in bytes) to allocate
- * @return jcuda pointer
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- public Pointer allocate(long size) throws DMLRuntimeException {
- return allocate(null, size, 1);
- }
-
- /**
- * Convenience method for {@link #allocate(String, long, int)}, defaults statsCount to 1.
- *
- * @param instructionName name of instruction for which to record per instruction performance statistics, null if don't want to record
- * @param size size of data (in bytes) to allocate
- * @return jcuda pointer
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- public Pointer allocate(String instructionName, long size) throws DMLRuntimeException {
- return allocate(instructionName, size, 1);
- }
-
- /**
- * Allocates temporary space on the device.
- * Does not update bookkeeping.
- * The caller is responsible for freeing up after usage.
- *
- * @param instructionName name of instruction for which to record per instruction performance statistics, null if don't want to record
- * @param size Size of data (in bytes) to allocate
- * @param statsCount amount to increment the cudaAllocCount by
- * @return jcuda Pointer
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- public Pointer allocate(String instructionName, long size, int statsCount) throws DMLRuntimeException {
- long t0 = 0, t1 = 0, end = 0;
- Pointer A;
- if (freeCUDASpaceMap.containsKey(size)) {
- LOG.trace("GPU : in allocate from instruction " + instructionName + ", found free block of size " + (size / 1024.0) + " Kbytes from previously allocated block on " + this);
- if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS) t0 = System.nanoTime();
- LinkedList<Pointer> freeList = freeCUDASpaceMap.get(size);
- A = freeList.pop();
- if (freeList.isEmpty())
- freeCUDASpaceMap.remove(size);
- if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
- GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_REUSE, System.nanoTime() - t0);
- } else {
- LOG.trace("GPU : in allocate from instruction " + instructionName + ", allocating new block of size " + (size / 1024.0) + " Kbytes on " + this);
- if (DMLScript.STATISTICS) t0 = System.nanoTime();
- ensureFreeSpace(instructionName, size);
- A = new Pointer();
- cudaMalloc(A, size);
- if (DMLScript.STATISTICS) GPUStatistics.cudaAllocTime.getAndAdd(System.nanoTime() - t0);
- if (DMLScript.STATISTICS) GPUStatistics.cudaAllocCount.getAndAdd(statsCount);
- if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
- GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_ALLOCATE, System.nanoTime() - t0);
- }
- // Set all elements to 0 since newly allocated space will contain garbage
- if (DMLScript.STATISTICS) t1 = System.nanoTime();
- LOG.trace("GPU : in allocate from instruction " + instructionName + ", setting block of size " + (size / 1024.0) + " Kbytes to zero on " + this);
- cudaMemset(A, 0, size);
- if (DMLScript.STATISTICS) end = System.nanoTime();
- if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
- GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_SET_ZERO, end - t1);
- if (DMLScript.STATISTICS) GPUStatistics.cudaMemSet0Time.getAndAdd(end - t1);
- if (DMLScript.STATISTICS) GPUStatistics.cudaMemSet0Count.getAndAdd(1);
- cudaBlockSizeMap.put(A, size);
- return A;
-
- }
-
- /**
- * Does lazy cudaFree calls
- *
- * @param toFree {@link Pointer} instance to be freed
- */
- public void cudaFreeHelper(final Pointer toFree) {
- cudaFreeHelper(null, toFree, false);
- }
-
- /**
- * does lazy/eager cudaFree calls
- *
- * @param toFree {@link Pointer} instance to be freed
- * @param eager true if to be done eagerly
- */
- public void cudaFreeHelper(final Pointer toFree, boolean eager) {
- cudaFreeHelper(null, toFree, eager);
- }
-
- /**
- * Does lazy cudaFree calls
- *
- * @param instructionName name of the instruction for which to record per instruction free time, null if do not want to record
- * @param toFree {@link Pointer} instance to be freed
- */
- public void cudaFreeHelper(String instructionName, final Pointer toFree) {
- cudaFreeHelper(instructionName, toFree, false);
- }
-
- /**
- * Does cudaFree calls, lazily
- *
- * @param instructionName name of the instruction for which to record per instruction free time, null if do not want to record
- * @param toFree {@link Pointer} instance to be freed
- * @param eager true if to be done eagerly
- */
- public void cudaFreeHelper(String instructionName, final Pointer toFree, boolean eager) {
- Pointer dummy = new Pointer();
- if (toFree == dummy) // trying to free a null pointer
- return;
- long t0 = 0;
- assert cudaBlockSizeMap.containsKey(toFree) : "ERROR : Internal state corrupted, cache block size map is not aware of a block it trying to free up";
- long size = cudaBlockSizeMap.get(toFree);
- if (eager) {
- LOG.trace("GPU : eagerly freeing cuda memory [ " + toFree + " ] for instruction " + instructionName + " on " + this);
- if (DMLScript.STATISTICS) t0 = System.nanoTime();
- cudaFree(toFree);
- cudaBlockSizeMap.remove(toFree);
- if (DMLScript.STATISTICS) GPUStatistics.cudaDeAllocTime.addAndGet(System.nanoTime() - t0);
- if (DMLScript.STATISTICS) GPUStatistics.cudaDeAllocCount.addAndGet(1);
- if (instructionName != null && GPUStatistics.DISPLAY_STATISTICS)
- GPUStatistics.maintainCPMiscTimes(instructionName, GPUInstruction.MISC_TIMER_CUDA_FREE, System.nanoTime() - t0);
- } else {
- LOG.trace("GPU : lazily freeing cuda memory for instruction " + instructionName + " on " + this);
- LinkedList<Pointer> freeList = freeCUDASpaceMap.get(size);
- if (freeList == null) {
- freeList = new LinkedList<Pointer>();
- freeCUDASpaceMap.put(size, freeList);
- }
- if (freeList.contains(toFree))
- throw new RuntimeException("GPU : Internal state corrupted, double free");
- freeList.add(toFree);
- }
- }
-
- /**
- * Thin wrapper over {@link GPUContext#evict(long)}
- *
- * @param size size to check
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- void ensureFreeSpace(long size) throws DMLRuntimeException {
- ensureFreeSpace(null, size);
- }
-
- /**
- * Thin wrapper over {@link GPUContext#evict(long)}
- *
- * @param instructionName instructionName name of the instruction for which performance measurements are made
- * @param size size to check
- * @throws DMLRuntimeException if DMLRuntimeException occurs
- */
- void ensureFreeSpace(String instructionName, long size) throws DMLRuntimeException {
- if (size >= getAvailableMemory()) {
- evict(instructionName, size);
- }
- }
-
- /**
- * Convenience wrapper over {@link GPUContext#evict(String, long)}
- *
- * @param GPUSize Desired size to be freed up on the GPU
- * @throws DMLRuntimeException If no blocks to free up or if not enough blocks with zero locks on them.
- */
- protected void evict(final long GPUSize) throws DMLRuntimeException {
- evict(null, GPUSize);
- }
-
- /**
- * Memory on the GPU is tried to be freed up until either a chunk of needed size is freed up
- * or it fails.
- * First the set of reusable blocks is freed up. If that isn't enough, the set of allocated matrix
- * blocks with zero locks on them is freed up.
- * The process cycles through the sorted list of allocated {@link GPUObject} instances. Sorting is based on
- * number of (read) locks that have been obtained on it (reverse order). It repeatedly frees up
- * blocks on which there are zero locks until the required size has been freed up.
- * // TODO: update it with hybrid policy
- *
- * @param instructionName name of the instruction for which performance measurements are made
- * @param neededSize desired size to be freed up on the GPU
- * @throws DMLRuntimeException If no reusable memory blocks to free up or if not enough matrix blocks with zero locks on them.
- */
- protected void evict(String instructionName, final long neededSize) throws DMLRuntimeException {
- LOG.trace("GPU : evict called from " + instructionName + " for size " + neededSize + " on " + this);
- GPUStatistics.cudaEvictionCount.addAndGet(1);
- // Release the set of free blocks maintained in a GPUObject.freeCUDASpaceMap
- // to free up space
- LRUCacheMap<Long, LinkedList<Pointer>> lruCacheMap = freeCUDASpaceMap;
- while (lruCacheMap.size() > 0) {
- if (neededSize <= getAvailableMemory())
- break;
- Map.Entry<Long, LinkedList<Pointer>> toFreeListPair = lruCacheMap.removeAndGetLRUEntry();
- LinkedList<Pointer> toFreeList = toFreeListPair.getValue();
- Long size = toFreeListPair.getKey();
- Pointer toFree = toFreeList.pop();
- if (toFreeList.isEmpty())
- lruCacheMap.remove(size);
- cudaFreeHelper(instructionName, toFree, true);
- }
-
- if (neededSize <= getAvailableMemory())
- return;
-
- if (allocatedGPUObjects.size() == 0) {
- throw new DMLRuntimeException("There is not enough memory on device for this matrix, request (" + neededSize + ")");
- }
-
- Collections.sort(allocatedGPUObjects, new Comparator<GPUObject>() {
- @Override
- public int compare(GPUObject p1, GPUObject p2) {
- long p1Val = p1.locks.get();
- long p2Val = p2.locks.get();
-
- if (p1Val > 0 && p2Val > 0) {
- // Both are locked, so don't sort
- return 0;
- } else if (p1Val > 0 || p2Val > 0) {
- // Put the unlocked one to RHS
- return Long.compare(p2Val, p1Val);
- } else {
- // Both are unlocked
-
- if (evictionPolicy == EvictionPolicy.MIN_EVICT) {
- long p1Size = 0;
- long p2Size = 0;
- try {
- p1Size = p1.getSizeOnDevice() - neededSize;
- p2Size = p2.getSizeOnDevice() - neededSize;
- } catch (DMLRuntimeException e) {
- throw new RuntimeException(e);
- }
-
- if (p1Size >= 0 && p2Size >= 0) {
- return Long.compare(p2Size, p1Size);
- } else {
- return Long.compare(p1Size, p2Size);
- }
- } else if (evictionPolicy == EvictionPolicy.LRU || evictionPolicy == EvictionPolicy.LFU) {
- return Long.compare(p2.timestamp.get(), p1.timestamp.get());
- } else {
- throw new RuntimeException("Unsupported eviction policy:" + evictionPolicy.name());
- }
- }
- }
- });
-
- while (neededSize > getAvailableMemory() && allocatedGPUObjects.size() > 0) {
- GPUObject toBeRemoved = allocatedGPUObjects.get(allocatedGPUObjects.size() - 1);
- if (toBeRemoved.locks.get() > 0) {
- throw new DMLRuntimeException("There is not enough memory on device for this matrix, request (" + neededSize + ")");
- }
- if (toBeRemoved.dirty) {
- toBeRemoved.copyFromDeviceToHost();
- }
-
- toBeRemoved.clearData(true);
- }
- }
-
- /**
- * Whether the GPU associated with this {@link GPUContext} has recorded the usage of a certain block
- *
- * @param o the block
- * @return true if present, false otherwise
- */
- public boolean isBlockRecorded(GPUObject o) {
- return allocatedGPUObjects.contains(o);
- }
-
- /**
- * @param o {@link GPUObject} instance to record
- * @see GPUContext#allocatedGPUObjects
- * Records the usage of a matrix block
- */
- public void recordBlockUsage(GPUObject o) {
- allocatedGPUObjects.add(o);
- }
-
- /**
- * @param o {@link GPUObject} instance to remove from the list of allocated GPU objects
- * @see GPUContext#allocatedGPUObjects
- * Records that a block is not used anymore
- */
- public void removeRecordedUsage(GPUObject o) {
- allocatedGPUObjects.remove(o);
- }
-
- /**
- * Gets the available memory on GPU that SystemML can use
- *
- * @return the available memory in bytes
- */
- public long getAvailableMemory() {
- long free[] = {0};
- long total[] = {0};
- cudaMemGetInfo(free, total);
- return (long) (free[0] * GPU_MEMORY_UTILIZATION_FACTOR);
- }
-
- /**
- * Makes sure that GPU that SystemML is trying to use has the minimum compute capability needed
- *
- * @throws DMLRuntimeException if the compute capability is less than what is required
- */
- public void ensureComputeCapability() throws DMLRuntimeException {
- int[] devices = {-1};
- cudaGetDeviceCount(devices);
- if (devices[0] == -1) {
- throw new DMLRuntimeException("Call to cudaGetDeviceCount returned 0 devices");
- }
- boolean isComputeCapable = true;
- for (int i = 0; i < devices[0]; i++) {
- cudaDeviceProp properties = GPUContextPool.getGPUProperties(i);
- int major = properties.major;
- int minor = properties.minor;
- if (major < MAJOR_REQUIRED) {
- isComputeCapable = false;
- } else if (major == MAJOR_REQUIRED && minor < MINOR_REQUIRED) {
- isComputeCapable = false;
- }
- }
- if (!isComputeCapable) {
- throw new DMLRuntimeException("One of the CUDA cards on the system has compute capability lower than " + MAJOR_REQUIRED + "." + MINOR_REQUIRED);
- }
- }
-
- public GPUObject createGPUObject(MatrixObject mo) {
- return new GPUObject(this, mo);
- }
-
- /**
- * Gets the device properties for the active GPU (set with cudaSetDevice())
- *
- * @return the device properties
- * @throws DMLRuntimeException ?
- */
- public cudaDeviceProp getGPUProperties() throws DMLRuntimeException {
- return GPUContextPool.getGPUProperties(deviceNum);
- }
-
- /**
- * Gets the maximum number of threads per block for "active" GPU
- *
- * @return the maximum number of threads per block
- * @throws DMLRuntimeException ?
- */
- public int getMaxThreadsPerBlock() throws DMLRuntimeException {
- cudaDeviceProp deviceProps = getGPUProperties();
- return deviceProps.maxThreadsPerBlock;
- }
-
- /**
- * Gets the maximum number of blocks supported by the active cuda device
- *
- * @return the maximum number of blocks supported
- * @throws DMLRuntimeException ?
- */
- public int getMaxBlocks() throws DMLRuntimeException {
- cudaDeviceProp deviceProp = getGPUProperties();
- return deviceProp.maxGridSize[0];
- }
-
- /**
- * Gets the shared memory per block supported by the active cuda device
- *
- * @return the shared memory per block
- * @throws DMLRuntimeException ?
- */
- public long getMaxSharedMemory() throws DMLRuntimeException {
- cudaDeviceProp deviceProp = getGPUProperties();
- return deviceProp.sharedMemPerBlock;
- }
-
- /**
- * Gets the warp size supported by the active cuda device
- *
- * @return the warp size
- * @throws DMLRuntimeException ?
- */
- public int getWarpSize() throws DMLRuntimeException {
- cudaDeviceProp deviceProp = getGPUProperties();
- return deviceProp.warpSize;
- }
-
- public cudnnHandle getCudnnHandle() {
- return cudnnHandle;
- }
-
- public cublasHandle getCublasHandle() {
- return cublasHandle;
- }
-
- public cusparseHandle getCusparseHandle() {
- return cusparseHandle;
- }
-
- public cusolverDnHandle getCusolverDnHandle() {
- return cusolverDnHandle;
- }
-
- public cusolverSpHandle getCusolverSpHandle() {
- return cusolverSpHandle;
- }
-
- public JCudaKernels getKernels() {
- return kernels;
- }
-
- /**
- * Destroys this GPUContext object
- *
- * @throws DMLRuntimeException if error
- */
- public void destroy() throws DMLRuntimeException {
- LOG.trace("GPU : this context was destroyed, this = " + this.toString());
- clearMemory();
- cudnnDestroy(cudnnHandle);
- cublasDestroy(cublasHandle);
- cusparseDestroy(cusparseHandle);
- cusolverDnDestroy(cusolverDnHandle);
- cusolverSpDestroy(cusolverSpHandle);
- cudnnHandle = null;
- cublasHandle = null;
- cusparseHandle = null;
-
- }
-
- /**
- * Clears all memory used by this {@link GPUContext}
- * Be careful to ensure that no memory is currently being used in the temporary memory before invoking this
- * If memory is being used between MLContext invocations, they are pointed to by a {@link GPUObject} instance
- * which would be part of the {@link MatrixObject}. The cleanup of that {@link MatrixObject} instance will
- * cause the memory associated with that block on the GPU to be freed up.
- * @throws DMLRuntimeException ?
- */
- public void clearMemory() throws DMLRuntimeException {
- clearTemporaryMemory();
- while (!allocatedGPUObjects.isEmpty()) {
- GPUObject o = allocatedGPUObjects.get(0);
- if (o.isDirty()){
- LOG.warn("Attempted to free GPU Memory when a block[" + o + "] is still on GPU memory, copying it back to host.");
- o.acquireHostRead();
- }
- o.clearData(true);
- }
- allocatedGPUObjects.clear();
- }
-
- /**
- * Clears up the memory used to optimize cudaMalloc/cudaFree calls
- */
- public void clearTemporaryMemory() {
- // To record the cuda block sizes needed by allocatedGPUObjects, others are cleared up.
- HashMap<Pointer, Long> tmpCudaBlockSizeMap = new HashMap<>();
- for (GPUObject o : allocatedGPUObjects) {
- if (o.isSparse()) {
- CSRPointer p = o.getSparseMatrixCudaPointer();
- if (p.rowPtr != null && cudaBlockSizeMap.containsKey(p.rowPtr)) {
- tmpCudaBlockSizeMap.put(p.rowPtr, cudaBlockSizeMap.get(p.rowPtr));
- }
- if (p.colInd != null && cudaBlockSizeMap.containsKey(p.colInd)) {
- tmpCudaBlockSizeMap.put(p.colInd, cudaBlockSizeMap.get(p.colInd));
- }
- if (p.val != null && cudaBlockSizeMap.containsKey(p.val)) {
- tmpCudaBlockSizeMap.put(p.val, cudaBlockSizeMap.get(p.val));
- }
-
- } else {
- Pointer p = o.getJcudaDenseMatrixPtr();
- tmpCudaBlockSizeMap.put(p, cudaBlockSizeMap.get(p));
- }
- }
-
- // garbage collect all temporarily allocated spaces
- for (LinkedList<Pointer> l : freeCUDASpaceMap.values()) {
- for (Pointer p : l) {
- cudaFreeHelper(p, true);
- }
- }
- cudaBlockSizeMap.clear();
- freeCUDASpaceMap.clear();
-
- // Restore only those entries for which there are still blocks on the GPU
- cudaBlockSizeMap.putAll(tmpCudaBlockSizeMap);
- }
-
- @Override
- public String toString() {
- return "GPUContext{" +
- "deviceNum=" + deviceNum +
- '}';
- }
+ /**
+ * @param o {@link GPUObject} instance to remove from the list of allocated GPU objects
+ * @see GPUContext#allocatedGPUObjects
+ * Records that a block is not used anymore
+ */
+ public void removeRecordedUsage(GPUObject o) {
+ allocatedGPUObjects.remove(o);
+ }
+
+ /**
+ * Gets the available memory on GPU that SystemML can use
+ *
+ * @return the available memory in bytes
+ */
+ public long getAvailableMemory() {
+ long free[] = { 0 };
+ long total[] = { 0 };
+ cudaMemGetInfo(free, total);
+ return (long) (free[0] * GPU_MEMORY_UTILIZATION_FACTOR);
+ }
+
+ /**
+ * Makes sure that GPU that SystemML is trying to use has the minimum compute capability needed
+ *
+ * @throws DMLRuntimeException if the compute capability is less than what is required
+ */
+ public void ensureComputeCapability() throws DMLRuntimeException {
+ int[] devices = { -1 };
+ cudaGetDeviceCount(devices);
+ if (devices[0] == -1) {
+ throw new DMLRuntimeException("Call to cudaGetDeviceCount returned 0 devices");
+ }
+ boolean isComputeCapable = true;
+ for (int i = 0; i < devices[0]; i++) {
+ cudaDeviceProp properties = GPUContextPool.getGPUProperties(i);
+ int major = properties.major;
+ int minor = properties.minor;
+ if (major < MAJOR_REQUIRED) {
+ isComputeCapable = false;
+ } else if (major == MAJOR_REQUIRED && minor < MINOR_REQUIRED) {
+ isComputeCapable = false;
+ }
+ }
+ if (!isComputeCapable) {
+ throw new DMLRuntimeException(
+ "One of the CUDA cards on the system has compute capability lower than " + MAJOR_REQUIRED + "."
+ + MINOR_REQUIRED);
+ }
+ }
+
+ public GPUObject createGPUObject(MatrixObject mo) {
+ return new GPUObject(this, mo);
+ }
+
+ /**
+ * Gets the device properties for the active GPU (set with cudaSetDevice())
+ *
+ * @return the device properties
+ * @throws DMLRuntimeException ?
+ */
+ public cudaDeviceProp getGPUProperties() throws DMLRuntimeException {
+ return GPUContextPool.getGPUProperties(deviceNum);
+ }
+
+ /**
+ * Gets the maximum number of threads per block for "active" GPU
+ *
+ * @return the maximum number of threads per block
+ * @throws DMLRuntimeException ?
+ */
+ public int getMaxThreadsPerBlock() throws DMLRuntimeException {
+ cudaDeviceProp deviceProps = getGPUProperties();
+ return deviceProps.maxThreadsPerBlock;
+ }
+
+ /**
+ * Gets the maximum number of blocks supported by the active cuda device
+ *
+ * @return the maximum number of blocks supported
+ * @throws DMLRuntimeException ?
+ */
+ public int getMaxBlocks() throws DMLRuntimeException {
+ cudaDeviceProp deviceProp = getGPUProperties();
+ return deviceProp.maxGridSize[0];
+ }
+
+ /**
+ * Gets the shared memory per block supported by the active cuda device
+ *
+ * @return the shared memory per block
+ * @throws DMLRuntimeException ?
+ */
+ public long getMaxSharedMemory() throws DMLRuntimeException {
+ cudaDeviceProp deviceProp = getGPUProperties();
+ return deviceProp.sharedMemPerBlock;
+ }
+
+ /**
+ * Gets the warp size supported by the active cuda device
+ *
+ * @return the warp size
+ * @throws DMLRuntimeException ?
+ */
+ public int getWarpSize() throws DMLRuntimeException {
+ cudaDeviceProp deviceProp = getGPUProperties();
+ return deviceProp.warpSize;
+ }
+
+ public cudnnHandle getCudnnHandle() {
+ return cudnnHandle;
+ }
+
+ public cublasHandle getCublasHandle() {
+ return cublasHandle;
+ }
+
+ public cusparseHandle getCusparseHandle() {
+ return cusparseHandle;
+ }
+
+ public cusolverDnHandle getCusolverDnHandle() {
+ return cusolverDnHandle;
+ }
+
+ public cusolverSpHandle getCusolverSpHandle() {
+ return cusolverSpHandle;
+ }
+
+ public JCudaKernels getKernels() {
+ return kernels;
+ }
+
+ /**
+ * Destroys this GPUContext object
+ *
+ * @throws DMLRuntimeException if error
+ */
+ public void destroy() throws DMLRuntimeException {
+ LOG.trace("GPU : this context was destroyed, this = " + this.toString());
+ clearMemory();
+ cudnnDestroy(cudnnHandle);
+ cublasDestroy(cublasHandle);
+ cusparseDestroy(cusparseHandle);
+ cusolverDnDestroy(cusolverDnHandle);
+ cusolverSpDestroy(cusolverSpHandle);
+ cudnnHandle = null;
+ cublasHandle = null;
+ cusparseHandle = null;
+
+ }
+
+ /**
+ * Clears all memory used by this {@link GPUContext}
+ * Be careful to ensure that no memory is currently being used in the temporary memory before invoking this
+ * If memory is being used between MLContext invocations, they are pointed to by a {@link GPUObject} instance
+ * which would be part of the {@link MatrixObject}. The cleanup of that {@link MatrixObject} instance will
+ * cause the memory associated with that block on the GPU to be freed up.
+ *
+ * @throws DMLRuntimeException ?
+ */
+ public void clearMemory() throws DMLRuntimeException {
+ clearTemporaryMemory();
+ while (!allocatedGPUObjects.isEmpty()) {
+ GPUObject o = allocatedGPUObjects.get(0);
+ if (o.isDirty()) {
+ LOG.warn("Attempted to free GPU Memory when a block[" + o
+ + "] is still on GPU memory, copying it back to host.");
+ o.acquireHostRead();
+ }
+ o.clearData(true);
+ }
+ allocatedGPUObjects.clear();
+ }
+
+ /**
+ * Clears up the memory used to optimize cudaMalloc/cudaFree calls
+ */
+ public void clearTemporaryMemory() {
+ // To record the cuda block sizes needed by allocatedGPUObjects, others are cleared up.
+ HashMap<Pointer, Long> tmpCudaBlockSizeMap = new HashMap<>();
+ for (GPUObject o : allocatedGPUObjects) {
+ if (o.isSparse()) {
+ CSRPointer p = o.getSparseMatrixCudaPointer();
+ if (p.rowPtr != null && cudaBlockSizeMap.containsKey(p.rowPtr)) {
+ tmpCudaBlockSizeMap.put(p.rowPtr, cudaBlockSizeMap.get(p.rowPtr));
+ }
+ if (p.colInd != null && cudaBlockSizeMap.containsKey(p.colInd)) {
+ tmpCudaBlockSizeMap.put(p.colInd, cudaBlockSizeMap.get(p.colInd));
+ }
+ if (p.val != null && cudaBlockSizeMap.containsKey(p.val)) {
+ tmpCudaBlockSizeMap.put(p.val, cudaBlockSizeMap.get(p.val));
+ }
+
+ } else {
+ Pointer p = o.getJcudaDenseMatrixPtr();
+ tmpCudaBlockSizeMap.put(p, cudaBlockSizeMap.get(p));
+ }
+ }
+
+ // garbage collect all temporarily allocated spaces
+ for (LinkedList<Pointer> l : freeCUDASpaceMap.values()) {
+ for (Pointer p : l) {
+ cudaFreeHelper(p, true);
+ }
+ }
+ cudaBlockSizeMap.clear();
+ freeCUDASpaceMap.clear();
+
+ // Restore only those entries for which there are still blocks on the GPU
+ cudaBlockSizeMap.putAll(tmpCudaBlockSizeMap);
+ }
+
+ @Override
+ public String toString() {
+ return "GPUContext{" + "deviceNum=" + deviceNum + '}';
+ }
+
+ /**
+ * Eviction policies for {@link GPUContext#evict(long)}
+ */
+ public enum EvictionPolicy {
+ LRU, LFU, MIN_EVICT
+ }
}
http://git-wip-us.apache.org/repos/asf/systemml/blob/f5871756/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContextPool.java
----------------------------------------------------------------------
diff --git a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContextPool.java b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContextPool.java
index 1d0b5c8..ac1c059 100644
--- a/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContextPool.java
+++ b/src/main/java/org/apache/sysml/runtime/instructions/gpu/context/GPUContextPool.java
@@ -23,7 +23,7 @@ import static jcuda.driver.JCudaDriver.cuInit;
import static jcuda.runtime.JCuda.cudaGetDeviceProperties;
import java.util.LinkedList;
-import java.util.Queue;
+import java.util.List;
import org.apache.commons.logging.Log;
import org.apache.commons.logging.LogFactory;
@@ -39,122 +39,152 @@ import jcuda.runtime.cudaDeviceProp;
public class GPUContextPool {
- protected static final Log LOG = LogFactory.getLog(GPUContextPool.class.getName());
-
- /** Maximum number of gpus to use, -1 for all */
- public static int PER_PROCESS_MAX_GPUS = -1;
-
- /** Whether cuda has been initialized */
- static boolean initialized = false;
-
- /** The total number of cuda devices on this machine */
- static int deviceCount = -1;
-
- /** Stores the cached deviceProperties */
- static cudaDeviceProp[] deviceProperties;
-
- /** Set of free GPUContexts */
- static Queue<GPUContext> freePool = new LinkedList<>();
-
- /**
- * Static initialization of the number of devices
- * Also sets behaviour for J{Cuda, Cudnn, Cublas, Cusparse} in case of error
- * Initializes the CUDA driver
- * All these need be done once, and not per GPU
- * @throws DMLRuntimeException ?
- */
- public synchronized static void initializeGPU() throws DMLRuntimeException {
- GPUContext.LOG.info("Initializing CUDA");
- long start = System.nanoTime();
- JCuda.setExceptionsEnabled(true);
- JCudnn.setExceptionsEnabled(true);
- JCublas2.setExceptionsEnabled(true);
- JCusparse.setExceptionsEnabled(true);
- JCudaDriver.setExceptionsEnabled(true);
- cuInit(0); // Initialize the driver
-
- int deviceCountArray[] = {0};
- cuDeviceGetCount(deviceCountArray); // Obtain the number of devices
- deviceCount = deviceCountArray[0];
- deviceProperties = new cudaDeviceProp[deviceCount];
-
- if (PER_PROCESS_MAX_GPUS > 0)
- deviceCount = Math.min(PER_PROCESS_MAX_GPUS, deviceCount);
-
- // Initialize the list of devices
- for (int i = 0; i < deviceCount; i++) {
- cudaDeviceProp properties = new cudaDeviceProp();
- cudaGetDeviceProperties(properties, i);
- deviceProperties[i] = properties;
- }
-
- // Initialize the pool of GPUContexts
- for (int i=0; i<deviceCount; i++){
- GPUContext gCtx = new GPUContext(i);
- freePool.add(gCtx);
- }
-
- GPUContext.LOG.info("Total number of GPUs on the machine: " + deviceCount);
- //int[] device = {-1};
- //cudaGetDevice(device);
- //cudaDeviceProp prop = getGPUProperties(device[0]);
- //int maxBlocks = prop.maxGridSize[0];
- //int maxThreadsPerBlock = prop.maxThreadsPerBlock;
- //long sharedMemPerBlock = prop.sharedMemPerBlock;
- //LOG.debug("Active CUDA device number : " + device[0]);
- //LOG.debug("Max Blocks/Threads/SharedMem on active device: " + maxBlocks + "/" + maxThreadsPerBlock + "/" + sharedMemPerBlock);
- initialized = true;
- GPUStatistics.cudaInitTime = System.nanoTime() - start;
- }
-
- /**
- * Gets an initialized GPUContext from a pool of GPUContexts, each linked to a GPU
- * @return null if not more GPUContexts in pool, a valid GPUContext otherwise
- * @throws DMLRuntimeException ?
- */
- public static synchronized GPUContext getFromPool() throws DMLRuntimeException {
- if (!initialized) initializeGPU();
- GPUContext gCtx = freePool.poll();
- LOG.trace("GPU : got GPUContext (" + gCtx + ") from freePool. New sizes - FreePool[" + freePool.size() + "]");
- return gCtx;
- }
-
- /**
- * Get the number of free GPUContexts
- * @return number of free GPUContexts
- */
- public static synchronized int getAvailableCount() {
- return freePool.size();
- }
-
- /**
- * Gets the device properties
- * @param device the device number (on a machine with more than 1 GPU)
- * @return the device properties
- * @throws DMLRuntimeException if there is problem initializing the GPUContexts
- */
- static cudaDeviceProp getGPUProperties(int device) throws DMLRuntimeException {
- // do once - initialization of GPU
- if (!initialized) initializeGPU();
- return deviceProperties[device];
- }
-
- public static int getDeviceCount() throws DMLRuntimeException {
- if (!initialized) initializeGPU();
- return deviceCount;
- }
-
- /**
- * Returns a {@link GPUContext} back to the pool of {@link GPUContext}s
- * @param gCtx the GPUContext instance to return. If null, nothing happens
- * @throws DMLRuntimeException if error
- */
- public static synchronized void returnToPool(GPUContext gCtx) throws DMLRuntimeException {
- if (gCtx == null)
- return;
- freePool.add(gCtx);
- LOG.trace("GPU : returned GPUContext (" + gCtx + ") to freePool. New sizes - FreePool[" + freePool.size() + "]");
-
- }
+ protected static final Log LOG = LogFactory.getLog(GPUContextPool.class.getName());
+
+ /**
+ * Maximum number of gpus to use, -1 for all
+ */
+ public static int PER_PROCESS_MAX_GPUS = -1;
+
+ /**
+ * Whether cuda has been initialized
+ */
+ static boolean initialized = false;
+
+ /**
+ * The total number of cuda devices on this machine
+ */
+ static int deviceCount = -1;
+
+ /**
+ * Stores the cached deviceProperties
+ */
+ static cudaDeviceProp[] deviceProperties;
+
+ /**
+ * Set of free GPUContexts
+ */
+ static List<GPUContext> pool = new LinkedList<>();
+
+ /**
+ * Whether the pool of GPUs is reserved or not
+ */
+ static boolean reserved = false;
+
+ /**
+ * Static initialization of the number of devices
+ * Also sets behaviour for J{Cuda, Cudnn, Cublas, Cusparse} in case of error
+ * Initializes the CUDA driver
+ * All these need be done once, and not per GPU
+ *
+ * @throws DMLRuntimeException ?
+ */
+ public synchronized static void initializeGPU() throws DMLRuntimeException {
+ GPUContext.LOG.info("Initializing CUDA");
+ long start = System.nanoTime();
+ JCuda.setExceptionsEnabled(true);
+ JCudnn.setExceptionsEnabled(true);
+ JCublas2.setExceptionsEnabled(true);
+ JCusparse.setExceptionsEnabled(true);
+ JCudaDriver.setExceptionsEnabled(true);
+ cuInit(0); // Initialize the driver
+
+ int deviceCountArray[] = { 0 };
+ cuDeviceGetCount(deviceCountArray); // Obtain the number of devices
+ deviceCount = deviceCountArray[0];
+ deviceProperties = new cudaDeviceProp[deviceCount];
+
+ if (PER_PROCESS_MAX_GPUS > 0)
+ deviceCount = Math.min(PER_PROCESS_MAX_GPUS, deviceCount);
+
+ // Initialize the list of devices
+ for (int i = 0; i < deviceCount; i++) {
+ cudaDeviceProp properties = new cudaDeviceProp();
+ cudaGetDeviceProperties(properties, i);
+ deviceProperties[i] = properties;
+ }
+
+ // Initialize the pool of GPUContexts
+ for (int i = 0; i < deviceCount; i++) {
+ GPUContext gCtx = new GPUContext(i);
+ pool.add(gCtx);
+ }
+
+ GPUContext.LOG.info("Total number of GPUs on the machine: " + deviceCount);
+ //int[] device = {-1};
+ //cudaGetDevice(device);
+ //cudaDeviceProp prop = getGPUProperties(device[0]);
+ //int maxBlocks = prop.maxGridSize[0];
+ //int maxThreadsPerBlock = prop.maxThreadsPerBlock;
+ //long sharedMemPerBlock = prop.sharedMemPerBlock;
+ //LOG.debug("Active CUDA device number : " + device[0]);
+ //LOG.debug("Max Blocks/Threads/SharedMem on active device: " + maxBlocks + "/" + maxThreadsPerBlock + "/" + sharedMemPerBlock);
+ initialized = true;
+ GPUStatistics.cudaInitTime = System.nanoTime() - start;
+ }
+
+ /**
+ * Reserves and gets an initialized list of GPUContexts
+ *
+ * @return null if no GPUContexts in pool, otherwise a valid list of GPUContext
+ * @throws DMLRuntimeException ?
+ */
+ public static synchronized List<GPUContext> reserveAllGPUContexts() throws DMLRuntimeException {
+ if (reserved)
+ throw new DMLRuntimeException("Trying to re-reserve GPUs");
+ if (!initialized)
+ initializeGPU();
+ reserved = true;
+ LOG.trace("GPU : Reserved all GPUs");
+ return pool;
+ }
+
+ /**
+ * Get the number of free GPUContexts
+ *
+ * @return number of free GPUContexts
+ */
+ public static synchronized int getAvailableCount() {
+ return pool.size();
+ }
+
+ /**
+ * Gets the device properties
+ *
+ * @param device the device number (on a machine with more than 1 GPU)
+ * @return the device properties
+ * @throws DMLRuntimeException if there is problem initializing the GPUContexts
+ */
+ static cudaDeviceProp getGPUProperties(int device) throws DMLRuntimeException {
+ // do once - initialization of GPU
+ if (!initialized)
+ initializeGPU();
+ return deviceProperties[device];
+ }
+
+ /**
+ * Number of available devices on this machine
+ *
+ * @return number of available GPUs on this machine
+ * @throws DMLRuntimeException if error
+ */
+ public static int getDeviceCount() throws DMLRuntimeException {
+ if (!initialized)
+ initializeGPU();
+ return deviceCount;
+ }
+
+ /**
+ * Unreserves all GPUContexts
+ *
+ * @throws DMLRuntimeException if error
+ */
+ public static synchronized void freeAllGPUContexts() throws DMLRuntimeException {
+ if (!reserved)
+ throw new DMLRuntimeException("Trying to free unreserved GPUs");
+ reserved = false;
+ LOG.trace("GPU : Unreserved all GPUs");
+
+ }
}