diff --git a/Expressions/src/main/scala/com/thoughtworks/compute/Expressions.scala b/Expressions/src/main/scala/com/thoughtworks/compute/Expressions.scala index 8039bc9f..6eaa8703 100644 --- a/Expressions/src/main/scala/com/thoughtworks/compute/Expressions.scala +++ b/Expressions/src/main/scala/com/thoughtworks/compute/Expressions.scala @@ -25,7 +25,7 @@ trait Expressions { type Term <: TermApi - protected trait TypeApi extends ExpressionApi {} + protected trait TypeApi extends ExpressionApi type Type <: TypeApi diff --git a/OpenCL/build.sbt b/OpenCL/build.sbt index a6d99e27..9d7f2862 100644 --- a/OpenCL/build.sbt +++ b/OpenCL/build.sbt @@ -13,9 +13,7 @@ val lwjglNatives: String = { libraryDependencies += "org.lwjgl" % "lwjgl-opencl" % "3.1.5" -libraryDependencies += "org.lwjgl" % "lwjgl" % "3.1.5" - -libraryDependencies += "org.lwjgl" % "lwjgl" % "3.1.5" % Test classifier lwjglNatives +libraryDependencies += ("org.lwjgl" % "lwjgl" % "3.1.5" % Test).classifier(lwjglNatives).jar() libraryDependencies += "com.thoughtworks.raii" %% "asynchronous" % "3.0.0-M8" diff --git a/OpenCL/src/main/scala/com/thoughtworks/compute/OpenCL.scala b/OpenCL/src/main/scala/com/thoughtworks/compute/OpenCL.scala index 20722c5f..3442597d 100644 --- a/OpenCL/src/main/scala/com/thoughtworks/compute/OpenCL.scala +++ b/OpenCL/src/main/scala/com/thoughtworks/compute/OpenCL.scala @@ -588,13 +588,13 @@ object OpenCL { * will cause memory access error. * */ - final def toHostBuffer(implicit witnessOwner: Witness.Aux[Owner], - memory: Memory[Element]): Do[memory.HostBuffer] = { + final def toHostBuffer(preconditionEvents: Event[Owner]*)(implicit witnessOwner: Witness.Aux[Owner], + memory: Memory[Element]): Do[memory.HostBuffer] = { Do(TryT(ResourceT(UnitContinuation.delay { val hostBuffer = memory.allocate(length) Resource(value = Success(hostBuffer), release = UnitContinuation.delay { memory.free(hostBuffer) }) }))).flatMap { hostBuffer => - enqueueReadBuffer[memory.HostBuffer](hostBuffer)(witnessOwner, memory) + enqueueReadBuffer[memory.HostBuffer](hostBuffer, preconditionEvents: _*)(witnessOwner, memory) .flatMap { event => Do.garbageCollected(event.waitForComplete()) } diff --git a/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala b/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala index b1f2a2de..448d1ef6 100644 --- a/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala +++ b/OpenCLKernelBuilder/src/main/scala/com/thoughtworks/compute/OpenCLKernelBuilder.scala @@ -42,7 +42,7 @@ object OpenCLKernelBuilder { final case class ClTypeSymbol(firstDefinition: ClTypeDefinition, typeCode: ClTypeCode) - final class GlobalContext { + final class GlobalContext extends Fastring { private var seed = 0 @@ -56,8 +56,8 @@ object OpenCLKernelBuilder { name } - val globalDeclarations = mutable.Buffer.empty[Fastring] - val globalDefinitions = mutable.Buffer.empty[Fastring] + protected[OpenCLKernelBuilder] val globalDeclarations = mutable.Buffer.empty[Fastring] + protected[OpenCLKernelBuilder] val globalDefinitions = mutable.Buffer.empty[Fastring] private val typeSymbolCache = mutable.HashMap.empty[ClTypeDefinition, ClTypeSymbol] val floatSymbol = cachedSymbol(FloatDefinition) @@ -73,6 +73,10 @@ object OpenCLKernelBuilder { typeSymbol } + def foreach[U](f: String => U): Unit = { + globalDeclarations.foreach(_.foreach(f)) + globalDefinitions.foreach(_.foreach(f)) + } } } @@ -98,7 +102,8 @@ trait OpenCLKernelBuilder extends FloatArrays { val (outputParameters, outputAssignments) = outputs.map { output => val outputTermCode = output.termCode val outputTypeCode = output.typeCode - val outputParameter = fast"global $outputTypeCode *output_$outputTermCode" + val outputId = freshName("output") + val outputParameter = fast"global $outputTypeCode *$outputId" def outputIndex(dimension: Int): Fastring = { if (dimension == 0) { fast"get_global_id(0)" @@ -108,11 +113,11 @@ trait OpenCLKernelBuilder extends FloatArrays { } val index = outputIndex(numberOfDimensions - 1) - val outputAssignment = fast"output_$outputTermCode[$index] = $outputTermCode;\n" + val outputAssignment = fast"$outputId[$index] = $outputTermCode;\n" (outputParameter, outputAssignment) }.unzip fastraw""" - kernel void $functionName(${parameterDeclarations.mkFastring(", ")}, ${outputParameters.mkFastring(", ")}) { + kernel void $functionName(${(parameterDeclarations.view ++ outputParameters).mkFastring(", ")}) { ${localDefinitions.mkFastring} ${outputAssignments.mkFastring} } diff --git a/OpenCLKernelBuilder/src/test/scala/com/thoughtworks/compute/OpenCLKernelBuilderSpec.scala b/OpenCLKernelBuilder/src/test/scala/com/thoughtworks/compute/OpenCLKernelBuilderSpec.scala index 22116476..12a424c3 100644 --- a/OpenCLKernelBuilder/src/test/scala/com/thoughtworks/compute/OpenCLKernelBuilderSpec.scala +++ b/OpenCLKernelBuilder/src/test/scala/com/thoughtworks/compute/OpenCLKernelBuilderSpec.scala @@ -63,8 +63,7 @@ class OpenCLKernelBuilderSpec extends FreeSpec with Matchers { ), Seq(f.tree.export(openCLFunctionContext, map))) - globalContext.globalDeclarations.foreach(print) - globalContext.globalDefinitions.foreach(print) + globalContext.foreach(print) sourceCode.foreach(print) // TODO: Convert this example to a test case diff --git a/Tensors/build.sbt b/Tensors/build.sbt index 90657c6d..27ce38ba 100644 --- a/Tensors/build.sbt +++ b/Tensors/build.sbt @@ -1,3 +1,22 @@ scalacOptions += "-Ypartial-unification" libraryDependencies += "com.google.guava" % "guava" % "23.6-jre" + +libraryDependencies += "org.scalatest" %% "scalatest" % "3.0.4" % Test + +val lwjglNatives: String = { + import scala.util.Properties._ + if (isMac) { + "natives-macos" + } else if (isLinux) { + "natives-linux" + } else if (isWin) { + "natives-windows" + } else { + throw new MessageOnlyException(s"lwjgl does not support $osName") + } +} + +libraryDependencies += ("org.lwjgl" % "lwjgl" % "3.1.5" % Test).classifier(lwjglNatives).jar() + +fork in Test := true diff --git a/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala b/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala index 1fe601dd..5368d490 100644 --- a/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala +++ b/Tensors/src/main/scala/com/thoughtworks/compute/Tensors.scala @@ -65,6 +65,19 @@ trait Tensors extends OpenCL { def event: Event def buffer: DeviceBuffer[Float] + + def toHostBuffer = { + buffer.toHostBuffer(event) + } + } + object Tensor { + def fill(value: Float, shape0: Array[Int]) = { + new InlineTensor { + val padding: Float = value + val shape: shape0.type = shape0 + val closure: trees.FloatTerm = float.literal(value) + } + } } sealed trait Tensor { thisTensor => @@ -75,6 +88,7 @@ trait Tensors extends OpenCL { def closure: ValueTerm + // TODO: rename to make buffer def enqueue: Do[PendingBuffer] def padding: Float @@ -96,7 +110,7 @@ trait Tensors extends OpenCL { }) } - protected val kernelCache: Cache[ValueTerm, CompiledKernel] = kernelCacheBuilder.build() + protected[compute] val kernelCache: Cache[ValueTerm, CompiledKernel] = kernelCacheBuilder.build() protected implicit val executionContext: ExecutionContext @@ -132,22 +146,21 @@ trait Tensors extends OpenCL { def call(): CompiledKernel = { val alphConversionContext = new AlphaConversionContext - val convertedTerm = closure.tree.alphaConversion(alphConversionContext).asInstanceOf[ValueTerm] + val convertedTree = closure.tree.alphaConversion(alphConversionContext) val sourceCode = { val globalContext = new GlobalContext val functionContext = Factory[OpenCLKernelBuilder].newInstance(globalContext) val exportContext = new ExportContext - val kernelBody = convertedTerm.tree.export(functionContext, exportContext) + val kernelBody = convertedTree.export(functionContext, exportContext).asInstanceOf[functionContext.Term] val kernelParameters = upvalues(closure.tree).map { upvalue: Parameter => exportContext.get(alphConversionContext.get(upvalue)).asInstanceOf[functionContext.Term] } fastraw""" - ${globalContext.globalDeclarations} - ${globalContext.globalDefinitions} - ${functionContext.generateKernelSourceCode("kernel", shape.length, kernelParameters, Seq(kernelBody))} + $globalContext + ${functionContext.generateKernelSourceCode("jit_kernel", shape.length, kernelParameters, Seq(kernelBody))} """ } @@ -188,7 +201,7 @@ trait Tensors extends OpenCL { } } - kernelCache.put(convertedTerm, compiledKernel) + kernelCache.put(float.factory.newInstance(convertedTree.asInstanceOf[float.Tree]), compiledKernel) compiledKernel } } diff --git a/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala b/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala index 8ee76bd0..a97d0677 100644 --- a/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala +++ b/Tensors/src/test/scala/com/thoughtworks/compute/TensorsSpec.scala @@ -1,27 +1,49 @@ package com.thoughtworks.compute -import java.nio.ByteBuffer +import java.nio.{ByteBuffer, FloatBuffer} import com.thoughtworks.feature.Factory import TensorsSpec._ +import com.thoughtworks.future._ import com.thoughtworks.raii.asynchronous._ import org.lwjgl.opencl.CLCapabilities +import scalaz.syntax.all._ import scala.language.existentials +import org.scalatest._ /** * @author 杨博 (Yang Bo) */ -class TensorsSpec { - private val hyperparameters = +class TensorsSpec extends AsyncFreeSpec with Matchers { + private val tensors: Tensors = Factory[ OpenCL.GlobalExecutionContext with OpenCL.UseAllDevices with OpenCL.UseFirstPlatform with OpenCL.CommandQueuePool with Tensors] .newInstance( handleOpenCLNotification = handleOpenCLNotification, numberOfCommandQueuesForDevice = { (deviceId: Long, capabilities: CLCapabilities) => - 1 + 5 } ) + + "create a tensor of a constant" in { + val shape = Array(2, 3, 5) + val element = 42.0f + val zeros = tensors.Tensor.fill(element, shape) + + for { + pendingBuffer <- zeros.enqueue + floatBuffer <- pendingBuffer.toHostBuffer + } yield { + for (i <- 0 until floatBuffer.capacity()) { + floatBuffer.get(i) should be(element) + } + floatBuffer.position() should be(0) + floatBuffer.limit() should be(shape.product) + floatBuffer.capacity() should be(shape.product) + } + }.run.toScalaFuture + } object TensorsSpec { diff --git a/Trees/src/main/scala/com/thoughtworks/compute/Trees.scala b/Trees/src/main/scala/com/thoughtworks/compute/Trees.scala index 50aaf43a..0372bc56 100644 --- a/Trees/src/main/scala/com/thoughtworks/compute/Trees.scala +++ b/Trees/src/main/scala/com/thoughtworks/compute/Trees.scala @@ -18,25 +18,58 @@ trait Trees extends Expressions { final class StructuralComparisonContext extends IdentityHashMap[TreeApi, TreeApi] + private def childHashCode(child: Any, context: HashCodeContext): Int = { + child match { + case childTree: TreeApi => + childTree.structuralHashCode(context) + case childArray: Array[_] => + arrayHashCode(childArray, context) + case _ => + child.## + } + } + + private def arrayHashCode[@specialized A](childArray: Array[A], context: HashCodeContext): Int = { + val length = childArray.length + if (length == 0) { + MurmurHash3.arraySeed + } else { + val last = length - 1 + @tailrec + def arrayLoop(h: Int, i: Int): Int = { + if (i < last) { + arrayLoop(h = MurmurHash3.mix(h, childHashCode(childArray(i), context)), i = i + 1) + } else { + MurmurHash3.finalizeHash(MurmurHash3.mixLast(h, childHashCode(childArray(i), context)), length) + } + } + arrayLoop(MurmurHash3.arraySeed, 0) + } + } + trait Operator extends TreeApi { thisOperator => def structuralHashCode(context: HashCodeContext): Int = { val productArity: Int = this.productArity - @tailrec - def generateHashCode(h: Int = productPrefix.hashCode, i: Int = 0): Int = { - if (i < productArity) { - val childHashCode = productElement(i) match { - case childTree: TreeApi => - childTree.structuralHashCode(context) - case leaf => - leaf.## + if (productArity == 0) { + productPrefix.## + } else { + context.asScala.getOrElseUpdate( + this, { + val last = productArity - 1 + @tailrec + def loop(h: Int = productPrefix.hashCode, i: Int = 0): Int = { + if (i < last) { + loop(h = MurmurHash3.mix(h, childHashCode(productElement(i), context)), i = i + 1) + } else { + MurmurHash3.finalizeHash(MurmurHash3.mixLast(h, childHashCode(productElement(i), context)), + productArity) + } + } + loop() } - generateHashCode(h = MurmurHash3.mix(h, childHashCode), i = i + 1) - } else { - MurmurHash3.finalizeHash(h, productArity) - } + ) } - context.asScala.getOrElseUpdate(this, generateHashCode()) } def isSameStructure(that: TreeApi, map: StructuralComparisonContext): Boolean = { @@ -80,14 +113,6 @@ trait Trees extends Expressions { val id: Any - def structuralHashCode(context: HashCodeContext): Int = { - context.asScala.getOrElseUpdate(this, { - val newId = context.numberOfParameters - context.numberOfParameters = newId + 1 - newId - }) - } - def isSameStructure(that: TreeApi, map: StructuralComparisonContext): Boolean = { map.get(this) match { case null => @@ -122,20 +147,29 @@ trait Trees extends Expressions { final class ExportContext extends IdentityHashMap[TreeApi, Any] - protected trait TermApi extends super.TermApi { thisTree: Term => + protected trait ExpressionApi extends super.ExpressionApi { thisExpression => + type Tree = TreeApi { + type TermIn[C <: Category] = thisExpression.TermIn[C] + } + } + + protected trait TermApi extends ExpressionApi with super.TermApi { thisTree: Term => + + def alphaConversion: ThisTerm + def in(foreignCategory: Category): TermIn[foreignCategory.type] = { tree.export(foreignCategory, new ExportContext) } - type Tree = TreeApi { - type TermIn[C <: Category] = thisTree.TermIn[C] - } val tree: Tree } type Term <: TermApi + protected trait TypeApi extends ExpressionApi with super.TypeApi + type Type <: TypeApi + } object Trees { @@ -146,6 +180,11 @@ object Trees { trait ValueTrees extends Values with Trees { protected trait ValueTypeApi extends super.ValueTypeApi { + + override def equals(that: scala.Any): Boolean = { + that != null && that.getClass == this.getClass + } + def in(foreignCategory: Category): TypeIn[foreignCategory.type] def factory: Factory1[TreeApi { type TermIn[C <: Category] = ThisTerm#TermIn[C] }, ThisTerm] @@ -157,6 +196,10 @@ object Trees { protected trait ValueTermApi extends TermApi with super.ValueTermApi { thisValue: ValueTerm => def factory: Factory1[TreeApi { type TermIn[C <: Category] = thisValue.TermIn[C] }, ThisTerm] + + def alphaConversion: ThisTerm = { + factory.newInstance(tree.alphaConversion(new AlphaConversionContext).asInstanceOf[Tree]) + } } type ValueTerm <: (Term with Any) with ValueTermApi @@ -204,6 +247,14 @@ object Trees { final case class FloatParameter(id: Any) extends TreeApi with Parameter { thisParameter => type TermIn[C <: Category] = C#FloatTerm + def structuralHashCode(context: HashCodeContext): Int = { + context.asScala.getOrElseUpdate(this, { + val newId = context.numberOfParameters + context.numberOfParameters = newId + 1 + newId + }) + } + def export(foreignCategory: Category, map: ExportContext): foreignCategory.FloatTerm = { map.asScala .getOrElseUpdate(this, foreignCategory.float.parameter(id)) @@ -235,7 +286,7 @@ object Trees { } } - protected trait FloatTypeApi extends super.FloatTypeApi with FloatExpressionApi { + protected trait FloatTypeApi extends ValueTypeApi with super.FloatTypeApi with FloatExpressionApi { def in(foreignCategory: Category): TypeIn[foreignCategory.type] = { foreignCategory.float } @@ -250,6 +301,10 @@ object Trees { @inject def factory: Factory1[TreeApi { type TermIn[C <: Category] = ThisTerm#TermIn[C] }, ThisTerm] + + override def hashCode(): Int = { + "float".## + } } type FloatType <: (ValueType with Any) with FloatTypeApi @@ -321,6 +376,13 @@ object Trees { protected trait ArrayTermApi extends super.ArrayTermApi with TermApi { thisArray: ArrayTerm => + def alphaConversion: ThisTerm = { + array + .factory[Element] + .newInstance(tree.alphaConversion(new AlphaConversionContext).asInstanceOf[Tree], valueFactory) + .asInstanceOf[ThisTerm] + } + val valueFactory: Factory1[TreeApi { type TermIn[C <: Category] = thisArray.Element#TermIn[C] }, @@ -333,7 +395,7 @@ object Trees { def transform(matrix: RealMatrix): ThisTerm = { val translatedTree = Transform[Element](tree, matrix) array - .parameterFactory[Element] + .factory[Element] .newInstance( translatedTree, valueFactory @@ -344,7 +406,7 @@ object Trees { def translate(offset: Array[Int]): ThisTerm = { val translatedTree = Translate[Element](tree, offset) array - .parameterFactory[Element] + .factory[Element] .newInstance( translatedTree, valueFactory @@ -393,7 +455,7 @@ object Trees { val fillTree = Fill[thisValue.ThisTerm]( tree.asInstanceOf[TreeApi { type TermIn[C <: Category] = thisValue.ThisTerm#TermIn[C] }]) array - .parameterFactory[ThisTerm] + .factory[ThisTerm] .newInstance( fillTree, thisValue.factory @@ -415,10 +477,32 @@ object Trees { shape: Array[Int]) extends TreeApi with Parameter { thisParameter => + type TermIn[C <: Category] = C#ArrayTerm { type Element = elementType.TermIn[C] } + def structuralHashCode(context: HashCodeContext): Int = { + context.asScala.getOrElseUpdate( + this, { + val h = context.numberOfParameters + context.numberOfParameters = h + 1 + + MurmurHash3.finalizeHash( + MurmurHash3.mixLast( + MurmurHash3.mix( + MurmurHash3.mix(h, elementType.##), + padding.## + ), + MurmurHash3.arrayHash(shape) + ), + 3 + ) + } + ) + + } + def export(foreignCategory: Category, map: ExportContext): TermIn[foreignCategory.type] = { map.asScala .getOrElseUpdate( @@ -446,7 +530,7 @@ object Trees { protected trait ArrayCompanionApi extends super.ArrayCompanionApi { - @inject def parameterFactory[LocalElement <: ValueTerm] + @inject def factory[LocalElement <: ValueTerm] : Factory2[ArrayTree[LocalElement], Factory1[TreeApi { type TermIn[C <: Category] = LocalElement#TermIn[C] @@ -464,7 +548,7 @@ object Trees { } = { val parameterTree = ArrayParameter[Padding, elementType.type](id, elementType, padding, shape) array - .parameterFactory[elementType.ThisTerm] + .factory[elementType.ThisTerm] .newInstance( parameterTree.asInstanceOf[ArrayTree[elementType.ThisTerm]], elementType.factory diff --git a/Trees/src/test/scala/com/thoughtworks/compute/TreesSpec.scala b/Trees/src/test/scala/com/thoughtworks/compute/TreesSpec.scala new file mode 100644 index 00000000..b94f775b --- /dev/null +++ b/Trees/src/test/scala/com/thoughtworks/compute/TreesSpec.scala @@ -0,0 +1,29 @@ +package com.thoughtworks.compute + +import com.thoughtworks.compute.Trees.FloatArrayTrees +import com.thoughtworks.feature.Factory +import org.scalatest.{FreeSpec, Matchers} + +/** + * @author 杨博 (Yang Bo) + */ +final class TreesSpec extends FreeSpec with Matchers { + + "hashCode" in { + + val trees: FloatArrayTrees = Factory[Trees.FloatArrayTrees with Trees.StructuralTrees].newInstance() + + trees.float.literal(42.0f).## should be(trees.float.literal(42.0f).##) + trees.float.literal(42.0f).## shouldNot be(trees.float.literal(41.0f).##) + trees.float.parameter("my_id").## should be(trees.float.parameter("my_id").##) + trees.float.parameter("my_id_1").## should be(trees.float.parameter("my_id_2").##) + + trees.array.parameter("my_id", trees.float, 42.0f, Array(12, 34)).## should be( + trees.array.parameter("my_id2", trees.float, 42.0f, Array(12, 34)).##) + + trees.array.parameter("my_id", trees.float, 0.1f, Array(12, 34)).## shouldNot be( + trees.array.parameter("my_id2", trees.float, 99.9f, Array(56, 78)).##) + + } + +}