@ExplodeLoop
@Override
public Object executeGeneric(VirtualFrame frame) {
    Object function = functionNode.executeGeneric(frame);

    /*
     * The number of arguments is constant for one invoke node. During compilation, the loop is
     * unrolled and the execute methods of all arguments are inlined. This is triggered by the
     * ExplodeLoop annotation on the method. The compiler assertion below illustrates that the
     * array length is really constant.
     */
    CompilerAsserts.compilationConstant(argumentNodes.length);

    Object[] argumentValues = new Object[argumentNodes.length];
    for (int i = 0; i < argumentNodes.length; i++) {
        argumentValues[i] = argumentNodes[i].executeGeneric(frame);
    }

    try {
        return library.execute(function, argumentValues);
    } catch (ArityException | UnsupportedTypeException | UnsupportedMessageException e) {
        /* Execute was not successful. */
        throw HashemUndefinedNameException.undefinedBebin(this, function);
    }
}
 

@Specialization
Object doDefault(VirtualFrame frame,
                @CachedLibrary(limit = "2") InteropLibrary interop,
                @CachedContext(GrCUDALanguage.class) GrCUDAContext context) {
    String[] functionName = identifier.getIdentifierName();
    Namespace namespace = context.getRootNamespace();
    Optional<Object> maybeFunction = namespace.lookup(functionName);
    if (!maybeFunction.isPresent() || !(maybeFunction.get() instanceof Function)) {
        CompilerDirectives.transferToInterpreter();
        throw new GrCUDAException("function '" + GrCUDAException.format(functionName) + "' not found", this);
    }
    Function function = (Function) maybeFunction.get();
    Object[] argumentValues = new Object[argumentNodes.length];
    for (int i = 0; i < argumentNodes.length; i++) {
        argumentValues[i] = argumentNodes[i].execute(frame);
    }
    try {
        return interop.execute(function, argumentValues);
    } catch (ArityException | UnsupportedTypeException | UnsupportedMessageException e) {
        CompilerDirectives.transferToInterpreter();
        throw new GrCUDAException(e.getMessage(), this);
    }
}
 

@ExportMessage
Object execute(Object[] arguments,
                @CachedLibrary(limit = "3") InteropLibrary gridSizeAccess,
                @CachedLibrary(limit = "3") InteropLibrary gridSizeElementAccess,
                @CachedLibrary(limit = "3") InteropLibrary blockSizeAccess,
                @CachedLibrary(limit = "3") InteropLibrary blockSizeElementAccess,
                @CachedLibrary(limit = "3") InteropLibrary sharedMemoryAccess) throws UnsupportedTypeException, ArityException {
    int dynamicSharedMemoryBytes;
    if (arguments.length == 2) {
        dynamicSharedMemoryBytes = 0;
    } else if (arguments.length == 3) {
        // dynamic shared memory specified
        dynamicSharedMemoryBytes = extractNumber(arguments[2], "dynamicSharedMemory", sharedMemoryAccess);
    } else {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(2, arguments.length);
    }

    Dim3 gridSize = extractDim3(arguments[0], "gridSize", gridSizeAccess, gridSizeElementAccess);
    Dim3 blockSize = extractDim3(arguments[1], "blockSize", blockSizeAccess, blockSizeElementAccess);
    KernelConfig config = new KernelConfig(gridSize, blockSize, dynamicSharedMemoryBytes);

    return new ConfiguredKernel(this, config);
}
 

@ExportMessage
@TruffleBoundary
Object execute(Object[] arguments,
                @CachedLibrary(limit = "3") InteropLibrary boolAccess,
                @CachedLibrary(limit = "3") InteropLibrary int8Access,
                @CachedLibrary(limit = "3") InteropLibrary int16Access,
                @CachedLibrary(limit = "3") InteropLibrary int32Access,
                @CachedLibrary(limit = "3") InteropLibrary int64Access,
                @CachedLibrary(limit = "3") InteropLibrary doubleAccess) throws UnsupportedTypeException, ArityException {
    kernel.incrementLaunchCount();
    try (KernelArguments args = kernel.createKernelArguments(arguments, boolAccess, int8Access, int16Access,
                    int32Access, int64Access, doubleAccess)) {
        kernel.getCudaRuntime().cuLaunchKernel(kernel, config, args);
    }
    return this;
}
 

@Specialization(guards = "VALUE.equals(member)")
static MapFunctionBase readMemberValue(MapFunction receiver, String member) {
    return new MapFunctionBase(arguments -> {
        if (arguments.length < 1) {
            throw ArityException.create(1, arguments.length);
        }
        String name = checkString(arguments[0], "name of created value expected");
        if (arguments.length == 1) {
            return receiver.value(name);
        } else {
            Object function = arguments[1];
            Object[] args = Arrays.copyOfRange(arguments, 2, arguments.length);
            return receiver.value(name, function, args);
        }
    });
}
 

@ExportMessage
@TruffleBoundary
Object execute(Object[] arguments) throws UnsupportedMessageException, UnsupportedTypeException, ArityException {
    if (value instanceof MapArgObjectMember) {
        MapArgObjectMember member = (MapArgObjectMember) value;
        if (MAP.equals(member.name)) {
            checkArity(arguments, 1);
            return new MapArgObject(member.parent).map(arguments[0]);
        } else if (SHRED.equals(member.name)) {
            checkArity(arguments, 0);
            CompilerAsserts.neverPartOfCompilation();
            return new MapArgObject(member.parent).shred();
        } else if (DESCRIBE.equals(member.name)) {
            checkArity(arguments, 0);
            CompilerAsserts.neverPartOfCompilation();
            return new MapArgObject(member.parent).describe();
        } else if (BIND.equals(member.name)) {
            checkArity(arguments, 3);
            return member.parent.bind(arguments[0], arguments[1], arguments[2]);
        }
    }
    CompilerDirectives.transferToInterpreter();
    throw UnsupportedMessageException.create();
}
 

@ExportMessage
Object execute(Object[] arguments,
                @CachedLibrary("this.parent") InteropLibrary parentInterop,
                @CachedLibrary(limit = "2") InteropLibrary memberInterop) throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    if (arguments.length != 3) {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(3, arguments.length);
    }
    Object value = parentInterop.execute(parent, arguments);
    try {
        return memberInterop.readMember(value, name);
    } catch (UnsupportedMessageException | UnknownIdentifierException e) {
        CompilerDirectives.transferToInterpreter();
        throw new MapException("cannot read member '" + name + "' from argument " + parent);
    }
}
 

@ExportMessage
Object execute(Object[] arguments,
                @CachedLibrary("this.parent") InteropLibrary parentInterop,
                @CachedLibrary(limit = "2") InteropLibrary elementInterop) throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    if (arguments.length != 3) {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(3, arguments.length);
    }
    Object value = parentInterop.execute(parent, arguments);
    try {
        return elementInterop.readArrayElement(value, index);
    } catch (UnsupportedMessageException | InvalidArrayIndexException e) {
        CompilerDirectives.transferToInterpreter();
        throw new MapException("cannot read element '" + index + "' from argument " + parent);
    }
}
 

@ExportMessage
Object execute(Object[] arguments,
                @CachedLibrary("this.parent") InteropLibrary parentInterop,
                @CachedLibrary("this.function") InteropLibrary mapInterop) throws UnsupportedTypeException, ArityException, UnsupportedMessageException {
    if (arguments.length != 3) {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(3, arguments.length);
    }
    Object value = parentInterop.execute(parent, arguments);
    try {
        return mapInterop.execute(function, value);
    } catch (UnsupportedMessageException | UnsupportedTypeException | ArityException e) {
        CompilerDirectives.transferToInterpreter();
        throw new MapException("cannot map argument " + parent);
    }
}
 

@Override
@TruffleBoundary
public Object call(Object[] arguments) throws ArityException, UnsupportedTypeException {
    if (arguments.length < 1) {
        throw ArityException.create(1, arguments.length);
    }
    String typeName = expectString(arguments[0], "first argument of DeviceArray must be string (type name)");
    Type elementType;
    try {
        elementType = Type.fromGrCUDATypeString(typeName);
    } catch (TypeException e) {
        throw new GrCUDAException(e.getMessage());
    }
    if (arguments.length == 1) {
        return new TypedDeviceArrayFunction(runtime, elementType);
    } else {
        return createArray(arguments, 1, elementType, runtime);
    }
}
 

@ExportMessage
Object execute(Object[] arguments,
                @CachedLibrary(limit = "3") InteropLibrary pointerAccess,
                @CachedLibrary(limit = "3") InteropLibrary numElementsAccess) throws UnsupportedTypeException, ArityException, IndexOutOfBoundsException {
    long numElements;
    if (arguments.length == 1) {
        numElements = deviceArray.getArraySize();
    } else if (arguments.length == 2) {
        numElements = extractNumber(arguments[1], "numElements", numElementsAccess);
    } else {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(1, arguments.length);
    }
    long pointer = extractPointer(arguments[0], "fromPointer", pointerAccess);
    if (direction == CopyDirection.FROM_POINTER) {
        deviceArray.copyFrom(pointer, numElements);
    }
    if (direction == CopyDirection.TO_POINTER) {
        deviceArray.copyTo(pointer, numElements);
    }
    return deviceArray;
}
 

@Specialization(rewriteOn = RespecializeException.class)
protected static final Object invokeMember(final AbstractSqueakObject receiver, final String member, final Object[] arguments,
                @Shared("lookupNode") @Cached final LookupMethodByStringNode lookupNode,
                @Shared("classNode") @Cached final SqueakObjectClassNode classNode,
                @Exclusive @Cached final WrapToSqueakNode wrapNode,
                @Exclusive @Cached final DispatchUneagerlyNode dispatchNode) throws ArityException {
    final int actualArity = arguments.length;
    final Object methodObject = lookupNode.executeLookup(classNode.executeLookup(receiver), toSelector(member, actualArity));
    if (methodObject == null) {
        CompilerDirectives.transferToInterpreterAndInvalidate();
        /* DoesNotUnderstand, rewrite this specialization. */
        throw new RespecializeException();
    }
    final CompiledMethodObject method = (CompiledMethodObject) methodObject;
    final int expectedArity = method.getNumArgs();
    if (actualArity == expectedArity) {
        return dispatchNode.executeDispatch(method, ArrayUtils.copyWithFirst(wrapNode.executeObjects(arguments), receiver), NilObject.SINGLETON);
    } else {
        throw ArityException.create(1 + expectedArity, 1 + actualArity);
    }
}
 

@Specialization(guards = "!values.isNull(obj)", limit = "3")
public Object newObject(Object obj, @CachedLibrary("obj") InteropLibrary values) {
    try {
        return values.instantiate(obj);
    } catch (UnsupportedTypeException | ArityException | UnsupportedMessageException e) {
        /* Foreign access was not successful. */
        throw HashemUndefinedNameException.undefinedBebin(this, obj);
    }
}
 

@Test
public void testArgCount() throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    Object mapped = map.map(new TestFunction());
    TestContainer.check(INTEROP.execute(mapped));
    mapped = map.map(new TestFunction(), map.arg("a"));
    TestContainer.check(INTEROP.execute(mapped, 1.0), v -> (double) v == 1);
    mapped = map.map(new TestFunction(), map.arg("a"), map.arg("b"));
    TestContainer.check(INTEROP.execute(mapped, 1, 2), v -> (int) v == 1, v -> (int) v == 2);
}
 

@Test
public void testArgReuse() throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    Object mapped = map.map(new TestFunction());
    TestContainer.check(INTEROP.execute(mapped));
    mapped = map.map(new TestFunction(), map.arg("a"), map.arg("a"));
    TestContainer.check(INTEROP.execute(mapped, 1.0), v -> (double) v == 1, v -> (double) v == 1);
    mapped = map.map(new TestFunction(), map.arg("a"), map.arg("b"));
    TestContainer.check(INTEROP.execute(mapped, 1, 2), v -> (int) v == 1, v -> (int) v == 2);
}
 

@Test
public void testSimple() throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    Object intArray = asTruffleObject(new int[]{0, 1, 2, 3, 4, 5, 6, 7, 8, 9});
    Object doubleArray = asTruffleObject(new double[]{0, 1, 2, 3, 4, 5, 6, 7, 8, 9});
    Object object = asTruffleObject(new TestClass(5, 10, new int[]{0, 1, 2, 3, 4}, new double[]{0, 1, 2, 3, 4}));

    MappedFunction mapped = map.map(new TestFunction(), map.arg("a").readArrayElement(2));
    TestContainer.check(INTEROP.execute(mapped, intArray), v -> (int) v == 2);
    mapped = map.map(new TestFunction(), map.arg("a").readArrayElement(2), map.arg("a").readArrayElement(4));
    TestContainer.check(INTEROP.execute(mapped, doubleArray), v -> (double) v == 2, v -> (double) v == 4);

    mapped = map.map(new TestFunction(), map.arg("a").readMember("a"), map.arg("a").readMember("c").readArrayElement(3));
    TestContainer.check(INTEROP.execute(mapped, object), v -> (int) v == 5, v -> (int) v == 3);
}
 

@Test
public void testShred() throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    Object object = asTruffleObject(new TestClass(5, 10, new int[]{0, 1, 2, 3, 4}, new double[]{0, 1, 2, 3, 4}));
    Object object2 = asTruffleObject(new Object[]{new TestClass2(5, 10), new TestClass2(6, 11), new TestClass2(7, 12), new TestClass2(8, 13)});

    MapArgObject argC = map.arg("a").shred().readMember("c");
    MapArgObject argD = map.arg("a").shred().readMember("d");
    MappedFunction mapped = map.map(new TestFunction(), argC, argD, map.size(argC, argD));
    TestContainer.check(INTEROP.execute(mapped, object), v -> checkArray(v, new int[]{0, 1, 2, 3, 4}), v -> checkArray(v, new double[]{0, 1, 2, 3, 4}), v -> (long) v == 5);

    TestContainer.check(INTEROP.execute(mapped, object2), v -> checkArray(v, new int[]{5, 6, 7, 8}), v -> checkArray(v, new double[]{10, 11, 12, 13}), v -> (long) v == 4);
}
 

@Test
public void testReturn() throws ArityException, UnsupportedTypeException, UnsupportedMessageException {
    Object object = asTruffleObject(new TestClass(5, 10, new int[]{0, 1, 2, 3, 4}, new double[]{0, 1, 2, 3, 4}));
    Object object2 = asTruffleObject(new Object[]{new TestClass2(5, 10), new TestClass2(6, 11), new TestClass2(7, 12), new TestClass2(8, 13)});

    MapArgObject argC = map.arg("a").shred().readMember("c");
    MapArgObject argD = map.arg("a").shred().readMember("d");
    MappedFunction mapped = map.ret(map.value("output")).map(new TestFunction(), argC, argD, map.value("output", new TestNewArrayFunction(), map.size(argC, argD)));
    checkArray(INTEROP.execute(mapped, object), new int[]{0, 0, 0, 0, 0});
    checkArray(INTEROP.execute(mapped, object2), new int[]{0, 0, 0, 0});
}
 

@ExportMessage
Object invokeMember(String memberName,
                Object[] arguments,
                @CachedLibrary("this") InteropLibrary interopRead,
                @CachedLibrary(limit = "1") InteropLibrary interopExecute)
                throws UnsupportedTypeException, ArityException, UnsupportedMessageException, UnknownIdentifierException {
    return interopExecute.execute(interopRead.readMember(this, memberName), arguments);
}
 

@ExportMessage
Object execute(Object[] arguments) throws ArityException {
    if (arguments.length != 0) {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(0, arguments.length);
    }
    freeMemory();
    return NoneValue.get();
}
 

/**
 * The same as {@link #toSelector(String)}, but may return alternative selector depending on
 * actualArity.
 *
 * @param identifier for interop
 * @param actualArity of selector
 * @return String for Smalltalk selector
 */
@TruffleBoundary
private static String toSelector(final String identifier, final int actualArity) throws ArityException {
    final String selector = identifier.replace('_', ':');
    final int expectedArity = (int) selector.chars().filter(ch -> ch == ':').count();
    if (expectedArity == actualArity) {
        return selector;
    } else {
        if (expectedArity + 1 == actualArity && !selector.endsWith(":")) {
            return selector + ":";
        } else {
            throw ArityException.create(1 + expectedArity, 1 + actualArity);
        }
    }
}
 

@ExportMessage
Object execute(Object[] arguments) throws ArityException {
    if (arguments.length != 0) {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(0, arguments.length);
    }
    freeMemory();
    return NoneValue.get();
}
 

@Specialization(guards = "supportsNFI")
protected final NativeObject doErrorMessageAt(@SuppressWarnings("unused") final Object receiver, final long index,
                @CachedLibrary("getSysCallObject()") final InteropLibrary lib,
                @CachedLibrary(limit = "1") final InteropLibrary resultLib,
                @CachedContext(SqueakLanguage.class) final SqueakImageContext image) {
    try {
        return image.asByteString(resultLib.asString(lib.execute(sysCallObject, (int) index)));
    } catch (final UnsupportedMessageException | UnsupportedTypeException | ArityException e) {
        throw PrimitiveFailed.andTransferToInterpreterWithError(e);
    }
}
 

@SuppressWarnings("static-method")
@ExportMessage
public Object execute(Object[] arguments) throws ArityException {
    if (arguments.length != 0) {
        CompilerDirectives.transferToInterpreter();
        throw ArityException.create(0, arguments.length);
    }
    return runtime.cudaGetDevice() == deviceId;
}
 

private static void initializeObject(final Object[] arguments, final InteropLibrary initializer, final AbstractSqueakObjectWithHash newObject) throws UnsupportedTypeException {
    try {
        initializer.invokeMember(newObject, "initialize");
    } catch (UnsupportedMessageException | ArityException | UnknownIdentifierException | UnsupportedTypeException e) {
        throw UnsupportedTypeException.create(arguments, "Failed to initialize new object");
    }
}
 

@ExportMessage
Object execute(Object[] arguments,
                @CachedLibrary(limit = "3") InteropLibrary gridSizeAccess,
                @CachedLibrary(limit = "3") InteropLibrary gridSizeElementAccess,
                @CachedLibrary(limit = "3") InteropLibrary blockSizeAccess,
                @CachedLibrary(limit = "3") InteropLibrary blockSizeElementAccess,
                @CachedLibrary(limit = "3") InteropLibrary sharedMemoryAccess) throws UnsupportedTypeException, ArityException {
    assertKernelLoaded();
    return kernel.execute(arguments, gridSizeAccess, gridSizeElementAccess, blockSizeAccess, blockSizeElementAccess, sharedMemoryAccess);
}
 

@Override
public Object call(CUDARuntime cudaRuntime, Object[] args) throws ArityException, UnsupportedTypeException, InteropException {
    checkArgumentLength(args, 2);
    int attributeCode = expectInt(args[0]);
    int deviceId = expectInt(args[1]);
    try (UnsafeHelper.Integer32Object value = UnsafeHelper.createInteger32Object()) {
        callSymbol(cudaRuntime, value.getAddress(), attributeCode, deviceId);
        return value.getValue();
    }
}
 

@Override
public Object call(CUDARuntime cudaRuntime, Object[] args) throws ArityException, UnsupportedTypeException, InteropException {
    checkArgumentLength(args, 1);
    Object pointerObj = args[0];
    long addr;
    if (pointerObj instanceof GPUPointer) {
        addr = ((GPUPointer) pointerObj).getRawPointer();
    } else if (pointerObj instanceof LittleEndianNativeArrayView) {
        addr = ((LittleEndianNativeArrayView) pointerObj).getStartAddress();
    } else {
        throw new GrCUDAException("expected GPUPointer or LittleEndianNativeArrayView");
    }
    callSymbol(cudaRuntime, addr);
    return NoneValue.get();
}
 

@Override
@TruffleBoundary
public Object call(CUDARuntime cudaRuntime, Object[] args) throws ArityException, InteropException {
    checkArgumentLength(args, 0);
    try (UnsafeHelper.Integer32Object deviceId = UnsafeHelper.createInteger32Object()) {
        callSymbol(cudaRuntime, deviceId.getAddress());
        return deviceId.getValue();
    }
}
 

@Override
@TruffleBoundary
public Object call(CUDARuntime cudaRuntime, Object[] args) throws ArityException, InteropException {
    checkArgumentLength(args, 0);
    try (UnsafeHelper.Integer32Object deviceCount = UnsafeHelper.createInteger32Object()) {
        callSymbol(cudaRuntime, deviceCount.getAddress());
        return deviceCount.getValue();
    }
}