From 1ad66a967abe9ac022ffa593e9968a7d24a7060b Mon Sep 17 00:00:00 2001 From: Yang Bo Date: Fri, 2 Feb 2018 18:40:11 +0800 Subject: [PATCH 1/2] Generate correct OpenCL C identifier --- .../scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala b/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala index 448d1ef6..35a29445 100644 --- a/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala +++ b/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala @@ -48,8 +48,8 @@ object OpenCLKernelBuilder { def freshName(prefix: String): String = { val encodedPrefix = prefix.map { - case c if c.isLetterOrDigit => c - case _ => '_' + case c if (c >= 'a' && c <= 'z') || (c >= 'A' && c <= 'Z') || (c >= '0' && c <= '9') => c + case _ => '_' } val name = raw"""${encodedPrefix}_${seed}""" seed += 1 From a52c185a8ad20f7478c755b4a777f1e80b25d72c Mon Sep 17 00:00:00 2001 From: Yang Bo Date: Fri, 2 Feb 2018 18:41:16 +0800 Subject: [PATCH 2/2] Test and fix Tensor.translate --- .../compute/OpenCLKernelBuilder.scala | 14 ++++--- .../com/thoughtworks/compute/Tensors.scala | 24 ++++++------ .../thoughtworks/compute/TensorsSpec.scala | 38 +++++++++++++++++-- 3 files changed, 56 insertions(+), 20 deletions(-) diff --git a/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala b/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala index 35a29445..8093f436 100644 --- a/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala +++ b/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala @@ -234,14 +234,18 @@ trait OpenCLKernelBuilder extends FloatArrays { if matrix.getEntry(y, x) != 0.0 } yield { if (x < originalShape.length) { - // TODO: check boundary - fast"get_global_id($x) * ${matrix.getEntry(y, x)}" + matrix.getEntry(y, x) match { + case 1.0 => + fast"get_global_id($x)" + case scale => + fast"get_global_id($x) * $scale" + } } else { fast"${matrix.getEntry(y, x)}" } } val indexId = freshName("index") - indexId -> fast"size_t $indexId = ${products.mkFastring(" + ")}" + indexId -> fast"size_t $indexId = ${products.mkFastring(" + ")};\n" }).unzip // fast" @@ -258,7 +262,7 @@ trait OpenCLKernelBuilder extends FloatArrays { fast"[$i]" }.mkFastring}" localDefinitions += fastraw""" - const ${elementType.typeSymbol.typeCode} $termId = (${bounds.mkFastring(" && ")}) ? $paddingCode : $dereferenceCode; + const ${elementType.typeSymbol.typeCode} $termId = (${bounds.mkFastring(" && ")}) ? $dereferenceCode : $paddingCode; """ elementType.factory.newInstance(termId, elementType.typeSymbol.typeCode).asInstanceOf[Element] } @@ -282,7 +286,7 @@ trait OpenCLKernelBuilder extends FloatArrays { val shape: Array[Int] def transform(matrix: RealMatrix): ThisTerm = { - if (matrix.getColumnDimension != shape.length) { + if (matrix.getRowDimension != shape.length) { throw new IllegalArgumentException } arrayViewFactory.newInstance(elementType, matrix, padding, shape, termCode, typeCode).asInstanceOf[ThisTerm] diff --git a/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala b/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala index bc9c4eec..05208078 100644 --- a/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala +++ b/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala @@ -73,9 +73,9 @@ trait Tensors extends OpenCL { } } object Tensor { - def fill(value: Float, shape0: Array[Int]) = { + def fill(value: Float, shape0: Array[Int], padding: Float = 0.0f) = { new InlineTensor { - val padding: Float = value + val padding: Float = padding val shape: shape0.type = shape0 val closure: trees.FloatTerm = float.literal(value) } @@ -112,18 +112,16 @@ trait Tensors extends OpenCL { val padding: Float = thisTensor.padding } case _ => + val newMatrix = MatrixUtils.createRealMatrix(shape.length, shape.length + 1) + for (i <- offset.indices) { + newMatrix.setEntry(i, i, 1.0) + newMatrix.setEntry(i, newMatrix.getColumnDimension - 1, offset(i)) + } new TransformedTensor { - val checkpoint: Tensor = thisTensor - val shape: Array[Int] = newShape + def checkpoint: Tensor = thisTensor + def shape: Array[Int] = newShape // val debuggingInformation: Implicitly[DebuggingInformation] = debuggingInformation0 - val matrix: RealMatrix = { - val newMatrix = MatrixUtils.createRealMatrix(shape.length, shape.length + 1) - for (i <- offset.indices) { - newMatrix.setEntry(i, i, 1.0) - newMatrix.setEntry(i, newMatrix.getColumnDimension - 1, offset(i)) - } - newMatrix - } + def matrix: RealMatrix = newMatrix def padding: Float = checkpoint.padding } @@ -212,6 +210,8 @@ trait Tensors extends OpenCL { """ } + sourceCode.foreach(print) + val program = createProgramWithSource(sourceCode) program.build() diff --git a/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala b/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala index 123f7b10..d43c80da 100644 --- a/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala +++ b/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala @@ -6,7 +6,9 @@ import com.thoughtworks.feature.Factory import TensorsSpec._ import com.thoughtworks.future._ import com.thoughtworks.raii.asynchronous._ +import org.lwjgl.BufferUtils import org.lwjgl.opencl.CLCapabilities +import org.lwjgl.system.MemoryUtil import scalaz.syntax.all._ import scala.language.existentials @@ -30,9 +32,9 @@ class TensorsSpec extends AsyncFreeSpec with Matchers { doTensors.flatMap { tensors => val shape = Array(2, 3, 5) val element = 42.0f - val zeros = tensors.Tensor.fill(element, shape) + val filled = tensors.Tensor.fill(element, shape) for { - pendingBuffer <- zeros.enqueue + pendingBuffer <- filled.enqueue floatBuffer <- pendingBuffer.toHostBuffer } yield { for (i <- 0 until floatBuffer.capacity()) { @@ -41,13 +43,43 @@ class TensorsSpec extends AsyncFreeSpec with Matchers { floatBuffer.position() should be(0) floatBuffer.limit() should be(shape.product) floatBuffer.capacity() should be(shape.product) - tensors.kernelCache.getIfPresent(zeros.closure) should not be null + tensors.kernelCache.getIfPresent(filled.closure) should not be null val zeros2 = tensors.Tensor.fill(element, shape) tensors.kernelCache.getIfPresent(zeros2.closure) should not be null } } }.run.toScalaFuture + "translate" in { + doTensors.flatMap { tensors => + val shape = Array(2, 3, 5) + val element = 42.0f + val translated = tensors.Tensor.fill(element, shape).translate(Array(1, 2, -3)) + for { + pendingBuffer <- translated.enqueue + floatBuffer <- pendingBuffer.toHostBuffer + } yield { + + floatBuffer.position() should be(0) + + val array = Array.ofDim[Float](shape.product) + floatBuffer.get(array) + val array3d = array.grouped(shape(2)).grouped(shape(1)) + for ((xi, i) <- array3d.zipWithIndex; (xij, j) <- xi.zipWithIndex; (xijk, k) <- xij.view.zipWithIndex) { + if (2 - i > 1 && 3 - j > 2 && k >= 3) { + xijk should be(element) + } else { + xijk should be(0.0f) + } + } + + floatBuffer.limit() should be(shape.product) + floatBuffer.capacity() should be(shape.product) + + } + } + }.run.toScalaFuture + "convolution" ignore { doTensors.flatMap { tensors => import tensors.Tensor