diff --git a/examples/concepts/sycl/dataAccessors/AccessorInNDRange.cpp b/examples/concepts/sycl/dataAccessors/AccessorInNDRange.cpp index ac3f06aba1..95e4506855 100644 --- a/examples/concepts/sycl/dataAccessors/AccessorInNDRange.cpp +++ b/examples/concepts/sycl/dataAccessors/AccessorInNDRange.cpp @@ -1,7 +1,7 @@ #include /*@ - requires \pointer(a, 10, write); + context \pointer(a, 10, write); */ void test(int* a) { sycl::queue myQueue; @@ -28,6 +28,8 @@ void test(int* a) { ); } // Leaving scope, which destroys aBuffer, which waits on the kernel to terminate as it uses aBuffer + //@ assert (\forall int i; 0 <= i && i < 10 ==> a[i] == a[sycl::linearize2(i/5, i%5, 2,5)]); + //@ assert (\forall int i; i >= 0 && i < 10; a[i] == 10); } \ No newline at end of file diff --git a/examples/concepts/sycl/dataAccessors/MissingRangeRequirements.cpp b/examples/concepts/sycl/dataAccessors/MissingRangeRequirements.cpp index a2d194b91c..5a40d88ae2 100644 --- a/examples/concepts/sycl/dataAccessors/MissingRangeRequirements.cpp +++ b/examples/concepts/sycl/dataAccessors/MissingRangeRequirements.cpp @@ -1,5 +1,9 @@ #include +/* + This example is obsolete. Due to fixes in the scoping, the missing annotation here is not needed anymore. + I call it, "missing by design" +*/ /*@ requires \pointer(a, 12, write); */ @@ -17,6 +21,8 @@ void test(int* a) { /*@ requires 1 < a_accessor.get_range().get(0); requires 2 < a_accessor.get_range().get(1); + requires (\forall* int x = 0 .. a_accessor.get_range().get(0), int y = 0 .. a_accessor.get_range().get(1), int z = 0 .. a_accessor.get_range().get(2); + Perm(a_accessor[x][y][z], read)); */ [=] (sycl::item<1> it) { // Bound requirement for the 3rd index is missing, so cannot verify int a = a_accessor[1][2][1]; diff --git a/res/universal/res/cpp/sycl/sycl.hpp b/res/universal/res/cpp/sycl/sycl.hpp index 040ff6f6b5..fc3e4723e2 100644 --- a/res/universal/res/cpp/sycl/sycl.hpp +++ b/res/universal/res/cpp/sycl/sycl.hpp @@ -1,5 +1,14 @@ namespace sycl { + namespace h { + /*@ + ensures (a>0 && b>0) ==> \result > 0; + ensures \result == a*b; + pure int mul(int a, int b) = a*b; + + */ + } + namespace event { void wait(); } @@ -27,7 +36,7 @@ namespace sycl { requires id0 >= 0 && id0 < r0 && id1 >= 0 && id1 < r1; requires r0 > 0 && r1 > 0; ensures \result == sycl::linearize2formula(id0, id1, r1); - ensures \result >= 0 && \result < r0 * r1; + ensures \result >= 0 && \result < sycl::h::mul(r0,r1); ensures (\forall int ida0, int ida1; ida0 >= 0 && ida0 < r0 && ida1 >= 0 && ida1 < r1 && @@ -36,7 +45,7 @@ namespace sycl { ); pure int linearize2(int id0, int id1, int r0, int r1); - pure int linearize2formula(int id0, int id1, int r1) = id1 + (id0 * r1); + pure int linearize2formula(int id0, int id1, int r1) = id1 + sycl::h::mul(id0, r1); requires id0 >= 0 && id0 < r0 && id1 >= 0 && id1 < r1 && id2 >= 0 && id2 < r2; requires r0 > 0 && r1 > 0 && r2 > 0; diff --git a/src/col/vct/col/ast/lang/sycl/SYCLTBufferImpl.scala b/src/col/vct/col/ast/lang/sycl/SYCLTBufferImpl.scala index 0b09c49414..0f954887da 100644 --- a/src/col/vct/col/ast/lang/sycl/SYCLTBufferImpl.scala +++ b/src/col/vct/col/ast/lang/sycl/SYCLTBufferImpl.scala @@ -117,6 +117,7 @@ trait SYCLTBufferImpl[G] extends SYCLTBufferOps[G] { new Variable[G](TPointer(this.typ))(o.where(name = "hostData")) val sizeVar = new Variable[G](TCInt())(o.where(name = "size")) val indexVar = new Variable[G](TCInt())(o.where(name = "i")) + val indexVar2 = new Variable[G](TCInt())(o.where(name = "i")) val copyHostdataToBufferBlame = PanicBlame( "The generated method 'copy_hostData_to_buffer' should not throw any errors." ) @@ -150,6 +151,29 @@ trait SYCLTBufferImpl[G] extends SYCLTBufferOps[G] { WritePerm(), copyHostdataToBufferBlame, ), + Forall( + bindings = Seq(indexVar2), + triggers = Seq( + Seq(ArraySubscript(result, Local[G](indexVar2.ref))( + copyHostdataToBufferBlame + )) + ), + body = + (Local[G](indexVar2.ref) >= c_const(0) && + Local[G](indexVar2.ref) < Local[G](sizeVar.ref)) ==> Eq( + PointerSubscript( + Local[G](hostDataVar.ref), + Local[G](indexVar2.ref), + )(copyHostdataToBufferBlame), + Old[G]( + PointerSubscript( + Local[G](hostDataVar.ref), + Local[G](indexVar2.ref), + )(copyHostdataToBufferBlame), + None + )(copyHostdataToBufferBlame) + ), + ), Forall( bindings = Seq(indexVar), triggers = Seq( diff --git a/src/col/vct/col/resolve/lang/CPP.scala b/src/col/vct/col/resolve/lang/CPP.scala index 515eb67262..aa6059e0c4 100644 --- a/src/col/vct/col/resolve/lang/CPP.scala +++ b/src/col/vct/col/resolve/lang/CPP.scala @@ -188,7 +188,6 @@ case object CPP { getBaseTypeFromSpecs(returnedSpecs) case x => x } - } def paramsFromDeclarator[G](declarator: CPPDeclarator[G]): Seq[CPPParam[G]] = diff --git a/src/col/vct/col/typerules/CoercionUtils.scala b/src/col/vct/col/typerules/CoercionUtils.scala index d03d6d8ebd..4d1bc62a73 100644 --- a/src/col/vct/col/typerules/CoercionUtils.scala +++ b/src/col/vct/col/typerules/CoercionUtils.scala @@ -34,9 +34,7 @@ case object CoercionUtils { case (TResource(), TResourceVal()) => CoerceResourceResourceVal() case (TResourceVal(), TResource()) => CoerceResourceValResource() case (TBool(), TResource()) => CoerceBoolResource() - case (TBool(), TResourceVal()) => - CoercionSequence(Seq(CoerceBoolResource(), CoerceResourceResourceVal())) - + case (TBool(), TResourceVal()) => CoercionSequence(Seq(CoerceBoolResource(), CoerceResourceResourceVal())) case (_, TAnyValue()) => CoerceSomethingAnyValue(source) case (source @ TOption(innerSource), target @ TOption(innerTarget)) => diff --git a/src/parsers/antlr4/SpecLexer.g4 b/src/parsers/antlr4/SpecLexer.g4 index c8a683e390..a92f6080a8 100644 --- a/src/parsers/antlr4/SpecLexer.g4 +++ b/src/parsers/antlr4/SpecLexer.g4 @@ -34,6 +34,9 @@ NEVER: EOF '='; // Must be able to contain identifiers from any frontend, so it's fine to over-approximate valid identifiers a bit. LANG_ID_ESCAPE: '`' ~[`]+ '`'; +VAL_PrependOp : '+:'; +VAL_AppendOp : ':+'; + VAL_RESOURCE: 'resource'; VAL_PROCESS: 'process'; VAL_FRAC: 'frac'; diff --git a/src/parsers/antlr4/SpecParser.g4 b/src/parsers/antlr4/SpecParser.g4 index d7b73b3d24..aa0852846b 100644 --- a/src/parsers/antlr4/SpecParser.g4 +++ b/src/parsers/antlr4/SpecParser.g4 @@ -100,8 +100,8 @@ valImpOp: '-*' | '==>'; valAndOp: '**'; valInOp: '\\in'; valMulOp: '\\'; -valPrependOp : '::'; -valAppendOp : '++'; // postfix issues? maybe disable in spec - no side effects? +valPrependOp : '+:'; +valAppendOp : ':+'; // postfix issues? maybe disable in spec - no side effects? valPostfix : '[' '..' langExpr ']' | '[' langExpr '..' langExpr? ']' diff --git a/src/rewrite/vct/rewrite/lang/LangCPPToCol.scala b/src/rewrite/vct/rewrite/lang/LangCPPToCol.scala index 3a0ccfa6e3..c1c2385bed 100644 --- a/src/rewrite/vct/rewrite/lang/LangCPPToCol.scala +++ b/src/rewrite/vct/rewrite/lang/LangCPPToCol.scala @@ -2,17 +2,15 @@ package vct.rewrite.lang import com.typesafe.scalalogging.LazyLogging import hre.util.ScopedStack -import vct.col.ast.util.ExpressionEqualityCheck.isConstantInt import vct.col.ast.{ CPPLocalDeclaration, Expr, FunctionInvocation, InstanceField, Perm, - ProcedureInvocation, - SeqSubscript, _, } +import vct.col.ast.util.ExpressionEqualityCheck.isConstantInt import vct.col.origin._ import vct.col.ref.Ref import vct.col.resolve.NotApplicable @@ -210,6 +208,14 @@ case object LangCPPToCol { inv.blame.blame(SYCLItemMethodPreconditionFailed(error.node)) } + private case class SYCLMulMethodInvocationBlame() + extends CPPInvocationBlame("SYCL item methods") { + override def preconditionFailed(error: PreconditionFailed): Unit = + PanicBlame( + "The mul helper method doesn't have preconditions. This should not happen." + ).blame(error) + } + private case class SYCLItemMethodSeqBoundFailureBlame(inv: CPPInvocation[_]) extends Blame[SeqBoundFailure] { private case class SYCLItemMethodSeqBoundNegativeError() extends UserError { @@ -332,9 +338,9 @@ case object LangCPPToCol { } } - private case class SYCLKernelRangeInvalidBlame() - extends CPPInvocationBlame("SYCL kernel constructors") { - private case class KernelRangeInvalidError(error: PreconditionFailed) + private case class SYCLKernelRangeInvalidBlame(expr: Expr[_]) + extends Blame[AssertFailed] { + private case class KernelRangeInvalidError(error: AssertFailed) extends UserError { override def code: String = "syclKernelRangeInvalid" @@ -352,13 +358,13 @@ case object LangCPPToCol { ) case _ => Message.messagesInContext( - (error.node.o, "Precondition may not hold, since ..."), + (error.node.o, "Assertion may not hold, since ..."), (error.failure.node.o, "... " + error.failure.descCompletion), ) } } - def preconditionFailed(error: PreconditionFailed): Unit = + override def blame(error: AssertFailed): Unit = throw KernelRangeInvalidError(error) } @@ -706,8 +712,10 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) val syclLocalAccessorVarToDimensions : mutable.Map[Variable[Post], Seq[Expr[Post]]] = mutable.Map.empty - val currentlyRunningKernels - : mutable.Map[Local[Post], Seq[SYCLAccessor[Post]]] = mutable.Map.empty + val currentlyRunningKernels: mutable.Map[Local[ + Post + ], (UnitAccountedPredicate[Post], Seq[SYCLAccessor[Post]])] = + mutable.Map.empty val currentDimensionIterVars : mutable.Map[KernelScopeLevel, mutable.Buffer[IterVariable[Post]]] = mutable.Map.empty @@ -720,58 +728,46 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) private case class GroupScope() extends KernelScopeLevel("GROUP_ID") sealed abstract class KernelType( - rangeFields: Seq[InstanceField[Post]], + rangeFields: Seq[Local[Post]], rangeValues: Seq[Expr[Post]], ) { def getRangeSizeChecksForConstructor( params: mutable.Buffer[Variable[Post]] ): Expr[Post] - def getConstructorPostConditions( - result: Result[Post], - params: mutable.Buffer[Variable[Post]], - ): Seq[Expr[Post]] + def getRangeParamAssignments( + params: mutable.Buffer[Variable[Post]] + ): Seq[Assign[Post]] - def getRangeFields: Seq[InstanceField[Post]] = rangeFields + def getRangeFields: Seq[Local[Post]] = rangeFields def getRangeValues: Seq[Expr[Post]] = rangeValues - def getRangeFieldPermissions(thisObj: Expr[Post]): Seq[Expr[Post]] = - rangeFields.map(f => { - implicit val o: Origin = f.o - Star( - Perm[Post](FieldLocation[Post](thisObj, f.ref), ReadPerm()), - Deref[Post](thisObj, f.ref)(new SYCLRangeDerefBlame(f)) >= c_const(0), - ) - }) } private case class BasicKernel( rangeO: Origin, - rangeFields: Seq[InstanceField[Post]], + rangeFields: Seq[Local[Post]], rangeValues: Seq[Expr[Post]], ) extends KernelType(rangeFields, rangeValues) { override def getRangeSizeChecksForConstructor( params: mutable.Buffer[Variable[Post]] - ): Expr[Post] = + ): Expr[Post] = { foldStar(params.indices.map(i => { implicit val o: Origin = RangeDimensionCheckOrigin(params(i).o) Local[Post](params(i).ref) >= c_const(0) }))(RangeDimensionCheckOrigin(rangeO)) - override def getConstructorPostConditions( - result: Result[Post], - params: mutable.Buffer[Variable[Post]], - ): Seq[Expr[Post]] = + // TODO move over the error messages from (SYCLKernelRangeInvalidBlame) to here by generating the asssertion here. + } + + override def getRangeParamAssignments( + params: mutable.Buffer[Variable[Post]] + ): Seq[Assign[Post]] = params.indices.map(i => { implicit val o: Origin = params(i).o - Star( - Perm(FieldLocation[Post](result, rangeFields(i).ref), ReadPerm()), - Deref[Post](result, rangeFields(i).ref)(new SYCLRangeDerefBlame( - rangeFields(i) - )) === Local[Post](params(i).ref), - ) + assignLocal(Local[Post](rangeFields(i).ref), Local[Post](params(i).ref)) }) } private case class NDRangeKernel( rangeO: Origin, - rangeFields: Seq[InstanceField[Post]], + rangeFields: Seq[Local[Post]], rangeValues: Seq[Expr[Post]], ) extends KernelType(rangeFields, rangeValues) { override def getRangeSizeChecksForConstructor( @@ -789,36 +785,31 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) }))(NDRangeDimensionCheckOrigin(rangeO)) } - override def getConstructorPostConditions( - result: Result[Post], - params: mutable.Buffer[Variable[Post]], - ): Seq[Expr[Post]] = { - // Order of rangeFields is group0, local0, group1, local1, ... - // Order of params is global0, local0, global1, local1, ... - Seq.range(0, rangeFields.size, 2).map(i => { + override def getRangeParamAssignments( + params: mutable.Buffer[Variable[Post]] + ): Seq[Assign[Post]] = { + // Order of rangeFields is group0, local0, group1, local1, ... + // Order of params is global0, local0, global1, local1, ... + Seq.range(0, rangeFields.size, 2).flatMap(i => { implicit val o: Origin = params(i).o - foldStar(Seq( - Perm(FieldLocation[Post](result, rangeFields(i).ref), ReadPerm()), - Deref[Post](result, rangeFields(i).ref)(new SYCLRangeDerefBlame( - rangeFields(i) - )) === FloorDiv( - Local[Post](params(i).ref), + Seq( + assignLocal( + Local[Post](rangeFields(i).ref), + AmbiguousDiv( + Local[Post](params(i).ref), + Local[Post](params(i + 1).ref), + )(ImpossibleDivByZeroBlame()), + ), + assignLocal( + Local[Post](rangeFields(i + 1).ref), Local[Post](params(i + 1).ref), - )(ImpossibleDivByZeroBlame()), - Perm(FieldLocation[Post](result, rangeFields(i + 1).ref), ReadPerm()), - Deref[Post](result, rangeFields(i + 1).ref)(new SYCLRangeDerefBlame( - rangeFields(i + 1) - )) === Local[Post](params(i + 1).ref), - Deref[Post](result, rangeFields(i).ref)(new SYCLRangeDerefBlame( - rangeFields(i) - )) * Deref[Post](result, rangeFields(i + 1).ref)( - new SYCLRangeDerefBlame(rangeFields(i + 1)) - ) === Local[Post](params(i).ref), - )) + ), + // TODO Do I somehow need to assert this or not? + // Local[Post](rangeFields(i).ref) * Local[Post](rangeFields(i + 1).ref) === Local[Post](params(i).ref), + ) }) } } - case class SYCLBuffer[Post]( hostData: Expr[Post], generatedVar: Variable[Post], @@ -828,8 +819,8 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) case class SYCLAccessor[Post]( buffer: SYCLBuffer[Post], accessMode: SYCLAccessMode[Post], - instanceField: InstanceField[Post], - rangeIndexFields: Seq[InstanceField[Post]], + local: Local[Post], + rangeLocals: Seq[Local[Post]], )(implicit val o: Origin) private def getFromAll[K, V]( @@ -1096,9 +1087,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) case _: SYCLTAccessor[Pre] => // Referencing an accessor variable can only be done in kernels, otherwise an error will already have been thrown val accessor = syclAccessorSuccessor(ref) - Deref[Post](currentThis.get, accessor.instanceField.ref)( - SYCLAccessorFieldInsufficientReferencePermissionBlame(local) - ) + accessor.local case _: SYCLTLocalAccessor[Pre] if currentKernelType.get.isInstanceOf[BasicKernel] => throw SYCLNoLocalAccessorsInBasicKernel(local) @@ -1142,14 +1131,15 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) .classInstance ) match { case localVar: Local[Post] => - val accessors: Seq[SYCLAccessor[Post]] = currentlyRunningKernels - .getOrElse( + val (kernelrunnerpost, accessors) + : (UnitAccountedPredicate[Post], Seq[SYCLAccessor[Post]]) = + currentlyRunningKernels.getOrElse( localVar, throw Unreachable(inv.o.messageInContext( "Could not find the event variable in the stored running kernels." )), ) - syclKernelTermination(localVar, accessors) + syclKernelTermination(localVar, kernelrunnerpost, accessors) case _ => throw Unreachable(inv.o.messageInContext( "The object on which the wait() method was called is not a locally declared SYCL event." @@ -1236,20 +1226,13 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) getGlobalWorkItemLinearId(inv) case "sycl::accessor::get_range" => classInstance match { - case Some(Deref(_, ref)) => + + case Some(Local(ref)) => + val a = 1 + 1 val accessor = syclAccessorSuccessor.values - .find(acc => ref.decl.equals(acc.instanceField)).get - LiteralSeq[Post]( - TCInt(), - accessor.rangeIndexFields.map(f => - Deref[Post](currentThis.get, f.ref)( - SYCLAccessorRangeIndexFieldInsufficientReferencePermissionBlame( - inv - ) - ) - ), - ) + .find(acc => ref.equals(acc.local.ref)).get + LiteralSeq[Post](TCInt(), accessor.rangeLocals) case _ => throw NotApplicable(inv) } case "sycl::local_accessor::get_range" => @@ -1262,12 +1245,16 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) case "sycl::range::get" => (classInstance, args) match { case (Some(seq: LiteralSeq[Post]), Seq(arg)) => - SeqSubscript(seq, rw.dispatch(arg))( - SYCLRequestedRangeIndexOutOfBoundsBlame(seq, arg) - ) // Range coming from calling get_range() on a (local)accessor + isConstantInt(arg) match { + case Some(i) if 0 <= i && i < seq.values.size => + seq.values(i.toInt) + case _ => + SeqSubscript(seq, rw.dispatch(arg))( + SYCLRequestedRangeIndexOutOfBoundsBlame(seq, arg) + ) // Range coming from calling get_range() on a (local)accessor + } case _ => throw NotApplicable(inv) } - case _ => val procedureRef: Ref[Post, Procedure[Post]] = cppFunctionDeclSuccessor .ref((decls, initIdx)) @@ -1348,14 +1335,6 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) ) } - // Create a class that can be used to create a 'this' object - // It will be linked to the class made near the end of this method. - val preEventClass: Class[Pre] = - new ByReferenceClass(Nil, Nil, Nil, tt)(commandGroup.o) - this.currentThis = Some( - rw.dispatch(ThisObject[Pre](preEventClass.ref)(preEventClass.o)) - ) - // Generate conditions and accessor objects for each accessor declared in the command group val collectedAccessorDeclarations: Seq[CPPLocalDeclaration[Pre]] = findStatements( @@ -1368,6 +1347,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) case _ => None }, ) + val ( accessors, accessorRunMethodConditions, @@ -1407,8 +1387,13 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) kernelDeclaration.declarator, ) - // Create a block of code for the kernel body based on what type of kernel it is - val (kernelParBlock, contractRequires, contractEnsures) = +// Create a block of code for the kernel body based on what type of kernel it is + val ( + kernelParBlock, + contractRequires, + contractEnsures, + contractContextEverywhere, + ) = rangeType match { case SYCLTRange(_) => createBasicKernelBody( @@ -1430,16 +1415,34 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) ) } + val rangeVars: mutable.Buffer[Variable[Post]] = mutable.Buffer.empty + + // Generate expressions that check the bounds of the given range + rangeVars.appendAll(currentKernelType.get.getRangeFields.map(f => + new Variable[Post](TCInt())(f.o) + )) + + val rangeChecks = currentKernelType.get + .getRangeSizeChecksForConstructor(rangeVars) + + val rangeAssignments = currentKernelType.get + .getRangeParamAssignments(rangeVars) + // Create the pre- and postconditions for the run-method that will hold the generated kernel code + val kernelRunnerContextEverywhere = + getKernelQuantifiedCondition( + kernelParBlock, + removeKernelAccessorPermissions(contractContextEverywhere, accessors), + )(kernelDeclaration.contract.contextEverywhere.o) + val kernelRunnerPreCondition = { implicit val o: Origin = kernelDeclaration.contract.requires.o UnitAccountedPredicate( foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - accessorRunMethodConditions :+ getKernelQuantifiedCondition( - kernelParBlock, - removeKernelClassInstancePermissions(contractRequires), - ) + accessorRunMethodConditions :+ getKernelQuantifiedCondition( + kernelParBlock, + removeKernelAccessorPermissions(contractRequires, accessors), + ) :+ kernelRunnerContextEverywhere )(commandGroupBody.o) ) } @@ -1447,95 +1450,66 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) implicit val o: Origin = kernelDeclaration.contract.ensures.o UnitAccountedPredicate( foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - accessorRunMethodConditions :+ getKernelQuantifiedCondition( - kernelParBlock, - removeKernelClassInstancePermissions(contractEnsures), - ) + accessorRunMethodConditions ++ Seq(getKernelQuantifiedCondition( + kernelParBlock, + removeKernelAccessorPermissions(contractEnsures, accessors), + )) ++ Seq(kernelRunnerContextEverywhere) )(commandGroupBody.o) ) } - // Declare the newly generated kernel code inside a run-method - val kernelRunnerContract = - ApplicableContract[Post]( - kernelRunnerPreCondition, - kernelRunnerPostCondition, - tt, - Nil, - Nil, - Nil, - None, - )(SYCLKernelRunMethodContractUnsatisfiableBlame( - kernelRunnerPreCondition - ))(commandGroup.o) - val kernelRunner = - new RunMethod[Post]( - body = Some( - ParStatement[Post](kernelParBlock)(kernelDeclaration.body.o) - ), - contract = kernelRunnerContract, - )(KernelLambdaRunMethodBlame(kernelDeclaration))(commandGroup.o) - // Create the surrounding class - val postEventClass = - new ByReferenceClass[Post]( - typeArgs = Seq(), - decls = - currentKernelType.get.getRangeFields ++ - accessors - .flatMap(acc => acc.instanceField +: acc.rangeIndexFields) ++ - Seq(kernelRunner), - supports = Seq(), - intrinsicLockInvariant = tt, - )(commandGroup.o.where(name = "SYCL_EVENT_CLASS")) - rw.globalDeclarations.succeed(preEventClass, postEventClass) + val postEventClass: Class[Post] = + new ByReferenceClass(Nil, Nil, Nil, tt)(commandGroup.o) + rw.globalDeclarations.declare(postEventClass) // Create a variable to refer to the class instance val eventClassRef = new Variable[Post](TByReferenceClass(postEventClass.ref, Seq()))( commandGroup.o.where(name = "sycl_event_ref") ) - // Store the class ref and read-write accessors to be used when the kernel is done running +// Store the class ref and read-write accessors to be used when the kernel is done running currentlyRunningKernels.put( eventClassRef.get(commandGroup.o), - accessors - .filter(acc => acc.accessMode.isInstanceOf[SYCLReadWriteAccess[Post]]), - ) - - // Declare a constructor for the class as a separate global method - val eventClassConstructor = createEventClassConstructor( - accessors, - preEventClass, - commandGroup.o, + ( + kernelRunnerPostCondition, + accessors + .filter(acc => acc.accessMode.isInstanceOf[SYCLReadWriteAccess[Post]]), + ), ) // Create a new class instance and assign it to the class instance variable, then fork that variable val result = ( Block[Post]( - bufferAccessStatements ++ Seq( - LocalDecl[Post](eventClassRef)(commandGroup.o), - assignLocal( - eventClassRef.get(commandGroup.o), - ProcedureInvocation[Post]( // Call constructor - ref = eventClassConstructor.ref, - args = - currentKernelType.get.getRangeValues ++ - accessors.flatMap(acc => - acc.buffer.generatedVar.get(acc.buffer.generatedVar.o) +: - acc.buffer.range.dimensions.map(dim => dim) - ), - Nil, - Nil, - Nil, - Nil, - )(SYCLKernelRangeInvalidBlame())(commandGroup.o), - )(commandGroup.o), - Fork(eventClassRef.get(eventClassRef.o))(SYCLKernelForkBlame( - kernelDeclaration - ))(invocation.o), - ) + bufferAccessStatements ++ + Seq(LocalDecl[Post](eventClassRef)(commandGroup.o)) ++ + currentKernelType.get.getRangeFields + .map(rf => LocalDecl[Post](rf.ref.decl)(rf.o)) ++ + Seq.range(0, rangeVars.size).flatMap(i => + Seq( + LocalDecl[Post](rangeVars(i))(rangeVars(i).o), + assignLocal( + rangeVars(i).get(rangeVars(i).o), + currentKernelType.get.getRangeValues(i), + )(rangeVars(i).o), + ) + ) ++ + Seq(Assert(rangeChecks)((SYCLKernelRangeInvalidBlame(rangeChecks)))( + rangeChecks.o + )) ++ rangeAssignments ++ Seq( + IndetBranch(Seq( + Block(Seq[Statement[Post]]( + ParStatement[Post](kernelParBlock)(kernelDeclaration.body.o), + Inhale(ff)(kernelParBlock.o), + ))(kernelParBlock.o), + Block(Seq( + Exhale[Post](kernelRunnerPreCondition.pred)( + kernelRunnerPreCondition.pred.o + )(kernelRunnerPreCondition.pred.o) + ))(kernelDeclaration.body.o), + ))(kernelDeclaration.body.o) + ) )(invocation.o), eventClassRef, ) @@ -1553,7 +1527,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) kernelDimensions: Expr[Pre], kernelDeclaration: CPPLambdaDefinition[Pre], accessorParblockConditions: Seq[Expr[Post]], - ): (ParBlock[Post], Expr[Post], Expr[Post]) = { + ): (ParBlock[Post], Expr[Post], Expr[Post], Expr[Post]) = { // Register the kernel dimensions val range: Seq[Expr[Post]] = rw.dispatch(kernelDimensions) match { @@ -1566,18 +1540,11 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) currentDimensionIterVars.clear() currentDimensionIterVars(GlobalScope()) = mutable.Buffer.empty - val rangeFields: mutable.Buffer[InstanceField[Post]] = mutable.Buffer.empty + val rangeFields: mutable.Buffer[Local[Post]] = mutable.Buffer.empty range.indices.foreach(index => { implicit val o: Origin = range(index).o.where(name = s"range$index") - val instanceField = new InstanceField[Post](TCInt(), Nil) - rangeFields.append(instanceField) - val iterVar = createRangeIterVar( - GlobalScope(), - index, - Deref[Post](currentThis.get, instanceField.ref)(new SYCLRangeDerefBlame( - instanceField - )), - ) + val iterVar = createRangeIterVar(GlobalScope(), index, range(index)) + rangeFields.append(iterVar.variable.get) currentDimensionIterVars(GlobalScope()).append(iterVar) }) currentKernelType = Some( @@ -1589,6 +1556,9 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) .dispatch(kernelDeclaration.contract.requires) val UnitAccountedPredicate(contractEnsures: Expr[Post]) = rw .dispatch(kernelDeclaration.contract.ensures) + val contractContextEverywhere: Expr[Post] = rw + .dispatch(kernelDeclaration.contract.contextEverywhere) + val idLimits = currentDimensionIterVars(GlobalScope()).map(iterVar => { implicit val o: Origin = iterVar.o Local[Post](iterVar.variable.ref) >= c_const(0) && @@ -1601,20 +1571,16 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) ParBlock[Post]( decl = new ParBlockDecl[Post]()(o.where(name = "SYCL_BASIC_KERNEL")), iters = currentDimensionIterVars(GlobalScope()).toSeq, - context_everywhere = rw - .dispatch(kernelDeclaration.contract.contextEverywhere), - requires = foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - idLimits ++ accessorParblockConditions :+ contractRequires - ), - ensures = foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - idLimits ++ accessorParblockConditions :+ contractEnsures + context_everywhere = foldStar( + idLimits.toSeq ++ accessorParblockConditions :+ + contractContextEverywhere ), + requires = contractRequires, + ensures = contractEnsures, content = rw.dispatch(kernelDeclaration.body), )(SYCLKernelParBlockFailureBlame(kernelDeclaration)) - (parBlock, contractRequires, contractEnsures) + (parBlock, contractRequires, contractEnsures, contractContextEverywhere) } private def createNDRangeKernelBody( @@ -1622,7 +1588,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) kernelDeclaration: CPPLambdaDefinition[Pre], accessorParblockConditions: Seq[Expr[Post]], localAccessorDeclarations: Seq[CPPLocalDeclaration[Pre]], - ): (ParBlock[Post], Expr[Post], Expr[Post]) = { + ): (ParBlock[Post], Expr[Post], Expr[Post], Expr[Post]) = { // Register the kernel dimensions val (globalRange, localRange): (Seq[Expr[Post]], Seq[Expr[Post]]) = rw.dispatch(kernelDimensions) match { @@ -1640,34 +1606,29 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) currentDimensionIterVars.clear() currentDimensionIterVars(LocalScope()) = mutable.Buffer.empty currentDimensionIterVars(GroupScope()) = mutable.Buffer.empty - val rangeFields: mutable.Buffer[InstanceField[Post]] = mutable.Buffer.empty + val rangeFields: mutable.Buffer[Local[Post]] = mutable.Buffer.empty + val gRangeVars: mutable.Buffer[Variable[Post]] = mutable.Buffer.empty + val lRangeVars: mutable.Buffer[Variable[Post]] = mutable.Buffer.empty localRange.indices.foreach(index => { { implicit val o: Origin = kernelDimensions.o .where(name = s"group_range$index") - val groupInstanceField = new InstanceField[Post](TCInt(), Nil) - rangeFields.append(groupInstanceField) - val groupIterVar = createRangeIterVar( - GroupScope(), - index, - Deref[Post](currentThis.get, groupInstanceField.ref)( - new SYCLRangeDerefBlame(groupInstanceField) - ), - ) + val grange = globalRange(index) + val grangeVar = new Variable[Post](TCInt()) + val grangeLocal = new Local[Post](grangeVar.ref) + rangeFields.append(grangeLocal) + val groupIterVar = createRangeIterVar(GroupScope(), index, grangeLocal) currentDimensionIterVars(GroupScope()).append(groupIterVar) } { implicit val o: Origin = localRange(index).o .where(name = s"local_range$index") - val localInstanceField = new InstanceField[Post](TCInt(), Nil) - rangeFields.append(localInstanceField) - val localIterVar = createRangeIterVar( - LocalScope(), - index, - Deref[Post](currentThis.get, localInstanceField.ref)( - new SYCLRangeDerefBlame(localInstanceField) - ), - ) + val lrange = globalRange(index) + val lrangeVar = new Variable[Post](TCInt()) + val lrangeLocal = new Local[Post](lrangeVar.ref) + rangeFields.append(lrangeLocal) + + val localIterVar = createRangeIterVar(LocalScope(), index, lrangeLocal) currentDimensionIterVars(LocalScope()).append(localIterVar) } }) @@ -1693,6 +1654,9 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) .dispatch(kernelDeclaration.contract.requires) val UnitAccountedPredicate(contractEnsures: Expr[Post]) = rw .dispatch(kernelDeclaration.contract.ensures) + val contractContextEverywhere: Expr[Post] = rw + .dispatch(kernelDeclaration.contract.contextEverywhere) + val localIdLimits = currentDimensionIterVars(LocalScope()).map(iterVar => { implicit val o: Origin = iterVar.o Local[Post](iterVar.variable.ref) >= c_const(0) && @@ -1714,18 +1678,12 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) o.where(name = "SYCL_ND_RANGE_KERNEL_WORKITEMS") ), iters = currentDimensionIterVars(LocalScope()).toSeq, - context_everywhere = rw - .dispatch(kernelDeclaration.contract.contextEverywhere), - requires = foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - groupIdLimits ++ localIdLimits ++ accessorParblockConditions :+ - contractRequires - ), - ensures = foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - groupIdLimits ++ localIdLimits ++ accessorParblockConditions :+ - contractEnsures + context_everywhere = foldStar( + (groupIdLimits ++ localIdLimits ++ accessorParblockConditions) + .toSeq :+ rw.dispatch(kernelDeclaration.contract.contextEverywhere) ), + requires = contractRequires, + ensures = contractEnsures, content = rw.dispatch(kernelDeclaration.body), )(SYCLKernelParBlockFailureBlame(kernelDeclaration)) ) @@ -1753,24 +1711,23 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) o.where(name = "SYCL_ND_RANGE_KERNEL_WORKGROUPS") ), iters = currentDimensionIterVars(GroupScope()).toSeq, - context_everywhere = tt, - requires = foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - groupIdLimits ++ accessorParblockConditions :+ - quantifiedContractRequires - ), - ensures = foldStar( - currentKernelType.get.getRangeFieldPermissions(currentThis.get) ++ - groupIdLimits ++ accessorParblockConditions :+ - quantifiedContractEnsures + context_everywhere = foldStar( + groupIdLimits.toSeq ++ accessorParblockConditions ), + requires = quantifiedContractRequires, + ensures = quantifiedContractEnsures, content = Scope( Nil, Block(localAccessorDecls.toSeq :+ workItemParBlock), ), )(SYCLKernelParBlockFailureBlame(kernelDeclaration)) - (workGroupParBlock, quantifiedContractRequires, quantifiedContractEnsures) + ( + workGroupParBlock, + quantifiedContractRequires, + quantifiedContractEnsures, + contractContextEverywhere, + ) } private def rewriteSYCLAccessorDeclarations( @@ -1828,8 +1785,8 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) SYCLAccessor[Post]( buffer, accessMode, - foundAcc.instanceField, - foundAcc.rangeIndexFields, + foundAcc.local, + foundAcc.rangeLocals, )(accDecl.o) syclAccessorSuccessor(RefCPPLocalDeclaration(decl, 0)) = newAcc @@ -1841,10 +1798,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) accessMode.isInstanceOf[SYCLReadWriteAccess[Post]] ) { val (basicPermissions, fieldArrayElementsPermission) = - getBasicAccessorPermissions( - newAcc, - this.currentThis.get, - ) + getBasicAccessorPermissions(newAcc) buffersToAccessorResults(buffer) = ( newAcc, @@ -1857,26 +1811,24 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) } case None => // No accessor for buffer exist in the command group, so make fields and permissions - val instanceField = - new InstanceField[Post](buffer.generatedVar.t, Nil)(accO) - val rangeIndexFields = Seq - .range(0, buffer.range.dimensions.size).map(i => - new InstanceField[Post](TCInt(), Nil)( - dimO.where(name = s"${accName}_r$i") - ) - ) + val localBuff = buffer.generatedVar.get(accO) + + val rangeIndexLocal = buffer.range.dimensions.map { d => + d.asInstanceOf[Local[Post]] + } + val newAcc = SYCLAccessor[Post]( buffer, accessMode, - instanceField, - rangeIndexFields, + localBuff, + rangeIndexLocal, )(accDecl.o) syclAccessorSuccessor(RefCPPLocalDeclaration(decl, 0)) = newAcc val (basicPermissions, fieldArrayElementsPermission) = - getBasicAccessorPermissions(newAcc, this.currentThis.get) + getBasicAccessorPermissions(newAcc) buffersToAccessorResults(buffer) = ( newAcc, @@ -1923,8 +1875,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) } private def getBasicAccessorPermissions( - acc: SYCLAccessor[Post], - classObj: Expr[Post], + acc: SYCLAccessor[Post] ): (Expr[Post], Perm[Post]) = { val perm: Expr[Post] = acc.accessMode match { @@ -1932,149 +1883,33 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) case SYCLReadOnlyAccess() => ReadPerm[Post]()(acc.accessMode.o) } - val rangeIndexDerefs: Seq[Expr[Post]] = acc.rangeIndexFields.map(f => - Deref[Post](classObj, f.ref)(new SYCLAccessorDimensionDerefBlame(f))(f.o) - ) - ( - foldStar( - acc.rangeIndexFields.map(f => - Perm( - FieldLocation[Post](classObj, f.ref)(f.o), - ReadPerm[Post]()(f.o), - )(f.o) - ) ++ Seq( - Perm( - FieldLocation[Post](classObj, acc.instanceField.ref)( - acc.instanceField.o - ), - ReadPerm()(acc.instanceField.o), - )(acc.instanceField.o), - ValidArray( - Deref[Post](classObj, acc.instanceField.ref)( - new SYCLAccessorDerefBlame(acc.instanceField) - )(acc.instanceField.o), - rangeIndexDerefs.reduce((e1, e2) => (e1 * e2)(acc.buffer.o)), - )(acc.instanceField.o), - ) - )(acc.instanceField.o), + foldStar(Seq( + ValidArray( + acc.local, + acc.rangeLocals.reduce[Expr[Post]] { (e1, e2) => + syclHelperFunctions("sycl_:_:h_:_:mul")( + Seq(e1, e2), + SYCLMulMethodInvocationBlame(), + acc.local.o, + ) + }, + )(acc.local.o) + ))(acc.local.o), Perm( ArrayLocation( - Deref[Post](classObj, acc.instanceField.ref)( - new SYCLAccessorDerefBlame(acc.instanceField) - )(acc.instanceField.o), + acc.local, Any()(PanicBlame( "The accessor field is not null as that was proven in the previous conditions." - ))(acc.instanceField.o), + ))(acc.local.o), )(PanicBlame("All indices of the accessor array should be accessible"))( - acc.instanceField.o + acc.local.o ), perm, )(acc.accessMode.o), ) } - private def createEventClassConstructor( - accessors: Seq[SYCLAccessor[Post]], - preClass: Class[Pre], - commandGroupO: Origin, - ): Procedure[Post] = { - val t = rw.dispatch(TByReferenceClass[Pre](preClass.ref, Seq())) - rw.globalDeclarations.declare( - withResult((result: Result[Post]) => { - val constructorPostConditions: mutable.Buffer[Expr[Post]] = - mutable.Buffer.empty - val constructorArgs: mutable.Buffer[Variable[Post]] = - mutable.Buffer.empty - - // Generate expressions that check the bounds of the given range - constructorArgs.appendAll(currentKernelType.get.getRangeFields.map(f => - new Variable[Post](TCInt())(f.o) - )) - constructorPostConditions.appendAll( - currentKernelType.get - .getConstructorPostConditions(result, constructorArgs) - ) - val preConditions = - UnitAccountedPredicate[Post]( - currentKernelType.get - .getRangeSizeChecksForConstructor(constructorArgs) - )(commandGroupO) - - accessors.foreach(acc => { - val newConstructorAccessorArg = - new Variable[Post](acc.buffer.generatedVar.t)(acc.instanceField.o) - val newConstructorAccessorDimensionArgs = acc.rangeIndexFields - .map(f => new Variable[Post](TCInt())(f.o)) - - val (basicPermissions, fieldArrayElementsPermission) = - getBasicAccessorPermissions(acc, result) - constructorPostConditions.append(basicPermissions) - constructorPostConditions.append(fieldArrayElementsPermission) - constructorPostConditions.append( - foldStar[Post]( - Eq[Post]( - Deref[Post](result, acc.instanceField.ref)( - new SYCLAccessorDerefBlame(acc.instanceField) - )(acc.instanceField.o), - Local[Post](newConstructorAccessorArg.ref)( - newConstructorAccessorArg.o - ), - )(newConstructorAccessorArg.o) +: - Seq.range(0, acc.rangeIndexFields.size).map(i => - Eq[Post]( - Deref[Post](result, acc.rangeIndexFields(i).ref)( - new SYCLAccessorDimensionDerefBlame( - acc.rangeIndexFields(i) - ) - )(acc.rangeIndexFields(i).o), - Local[Post](newConstructorAccessorDimensionArgs(i).ref)( - newConstructorAccessorDimensionArgs(i).o - ), - )(newConstructorAccessorDimensionArgs(i).o) - ) - )(acc.o) - ) - - constructorArgs.append(newConstructorAccessorArg) - constructorArgs.appendAll(newConstructorAccessorDimensionArgs) - }) - - { - implicit val o: Origin = commandGroupO - new Procedure[Post]( - returnType = t, - args = constructorArgs.toSeq, - outArgs = Seq(), - typeArgs = Seq(), - body = None, - contract = - ApplicableContract[Post]( - preConditions, - SplitAccountedPredicate( - left = UnitAccountedPredicate( - (result !== Null()) && (TypeOf(result) === TypeValue(t)) - ), - right = UnitAccountedPredicate[Post](foldStar[Post]( - constructorPostConditions.toSeq :+ IdleToken[Post](result) - )), - ), - tt, - Seq(), - Seq(), - Seq(), - None, - )(SYCLKernelConstructorNontrivialUnsatisfiableBlame( - commandGroupO - )), - )(SYCLKernelConstructorCallableFailureBlame())( - o.where(name = "event_constructor") - ) - } - })(commandGroupO) - ) - } - // Returns what kind of kernel we are working with and check that the kernel ranges match with the (nd_)item ranges private def getKernelRangeType( complicatedKernelRangeType: Type[Pre], @@ -2174,27 +2009,48 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) inv: CPPInvocation[Pre], level: KernelScopeLevel, )(implicit o: Origin): Expr[Post] = { - SeqSubscript[Post]( - LiteralSeq( - TCInt(), - currentDimensionIterVars(level).map(iterVar => iterVar.variable.get) - .toSeq, - ), - rw.dispatch(inv.args.head), - )(SYCLItemMethodSeqBoundFailureBlame(inv)) + val dim = + inv.args match { + case Seq(dim) => dim + case _ => ??? + } + isConstantInt(dim) match { + case Some(i) if 0 <= i && i < currentDimensionIterVars(level).size => + currentDimensionIterVars(level)(i.toInt).variable.get + case _ => + SeqSubscript[Post]( + LiteralSeq( + TCInt(), + currentDimensionIterVars(level).map(iterVar => iterVar.variable.get) + .toSeq, + ), + rw.dispatch(inv.args.head), + )(SYCLItemMethodSeqBoundFailureBlame(inv)) + } } private def getSimpleWorkItemRange( inv: CPPInvocation[Pre], level: KernelScopeLevel, )(implicit o: Origin): Expr[Post] = { - SeqSubscript[Post]( - LiteralSeq( - TCInt(), - currentDimensionIterVars(level).map(iterVar => iterVar.to).toSeq, - ), - rw.dispatch(inv.args.head), - )(SYCLItemMethodSeqBoundFailureBlame(inv)) + val dim = + inv.args match { + case Seq(dim) => dim + case _ => ??? + } + isConstantInt(dim) match { + case Some(i) if 0 <= i && i < currentDimensionIterVars(level).size => + currentDimensionIterVars(level)(i.toInt).to + case _ => + SeqSubscript[Post]( + LiteralSeq( + TCInt(), + currentDimensionIterVars(level).map(iterVar => iterVar.to).toSeq, + ), + rw.dispatch(inv.args.head), + )(SYCLItemMethodSeqBoundFailureBlame(inv)) + + } } private def getSimpleWorkItemLinearId( @@ -2219,43 +2075,77 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) private def getGlobalWorkItemId( inv: CPPInvocation[Pre] )(implicit o: Origin): Expr[Post] = { - val args: Seq[Expr[Post]] = Seq( - SeqSubscript[Post]( - LiteralSeq( - TCInt(), + val dim = rw.dispatch(inv.args.head) + + val arg0: Expr[Post] = + isConstantInt(dim) match { + case Some(i) + if 0 <= i && i < currentDimensionIterVars(LocalScope()).size => currentDimensionIterVars(LocalScope()) - .map(iterVar => iterVar.variable.get).toSeq, - ), - rw.dispatch(inv.args.head), - )(SYCLItemMethodSeqBoundFailureBlame(inv)), - SeqSubscript[Post]( - LiteralSeq( - TCInt(), + .map(iterVar => iterVar.variable.get)(i.toInt) + case _ => + SeqSubscript[Post]( + LiteralSeq( + TCInt(), + currentDimensionIterVars(LocalScope()) + .map(iterVar => iterVar.variable.get).toSeq, + ), + dim, + )(SYCLItemMethodSeqBoundFailureBlame(inv)) + } + val arg1: Expr[Post] = + isConstantInt(dim) match { + case Some(i) + if 0 <= i && i < currentDimensionIterVars(GroupScope()).size => currentDimensionIterVars(GroupScope()) - .map(iterVar => iterVar.variable.get).toSeq, - ), - rw.dispatch(inv.args.head), - )(SYCLItemMethodSeqBoundFailureBlame(inv)), - SeqSubscript[Post]( - LiteralSeq( - TCInt(), - currentDimensionIterVars(LocalScope()).map(iterVar => iterVar.to) - .toSeq, - ), - rw.dispatch(inv.args.head), - )(SYCLItemMethodSeqBoundFailureBlame(inv)), - SeqSubscript[Post]( - LiteralSeq( - TCInt(), - currentDimensionIterVars(GroupScope()).map(iterVar => iterVar.to) - .toSeq, - ), - rw.dispatch(inv.args.head), - )(SYCLItemMethodSeqBoundFailureBlame(inv)), - ) + .map(iterVar => iterVar.variable.get)(i.toInt) + case _ => + SeqSubscript[Post]( + LiteralSeq( + TCInt(), + currentDimensionIterVars(GroupScope()) + .map(iterVar => iterVar.variable.get).toSeq, + ), + dim, + )(SYCLItemMethodSeqBoundFailureBlame(inv)) + } + val arg2: Expr[Post] = + isConstantInt(dim) match { + case Some(i) + if 0 <= i && i < currentDimensionIterVars(LocalScope()).size => + currentDimensionIterVars(LocalScope()) + .map(iterVar => iterVar.to)(i.toInt) + case _ => + SeqSubscript[Post]( + LiteralSeq( + TCInt(), + currentDimensionIterVars(LocalScope()).map(iterVar => iterVar.to) + .toSeq, + ), + dim, + )(SYCLItemMethodSeqBoundFailureBlame(inv)) + } + + val arg3: Expr[Post] = + isConstantInt(dim) match { + case Some(i) + if 0 <= i && i < currentDimensionIterVars(GroupScope()).size => + currentDimensionIterVars(GroupScope()) + .map(iterVar => iterVar.to)(i.toInt) + case _ => + SeqSubscript[Post]( + LiteralSeq( + TCInt(), + currentDimensionIterVars(GroupScope()).map(iterVar => iterVar.to) + .toSeq, + ), + dim, + )(SYCLItemMethodSeqBoundFailureBlame(inv)) + + } syclHelperFunctions("sycl_:_:linearize_2")( - args, + Seq(arg0, arg1, arg2, arg3), SYCLItemMethodInvocationBlame(inv), o, ) @@ -2372,9 +2262,25 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) // Get arguments val hostData = rw.dispatch(inv.args.head) - val range = + val (range, dimdecls) = rw.dispatch(inv.args(1)) match { - case r: SYCLRange[Post] => r + case r: SYCLRange[Post] => + r + val vars = Seq.range(0, r.dimensions.size).map(i => + new Variable[Post](TCInt())( + r.dimensions(i).o.where(name = s"${hostData}_r$i") + ) + ) + val updatedR = + SYCLRange[Post]( + Seq.range(0, r.dimensions.size).map(i => vars(i).get) + )(r.o) + + val dimdecls = Seq.range(0, vars.size).flatMap(i => + Seq(LocalDecl(vars(i)), assignLocal(vars(i).get, r.dimensions(i))) + ) + + (updatedR, dimdecls) case _ => throw Unreachable( "The range parameter of the buffer was not rewritten to a range." @@ -2405,7 +2311,8 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) .put(v, SYCLBuffer[Post](hostData, v, range, postType)) val result: Statement[Post] = Block( - Seq(LocalDecl(v), assignLocal(v.get, copyInv), gainExclusiveAccess) + dimdecls ++ + Seq(LocalDecl(v), assignLocal(v.get, copyInv), gainExclusiveAccess) ) (result, v) } @@ -2418,13 +2325,14 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) // Wait for SYCL kernels that access the buffer to finish executing val kernelsToTerminate = currentlyRunningKernels.filter(tuple => - tuple._2.exists(acc => + tuple._2._2.exists(acc => acc.buffer.equals(buffer) && acc.accessMode.isInstanceOf[SYCLReadWriteAccess[Post]] ) ) val kernelTerminations = - kernelsToTerminate.map(tuple => syclKernelTermination(tuple._1, tuple._2)) + kernelsToTerminate + .map(tuple => syclKernelTermination(tuple._1, tuple._2._1, tuple._2._2)) .toSeq // Get the exclusive access predicate and copy procedures @@ -2460,22 +2368,26 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) ): Statement[Post] = { implicit val o: Origin = source.o - val kernelsToWaitFor: mutable.Map[Local[Post], Seq[SYCLAccessor[Post]]] = + val kernelsToWaitFor: mutable.Map[Local[ + Post + ], (UnitAccountedPredicate[Post], Seq[SYCLAccessor[Post]])] = mode match { case SYCLReadOnlyAccess() => currentlyRunningKernels.filter(tuple => - tuple._2.exists(acc => + tuple._2._2.exists(acc => acc.accessMode.isInstanceOf[SYCLReadWriteAccess[Post]] && acc.buffer.equals(buffer) ) ) case SYCLReadWriteAccess() => - currentlyRunningKernels - .filter(tuple => tuple._2.exists(acc => acc.buffer.equals(buffer))) + currentlyRunningKernels.filter(tuple => + tuple._2._2.exists(acc => acc.buffer.equals(buffer)) + ) } Block[Post]( - kernelsToWaitFor.map(tuple => syclKernelTermination(tuple._1, tuple._2)) + kernelsToWaitFor + .map(tuple => syclKernelTermination(tuple._1, tuple._2._1, tuple._2._2)) .toSeq ) } @@ -2486,15 +2398,10 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) if CPP.unwrappedType(base.t).isInstanceOf[SYCLTAccessor[Pre]] => CPP.unwrappedType(base.t) match { case SYCLTAccessor(_, 1, _) => - ArraySubscript[Post]( - Deref[Post]( - currentThis.get, - syclAccessorSuccessor(base.ref.get).instanceField.ref, - )(new SYCLAccessorDerefBlame( - syclAccessorSuccessor(base.ref.get).instanceField - ))(sub.o), - rw.dispatch(index), - )(SYCLAccessorArraySubscriptErrorBlame(sub))(sub.o) + val accessor = syclAccessorSuccessor(base.ref.get) + ArraySubscript[Post](accessor.local, rw.dispatch(index))( + SYCLAccessorArraySubscriptErrorBlame(sub) + )(sub.o) case t: SYCLTAccessor[Pre] => throw SYCLWrongNumberOfSubscriptsForAccessor(sub, 1, t.dimCount) case _ => ??? @@ -2508,19 +2415,13 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) val linearizeArgs = Seq( rw.dispatch(indexX), rw.dispatch(indexY), - Deref[Post](currentThis.get, accessor.rangeIndexFields(0).ref)( - new SYCLAccessorDimensionDerefBlame(accessor.rangeIndexFields(0)) - ), - Deref[Post](currentThis.get, accessor.rangeIndexFields(1).ref)( - new SYCLAccessorDimensionDerefBlame(accessor.rangeIndexFields(1)) - ), + accessor.rangeLocals(0), + accessor.rangeLocals(1), ) CPP.unwrappedType(base.t) match { case SYCLTAccessor(_, 2, _) => ArraySubscript[Post]( - Deref[Post](currentThis.get, accessor.instanceField.ref)( - new SYCLAccessorDerefBlame(accessor.instanceField) - ), + accessor.local, syclHelperFunctions("sycl_:_:linearize_2")( linearizeArgs, SYCLAccessorArraySubscriptLinearizeInvocationBlame( @@ -2548,22 +2449,14 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) rw.dispatch(indexX), rw.dispatch(indexY), rw.dispatch(indexZ), - Deref[Post](currentThis.get, accessor.rangeIndexFields(0).ref)( - new SYCLAccessorDimensionDerefBlame(accessor.rangeIndexFields(0)) - ), - Deref[Post](currentThis.get, accessor.rangeIndexFields(1).ref)( - new SYCLAccessorDimensionDerefBlame(accessor.rangeIndexFields(1)) - ), - Deref[Post](currentThis.get, accessor.rangeIndexFields(2).ref)( - new SYCLAccessorDimensionDerefBlame(accessor.rangeIndexFields(2)) - ), + accessor.rangeLocals(0), + accessor.rangeLocals(1), + accessor.rangeLocals(2), ) CPP.unwrappedType(base.t) match { case SYCLTAccessor(_, 3, _) => ArraySubscript[Post]( - Deref[Post](currentThis.get, accessor.instanceField.ref)( - new SYCLAccessorDerefBlame(accessor.instanceField) - ), + accessor.local, syclHelperFunctions("sycl_:_:linearize_3")( linearizeArgs, SYCLAccessorArraySubscriptLinearizeInvocationBlame( @@ -2605,41 +2498,98 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) } } - private def removeKernelClassInstancePermissions( - e: Expr[Post] + private def removeKernelAccessorPermissions( + e: Expr[Post], + accs: Seq[SYCLAccessor[Post]], ): Expr[Post] = { implicit val o = e.o e match { case s @ Starall(bindings, triggers, body) => - Starall(bindings, triggers, removeKernelClassInstancePermissions(body))( - s.blame - ) + Starall( + bindings, + triggers, + removeKernelAccessorPermissions(body, accs), + )(s.blame) case ModelAbstractState(model, state) => ModelAbstractState( - removeKernelClassInstancePermissions(model), - removeKernelClassInstancePermissions(state), + removeKernelAccessorPermissions(model, accs), + removeKernelAccessorPermissions(state, accs), ) case ModelState(model, perm, state) => ModelState( - removeKernelClassInstancePermissions(model), - removeKernelClassInstancePermissions(perm), - removeKernelClassInstancePermissions(state), + removeKernelAccessorPermissions(model, accs), + removeKernelAccessorPermissions(perm, accs), + removeKernelAccessorPermissions(state, accs), ) case s @ Scale(expr, res) => Scale( - removeKernelClassInstancePermissions(expr), - removeKernelClassInstancePermissions(res), + removeKernelAccessorPermissions(expr, accs), + removeKernelAccessorPermissions(res, accs), )(s.blame) case Star(left, right) => Star( - removeKernelClassInstancePermissions(left), - removeKernelClassInstancePermissions(right), + removeKernelAccessorPermissions(left, accs), + removeKernelAccessorPermissions(right, accs), ) case Wand(left, right) => Wand( - removeKernelClassInstancePermissions(left), - removeKernelClassInstancePermissions(right), + removeKernelAccessorPermissions(left, accs), + removeKernelAccessorPermissions(right, accs), + ) + + case ActionPerm(Local(obj), _) + if accs.exists(acc => acc.local.ref.decl.equals(obj.decl)) => + tt + case ModelPerm(Local(obj), _) + if accs.exists(acc => acc.local.ref.decl.equals(obj.decl)) => + tt + case Perm(FieldLocation(obj, _), _) if obj.equals(this.currentThis.get) => + tt + case PointsTo(FieldLocation(obj, _), _, _) + if obj.equals(this.currentThis.get) => + tt + case Value(FieldLocation(obj, _)) if obj.equals(this.currentThis.get) => + tt + case Perm(AmbiguousLocation(ArraySubscript(Local(obj), _)), _) + if accs.exists(acc => acc.local.ref.decl == obj.decl) => + tt + case PointsTo(AmbiguousLocation(ArraySubscript(Local(obj), _)), _, _) + if accs.exists(acc => acc.local.ref.decl.equals(obj.decl)) => + tt + case Value(AmbiguousLocation(ArraySubscript(Local(obj), _))) + if accs.exists(acc => acc.local.ref.decl.equals(obj.decl)) => + tt + case Implies(left, right) => + Implies( + removeKernelAccessorPermissions(left, accs), + removeKernelAccessorPermissions(right, accs), + ) + case e => e + } + } + + private def removeRangeAccesses(e: Expr[Post]): Expr[Post] = { + implicit val o = e.o + e match { + case s @ Starall(bindings, triggers, body) => + Starall(bindings, triggers, removeRangeAccesses(body))(s.blame) + case ModelAbstractState(model, state) => + ModelAbstractState( + removeRangeAccesses(model), + removeRangeAccesses(state), ) + case ModelState(model, perm, state) => + ModelState( + removeRangeAccesses(model), + removeRangeAccesses(perm), + removeRangeAccesses(state), + ) + case s @ Scale(expr, res) => + Scale(removeRangeAccesses(expr), removeRangeAccesses(res))(s.blame) + case Star(left, right) => + Star(removeRangeAccesses(left), removeRangeAccesses(right)) + case Wand(left, right) => + Wand(removeRangeAccesses(left), removeRangeAccesses(right)) case ActionPerm(Deref(obj, _), _) if obj.equals(this.currentThis.get) => tt @@ -2661,11 +2611,7 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) if obj.equals(this.currentThis.get) => tt case Implies(left, right) => - Implies( - removeKernelClassInstancePermissions(left), - removeKernelClassInstancePermissions(right), - ) - + Implies(removeRangeAccesses(left), removeRangeAccesses(right)) case e => e } } @@ -2682,18 +2628,14 @@ case class LangCPPToCol[Pre <: Generation](rw: LangSpecificToCol[Pre]) private def syclKernelTermination( variable: Local[Post], + kernelRunnerPostCondition: UnitAccountedPredicate[Post], accessors: Seq[SYCLAccessor[Post]], )(implicit o: Origin): Statement[Post] = { currentlyRunningKernels.remove(variable) - Block(Join(variable)(SYCLKernelJoinBlame()) +: accessors.collect({ - case SYCLAccessor(buffer, SYCLReadWriteAccess(), instanceField, _) => - assignLocal( - buffer.generatedVar.get, - Deref[Post](variable, instanceField.ref)(new SYCLAccessorDerefBlame( - instanceField - )), - ) - })) + Block( + Seq(Inhale(kernelRunnerPostCondition.pred)(kernelRunnerPostCondition.o)) + // TODO I need to do something related to readonly accessors. I am currently exhaling all permission, maybe that should only happen if it read/write. + ) } private def findStatements[S]( diff --git a/test/main/vct/test/integration/examples/SYCLSpec.scala b/test/main/vct/test/integration/examples/SYCLSpec.scala index 2d7d7b3a7a..a30aca9909 100644 --- a/test/main/vct/test/integration/examples/SYCLSpec.scala +++ b/test/main/vct/test/integration/examples/SYCLSpec.scala @@ -54,7 +54,7 @@ class SYCLSpec extends VercorsSpec { vercors should verify using silicon flag "--no-infer-heap-context-into-frame" example "concepts/sycl/dataAccessors/GetKernelResult.cpp" vercors should error withCode "syclAccessorInsufficientReferencePermission" example "concepts/sycl/dataAccessors/GetRangeDimensionOutOfBounds1.cpp" vercors should error withCode "syclAccessorInsufficientReferencePermission" example "concepts/sycl/dataAccessors/GetRangeDimensionOutOfBounds2.cpp" - vercors should error withCode "syclAccessorArraySubscriptLinearizePreconditionFailed" example "concepts/sycl/dataAccessors/MissingRangeRequirements.cpp" +// vercors should error withCode "syclAccessorArraySubscriptLinearizePreconditionFailed" example "concepts/sycl/dataAccessors/MissingRangeRequirements.cpp" vercors should error withCode "syclBufferOutOfScope" example "concepts/sycl/dataAccessors/PassBufferToMethod.cpp" vercors should error withCode "syclAccessorArraySubscriptArrayBounds" example "concepts/sycl/dataAccessors/SubscriptOutOfBounds1.cpp" vercors should error withCode "syclAccessorArraySubscriptArrayBounds" example "concepts/sycl/dataAccessors/SubscriptOutOfBounds2.cpp"