From bbf828473fb882dbb810198984d3fa6ea8509b30 Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Mon, 18 Mar 2024 19:28:42 +0100 Subject: [PATCH 1/8] ACQUIRE_BUFFERn / RELEASE_BUFFERn functionality required for EA31337/EA31337-indicators-stats#1 --- Std.h | 154 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 154 insertions(+) diff --git a/Std.h b/Std.h index a680db7d2..f46504330 100644 --- a/Std.h +++ b/Std.h @@ -298,3 +298,157 @@ inline _NULL_VALUE::operator const std::string() const { #else #define NULL_VALUE NULL #endif + +/** + * Standarization of specifying ArraySetAsSeries for OnCalculate(). + * Automatically determines required ArraySetAsSeries for buffers on start and + * end of given code block that uses given buffers. + */ + +#define SET_BUFFER_AS_SERIES_FOR_TARGET(A) ArraySetAsSeries(A, false); + +#ifdef __MQL4__ +#define SET_BUFFER_AS_SERIES_FOR_HOST(A) ArraySetAsSeries(A, true); +#else +#define SET_BUFFER_AS_SERIES_FOR_HOST(A) ArraySetAsSeries(A, false); +#endif + +// Ensures that we do RELEASE_BUFFERx after ACQUIRE_BUFFERx. +struct AsSeriesReleaseEnsurer { + bool released; + int num_buffs; + AsSeriesReleaseEnsurer(int _num_buffs) : released(false), num_buffs(_num_buffs) {} + void done(int _num_buffs) { + if (_num_buffs != num_buffs) { + Alert("You have acquired ", num_buffs, " buffers via ACQUIRE_BUFFER", num_buffs, + "(), but now trying to release with mismatched RELEASE_BUFFER", _num_buffs, "()!"); + DebugBreak(); + } + + if (released) { + Alert("You have used RELEASE_BUFFER", num_buffs, "() again which is not required!"); + DebugBreak(); + } + + released = true; + } + ~AsSeriesReleaseEnsurer() { + if (!released) { + Alert("You have used ACQUIRE_BUFFER", num_buffs, "() but didn't release buffer(s) via RELEASE_BUFFER", num_buffs, + "() before returning from the scope!"); + DebugBreak(); + } + } +}; + +#define SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(NUM_BUFFS) \ + AsSeriesReleaseEnsurer _as_series_release_ensurer(NUM_BUFFS); +#define SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(NUM_BUFFS) _as_series_release_ensurer.done(NUM_BUFFS); + +// Acquiring buffer is preparing it to be used as in MQL5. +#define ACQUIRE_BUFFER1(A) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(1); +#define ACQUIRE_BUFFER2(A, B) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(2); +#define ACQUIRE_BUFFER3(A, B, C) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(C); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(3); +#define ACQUIRE_BUFFER4(A, B, C, D) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(C); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(D); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(4); +#define ACQUIRE_BUFFER5(A, B, C, D, E) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(C); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(D); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(E); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(5); +#define ACQUIRE_BUFFER6(A, B, C, D, E, F) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(C); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(D); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(E); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(F); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(6); +#define ACQUIRE_BUFFER7(A, B, C, D, E, F, G) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(C); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(D); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(E); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(F); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(G); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(7); +#define ACQUIRE_BUFFER8(A, B, C, D, E, F, G, H) \ + SET_BUFFER_AS_SERIES_FOR_TARGET(A); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(B); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(C); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(D); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(E); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(F); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(G); \ + SET_BUFFER_AS_SERIES_FOR_TARGET(H); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_BEGIN(8); + +// Releasing buffer is setting its AsSeries as the default in the host language. +#define RELEASE_BUFFER1(A) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(1); +#define RELEASE_BUFFER2(A, B) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(2); +#define RELEASE_BUFFER3(A, B, C) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_FOR_HOST(C); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(3); +#define RELEASE_BUFFER4(A, B, C, D) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_FOR_HOST(C); \ + SET_BUFFER_AS_SERIES_FOR_HOST(D); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(4); +#define RELEASE_BUFFER5(A, B, C, D, E) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_FOR_HOST(C); \ + SET_BUFFER_AS_SERIES_FOR_HOST(D); \ + SET_BUFFER_AS_SERIES_FOR_HOST(E); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(5); +#define RELEASE_BUFFER6(A, B, C, D, E, F) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_FOR_HOST(C); \ + SET_BUFFER_AS_SERIES_FOR_HOST(D); \ + SET_BUFFER_AS_SERIES_FOR_HOST(E); \ + SET_BUFFER_AS_SERIES_FOR_HOST(F); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(6); +#define RELEASE_BUFFER7(A, B, C, D, E, F, G) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_FOR_HOST(C); \ + SET_BUFFER_AS_SERIES_FOR_HOST(D); \ + SET_BUFFER_AS_SERIES_FOR_HOST(E); \ + SET_BUFFER_AS_SERIES_FOR_HOST(F); \ + SET_BUFFER_AS_SERIES_FOR_HOST(G); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(7); +#define RELEASE_BUFFER8(A, B, C, D, E, F, G, H) \ + SET_BUFFER_AS_SERIES_FOR_HOST(A); \ + SET_BUFFER_AS_SERIES_FOR_HOST(B); \ + SET_BUFFER_AS_SERIES_FOR_HOST(C); \ + SET_BUFFER_AS_SERIES_FOR_HOST(D); \ + SET_BUFFER_AS_SERIES_FOR_HOST(E); \ + SET_BUFFER_AS_SERIES_FOR_HOST(F); \ + SET_BUFFER_AS_SERIES_FOR_HOST(G); \ + SET_BUFFER_AS_SERIES_FOR_HOST(H); \ + SET_BUFFER_AS_SERIES_RELEASE_ENSURER_END(8); From 20a08a785353e183de99145f46d973538fdac24d Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Fri, 22 Mar 2024 20:21:07 +0100 Subject: [PATCH 2/8] Refs EA31337/EA31337-indicators-stats#2. Closed to finish Indi_AccountStats indicator. --- Account/Account.h | 1 + Account/AccountBase.h | 40 ++++ Account/AccountMt.h | 15 +- Indicator.enum.h | 25 ++- Indicator/IndicatorCandle.h | 28 ++- Indicator/IndicatorTick.h | 2 +- Indicator/tests/classes/IndicatorTfDummy.h | 12 +- Indicator/tests/classes/IndicatorTickDummy.h | 16 +- IndicatorData.mqh | 10 +- Indicators/Account/Indi_AccountStats.mqh | 207 ++++++++++++++++++ .../Account/tests/Indi_AccountStats.test.mq4 | 27 +++ .../Account/tests/Indi_AccountStats.test.mq5 | 61 ++++++ Indicators/Indi_AMA.mqh | 8 - Indicators/Tick/Indi_TickMt.mqh | 6 +- Storage/IValueStorage.h | 5 + Storage/ValueStorage.h | 9 + Storage/ValueStorage.native.h | 19 +- 17 files changed, 444 insertions(+), 47 deletions(-) create mode 100644 Indicators/Account/Indi_AccountStats.mqh create mode 100644 Indicators/Account/tests/Indi_AccountStats.test.mq4 create mode 100644 Indicators/Account/tests/Indi_AccountStats.test.mq5 diff --git a/Account/Account.h b/Account/Account.h index 5e50510d5..a42e5822c 100644 --- a/Account/Account.h +++ b/Account/Account.h @@ -56,4 +56,5 @@ class Account : public AccountBase { */ ~Account() {} }; + #endif // ACCOUNT_H diff --git a/Account/AccountBase.h b/Account/AccountBase.h index 9d3ae7253..187649570 100644 --- a/Account/AccountBase.h +++ b/Account/AccountBase.h @@ -56,6 +56,46 @@ class AccountBase : public Dynamic { * Class deconstructor. */ ~AccountBase() {} + + /** + * Returns balance value of the current account. + */ + virtual datetime GetDateTime() { return TimeCurrent(); }; + + /** + * Returns balance value of the current account. + */ + virtual float GetBalance() = 0; + + /** + * Returns credit value of the current account. + */ + virtual float GetCredit() = 0; + + /** + * Returns profit value of the current account. + */ + virtual float GetProfit() = 0; + + /** + * Returns equity value of the current account. + */ + virtual float GetEquity() = 0; + + /** + * Returns margin value of the current account. + */ + virtual float GetMarginUsed() = 0; + + /** + * Returns free margin value of the current account. + */ + virtual float GetMarginFree() = 0; + + /** + * Get account available margin. + */ + virtual float GetMarginAvail() = 0; }; #endif // ACCOUNTBASE_H diff --git a/Account/AccountMt.h b/Account/AccountMt.h index c03c300f9..66f053591 100644 --- a/Account/AccountMt.h +++ b/Account/AccountMt.h @@ -42,12 +42,13 @@ class AccountMt; #include "Account.define.h" #include "Account.enum.h" #include "Account.extern.h" +#include "Account.h" #include "Account.struct.h" /** * Class to provide functions that return parameters of the current account. */ -class AccountMt { +class AccountMt : public AccountBase { protected: // Struct variables. BufferStruct entries; @@ -136,7 +137,7 @@ class AccountMt { * Returns balance value of the current account. */ static double AccountBalance() { return AccountInfoDouble(ACCOUNT_BALANCE); } - float GetBalance() { + float GetBalance() override { // @todo: Adds caching. // return UpdateStats(ACC_BALANCE, AccountBalance()); return (float)AccountMt::AccountBalance(); @@ -146,7 +147,7 @@ class AccountMt { * Returns credit value of the current account. */ static double AccountCredit() { return AccountInfoDouble(ACCOUNT_CREDIT); } - float GetCredit() { + float GetCredit() override { // @todo: Adds caching. // return UpdateStats(ACC_CREDIT, AccountCredit()); return (float)AccountMt::AccountCredit(); @@ -156,7 +157,7 @@ class AccountMt { * Returns profit value of the current account. */ static double AccountProfit() { return AccountInfoDouble(ACCOUNT_PROFIT); } - float GetProfit() { + float GetProfit() override { // @todo: Adds caching. // return UpdateStats(ACC_PROFIT, AccountProfit()); return (float)AccountMt::AccountProfit(); @@ -166,7 +167,7 @@ class AccountMt { * Returns equity value of the current account. */ static double AccountEquity() { return AccountInfoDouble(ACCOUNT_EQUITY); } - float GetEquity() { + float GetEquity() override { // @todo: Adds caching. // return UpdateStats(ACC_EQUITY, AccountEquity()); return (float)AccountMt::AccountEquity(); @@ -198,7 +199,7 @@ class AccountMt { * Returns free margin value of the current account. */ static double AccountFreeMargin() { return AccountInfoDouble(ACCOUNT_MARGIN_FREE); } - float GetMarginFree() { + float GetMarginFree() override { // @todo: Adds caching. // return UpdateStats(ACC_MARGIN_FREE, AccountFreeMargin()); return (float)AccountMt::AccountFreeMargin(); @@ -267,7 +268,7 @@ class AccountMt { * Get account available margin. */ static double AccountAvailMargin() { return fmin(AccountFreeMargin(), AccountTotalBalance()); } - float GetMarginAvail() { return (float)AccountAvailMargin(); } + float GetMarginAvail() override { return (float)AccountAvailMargin(); } /** * Returns the calculation mode of free margin allowed to open orders on the current account. diff --git a/Indicator.enum.h b/Indicator.enum.h index b2b07a3a4..c20b4e8b6 100644 --- a/Indicator.enum.h +++ b/Indicator.enum.h @@ -44,6 +44,7 @@ enum ENUM_INDICATOR_ACTION { enum ENUM_INDICATOR_TYPE { INDI_NONE = 0, // (None) INDI_AC, // Accelerator Oscillator + INDI_ACCOUNT_STATS, // Account Stats INDI_AD, // Accumulation/Distribution INDI_ADX, // Average Directional Index INDI_ADXW, // ADX by Welles Wilder @@ -237,9 +238,24 @@ enum ENUM_INDI_VS_TYPE { INDI_VS_TYPE_INDEX_8, INDI_VS_TYPE_INDEX_9, INDI_VS_TYPE_INDEX_FIRST = INDI_VS_TYPE_INDEX_0, - INDI_VS_TYPE_INDEX_LAST = INDI_VS_TYPE_INDEX_9 + INDI_VS_TYPE_INDEX_LAST = INDI_VS_TYPE_INDEX_9, + + // Account Stats. + INDI_VS_TYPE_ACCOUNT_STATS_DATE_TIME, + INDI_VS_TYPE_ACCOUNT_STATS_BALANCE, + INDI_VS_TYPE_ACCOUNT_STATS_CREDIT, + INDI_VS_TYPE_ACCOUNT_STATS_EQUITY, + INDI_VS_TYPE_ACCOUNT_STATS_PROFIT, + INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_USED, + INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_FREE, + INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_AVAIL, + INDI_VS_TYPE_ACCOUNT_STATS_INDEX_FIRST = INDI_VS_TYPE_ACCOUNT_STATS_DATE_TIME, + INDI_VS_TYPE_ACCOUNT_STATS_INDEX_LAST = INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_AVAIL, }; +#define INDI_VS_TYPE_ACCOUNT_STATS_BUFFERS_COUNT \ + (INDI_VS_TYPE_ACCOUNT_STATS_INDEX_LAST - INDI_VS_TYPE_ACCOUNT_STATS_INDEX_FIRST + 1) + // Indicator flags. enum ENUM_INDI_FLAGS { INDI_FLAG_INDEXABLE_BY_SHIFT, // Indicator supports indexation by shift. @@ -269,4 +285,11 @@ enum ENUM_INDI_DS_MODE_KIND { INDI_DS_MODE_KIND_AP, // Mode is a value from ENUM_APPLIED_PRICE enumeration. It is used to retrieve value storage // based on ENUM_INDI_VS_TYPE enumeration, e.g., PRICE_OPEN becomes ENUM_INDI_VS_PRICE_OPEN. }; + +// Type of entry +enum ENUM_INDI_EMITTED_ENTRY_TYPE { + INDI_EMITTED_ENTRY_TYPE_PARENT, // Undetermined type of entry from direct parent indicator. + INDI_EMITTED_ENTRY_TYPE_TICK, + INDI_EMITTED_ENTRY_TYPE_CANDLE, +}; //+------------------------------------------------------------------+ diff --git a/Indicator/IndicatorCandle.h b/Indicator/IndicatorCandle.h index 8221afe82..da34844ca 100644 --- a/Indicator/IndicatorCandle.h +++ b/Indicator/IndicatorCandle.h @@ -32,6 +32,7 @@ // Includes. #include "../Buffer/BufferCandle.h" #include "../Candle.struct.h" +#include "../Indicator.enum.h" #include "../Indicator.mqh" #include "../Storage/ValueStorage.price_median.h" #include "../Storage/ValueStorage.price_typical.h" @@ -310,7 +311,7 @@ class IndicatorCandle : public Indicator { void EmitHistory() override { for (DictStructIterator> iter(icdata.Begin()); iter.IsValid(); ++iter) { IndicatorDataEntry _entry = CandleToEntry(iter.Key(), iter.Value()); - EmitEntry(_entry); + EmitEntry(_entry, INDI_EMITTED_ENTRY_TYPE_CANDLE); } } @@ -342,7 +343,7 @@ class IndicatorCandle : public Indicator { /** * Adds tick's price to the matching candle and updates its OHLC values. */ - void UpdateCandle(long _tick_timestamp, double _price) { + CandleOCTOHLC UpdateCandle(long _tick_timestamp, double _price) { long _candle_timestamp = CalcCandleTimestamp(_tick_timestamp); #ifdef __debug_verbose__ @@ -368,6 +369,8 @@ class IndicatorCandle : public Indicator { } icdata.Add(_candle, _candle_timestamp); + + return _candle; } /** @@ -380,12 +383,25 @@ class IndicatorCandle : public Indicator { /** * Called when data source emits new entry (historic or future one). */ - void OnDataSourceEntry(IndicatorDataEntry& entry) override { - // Updating candle from bid price. - UpdateCandle(entry.timestamp, entry[1]); + void OnDataSourceEntry(IndicatorDataEntry& entry, + ENUM_INDI_EMITTED_ENTRY_TYPE type = INDI_EMITTED_ENTRY_TYPE_PARENT) override { + Indicator::OnDataSourceEntry(entry, type); + + if (type != INDI_EMITTED_ENTRY_TYPE_TICK) { + return; + } + + long _candle_timestamp = CalcCandleTimestamp(entry.timestamp); // Updating tick & bar indices. - counter.OnTick(CalcCandleTimestamp(entry.timestamp)); + counter.OnTick(_candle_timestamp); + + // Updating candle from bid price. + CandleOCTOHLC _candle = UpdateCandle(entry.timestamp, entry[1]); + + // Emitting candle for children. + IndicatorDataEntry _candle_entry = CandleToEntry(_candle_timestamp, _candle); + EmitEntry(_candle_entry, INDI_EMITTED_ENTRY_TYPE_CANDLE); }; /** diff --git a/Indicator/IndicatorTick.h b/Indicator/IndicatorTick.h index 5438923cb..cf01dfe83 100644 --- a/Indicator/IndicatorTick.h +++ b/Indicator/IndicatorTick.h @@ -164,7 +164,7 @@ class IndicatorTick : public Indicator { void EmitHistory() override { for (DictStructIterator> iter(itdata.Begin()); iter.IsValid(); ++iter) { IndicatorDataEntry _entry = TickToEntry(iter.Key(), iter.Value()); - EmitEntry(_entry); + EmitEntry(_entry, INDI_EMITTED_ENTRY_TYPE_TICK); } } diff --git a/Indicator/tests/classes/IndicatorTfDummy.h b/Indicator/tests/classes/IndicatorTfDummy.h index 019b800d8..f49a8161d 100644 --- a/Indicator/tests/classes/IndicatorTfDummy.h +++ b/Indicator/tests/classes/IndicatorTfDummy.h @@ -48,11 +48,13 @@ class IndicatorTfDummy : public IndicatorTf { string GetName() override { return "IndicatorTfDummy(" + IntegerToString(iparams.spc) + ")"; } - void OnDataSourceEntry(IndicatorDataEntry& entry) override { - // When overriding OnDataSourceEntry() we have to remember to call parent - // method, because IndicatorCandle also need to invoke it in order to - // create/update matching candle. - IndicatorTf::OnDataSourceEntry(entry); + void OnDataSourceEntry(IndicatorDataEntry& entry, + ENUM_INDI_EMITTED_ENTRY_TYPE type = INDI_EMITTED_ENTRY_TYPE_PARENT) override { + IndicatorTf::OnDataSourceEntry(entry, type); + + if (type != INDI_EMITTED_ENTRY_TYPE_TICK) { + return; + } #ifdef __debug_indicator__ Print(GetFullName(), " got new tick at ", entry.timestamp, diff --git a/Indicator/tests/classes/IndicatorTickDummy.h b/Indicator/tests/classes/IndicatorTickDummy.h index dfd59e9a3..3e6bab0ab 100644 --- a/Indicator/tests/classes/IndicatorTickDummy.h +++ b/Indicator/tests/classes/IndicatorTickDummy.h @@ -59,13 +59,13 @@ class IndicatorTickDummy : public IndicatorTick _t7(4.2f, 4.21f); TickAB _t8(4.8f, 4.81f); - EmitEntry(TickToEntry(1000, _t1)); - EmitEntry(TickToEntry(1500, _t2)); - EmitEntry(TickToEntry(2000, _t3)); - EmitEntry(TickToEntry(3000, _t4)); - EmitEntry(TickToEntry(4000, _t5)); - EmitEntry(TickToEntry(4100, _t6)); - EmitEntry(TickToEntry(4200, _t7)); - EmitEntry(TickToEntry(4800, _t8)); + EmitEntry(TickToEntry(1000, _t1), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(1500, _t2), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(2000, _t3), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(3000, _t4), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(4000, _t5), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(4100, _t6), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(4200, _t7), INDI_EMITTED_ENTRY_TYPE_TICK); + EmitEntry(TickToEntry(4800, _t8), INDI_EMITTED_ENTRY_TYPE_TICK); }; }; diff --git a/IndicatorData.mqh b/IndicatorData.mqh index 46225b8d3..bc0c498b1 100644 --- a/IndicatorData.mqh +++ b/IndicatorData.mqh @@ -1616,10 +1616,10 @@ class IndicatorData : public IndicatorBase { /** * Sends entry to listening indicators. */ - void EmitEntry(IndicatorDataEntry& entry) { + void EmitEntry(IndicatorDataEntry& entry, ENUM_INDI_EMITTED_ENTRY_TYPE type = INDI_EMITTED_ENTRY_TYPE_PARENT) { for (int i = 0; i < ArraySize(listeners); ++i) { if (listeners[i].ObjectExists()) { - listeners[i].Ptr().OnDataSourceEntry(entry); + listeners[i].Ptr().OnDataSourceEntry(entry, type); } } } @@ -1683,7 +1683,11 @@ class IndicatorData : public IndicatorBase { /** * Called when data source emits new entry (historic or future one). */ - virtual void OnDataSourceEntry(IndicatorDataEntry& entry){}; + virtual void OnDataSourceEntry(IndicatorDataEntry& entry, + ENUM_INDI_EMITTED_ENTRY_TYPE type = INDI_EMITTED_ENTRY_TYPE_PARENT) { + // Sending entry to all chilren listeners (from highest parent to lowest child). + EmitEntry(entry, type); + }; virtual void OnTick() {} diff --git a/Indicators/Account/Indi_AccountStats.mqh b/Indicators/Account/Indi_AccountStats.mqh new file mode 100644 index 000000000..da45fef7d --- /dev/null +++ b/Indicators/Account/Indi_AccountStats.mqh @@ -0,0 +1,207 @@ +//+------------------------------------------------------------------+ +//| EA31337 framework | +//| Copyright 2016-2021, EA31337 Ltd | +//| https://github.com/EA31337 | +//+------------------------------------------------------------------+ + +/* + * This file is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + * + */ + +// Includes. +#include "../../Account/AccountBase.h" +#include "../../BufferStruct.mqh" +#include "../../Indicator.mqh" +#include "../../Platform.h" +#include "../../Storage/Objects.h" + +// Structs. +struct Indi_AccountStats_Params : IndicatorParams { + // Applied price. + ENUM_APPLIED_PRICE ap; + + // Account to use. + Ref account; + + // Struct constructor. + Indi_AccountStats_Params(AccountBase *_account = nullptr, int _shift = 0) + : IndicatorParams(INDI_ACCOUNT_STATS), account(_account) { + SetShift(_shift); + }; + Indi_AccountStats_Params(Indi_AccountStats_Params &_params) { THIS_REF = _params; }; + + // Getters. + AccountBase *GetAccount() { return account.Ptr(); } + ENUM_APPLIED_PRICE GetAppliedPrice() override { return ap; } + + // Setters. + void SetAccount(AccountBase *_account) { account = _account; } + void SetAppliedPrice(ENUM_APPLIED_PRICE _ap) { ap = _ap; } +}; + +/** + * Price Indicator. + */ +class Indi_AccountStats : public Indicator { + Ref> buffer_date_time; + Ref> buffer_balance; + Ref> buffer_credit; + Ref> buffer_equity; + Ref> buffer_profit; + Ref> buffer_margin_used; + Ref> buffer_margin_free; + Ref> buffer_margin_avail; + + public: + /** + * Class constructor. + */ + Indi_AccountStats(Indi_AccountStats_Params &_p, ENUM_IDATA_SOURCE_TYPE _idstype = IDATA_BUILTIN, + IndicatorData *_indi_src = NULL, int _indi_src_mode = 0) + : Indicator(_p, + IndicatorDataParams::GetInstance(INDI_VS_TYPE_ACCOUNT_STATS_BUFFERS_COUNT, TYPE_DOUBLE, _idstype, + IDATA_RANGE_PRICE, _indi_src_mode), + _indi_src) { + InitAccountStats(); + }; + Indi_AccountStats(int _shift = 0, ENUM_IDATA_SOURCE_TYPE _idstype = IDATA_BUILTIN, IndicatorData *_indi_src = NULL, + int _indi_src_mode = 0) + : Indicator(Indi_AccountStats_Params(), + IndicatorDataParams::GetInstance(INDI_VS_TYPE_ACCOUNT_STATS_BUFFERS_COUNT, TYPE_DOUBLE, _idstype, + IDATA_RANGE_PRICE, _indi_src_mode), + _indi_src) { + InitAccountStats(); + }; + void InitAccountStats() { + buffer_date_time = new NativeValueStorage(); + buffer_balance = new NativeValueStorage(); + buffer_credit = new NativeValueStorage(); + buffer_equity = new NativeValueStorage(); + buffer_profit = new NativeValueStorage(); + buffer_margin_used = new NativeValueStorage(); + buffer_margin_free = new NativeValueStorage(); + buffer_margin_avail = new NativeValueStorage(); + } + + /** + * Returns possible data source types. It is a bit mask of ENUM_INDI_SUITABLE_DS_TYPE. + */ + virtual unsigned int GetSuitableDataSourceTypes() { + // We require that candle indicator is attached. + return INDI_SUITABLE_DS_TYPE_CANDLE; + } + + /** + * Returns possible data source modes. It is a bit mask of ENUM_IDATA_SOURCE_TYPE. + */ + unsigned int GetPossibleDataModes() override { return IDATA_BUILTIN; } + + /** + * Checks whether indicator has a valid value for a given shift. + */ + virtual bool HasValidEntry(int _shift = 0) { return GetBarTime(_shift) != 0; } + + /** + * Returns the indicator's value. + */ + virtual IndicatorDataEntryValue GetEntryValue(int _mode = 0, int _shift = -1) { + int _ishift = _shift >= 0 ? _shift : iparams.GetShift(); + + // Converting mode into value storage type. + ENUM_INDI_VS_TYPE _vs_type = (ENUM_INDI_VS_TYPE)(INDI_VS_TYPE_ACCOUNT_STATS_INDEX_FIRST + _mode); + + // Retrieving data from specific value storage. + switch (_vs_type) { + case INDI_VS_TYPE_ACCOUNT_STATS_DATE_TIME: + return ((ValueStorage *)GetSpecificValueStorage(_vs_type))PTR_DEREF FetchSeries(_ishift); + case INDI_VS_TYPE_ACCOUNT_STATS_BALANCE: + case INDI_VS_TYPE_ACCOUNT_STATS_CREDIT: + case INDI_VS_TYPE_ACCOUNT_STATS_EQUITY: + case INDI_VS_TYPE_ACCOUNT_STATS_PROFIT: + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_USED: + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_FREE: + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_AVAIL: + return ((ValueStorage *)GetSpecificValueStorage(_vs_type))PTR_DEREF FetchSeries(_ishift); + default: + Alert("Error: Indi_AccountStats: Invalid mode passed to GetEntryValue()!"); + DebugBreak(); + return EMPTY_VALUE; + } + } + + /** + * Returns value storage of given kind. + */ + IValueStorage *GetSpecificValueStorage(ENUM_INDI_VS_TYPE _type) override { + // Returning Price indicator which provides applied price in the only buffer #0. + switch (_type) { + case INDI_VS_TYPE_ACCOUNT_STATS_DATE_TIME: + return buffer_date_time.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_BALANCE: + return buffer_balance.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_CREDIT: + return buffer_credit.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_EQUITY: + return buffer_equity.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_PROFIT: + return buffer_profit.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_USED: + return buffer_margin_used.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_FREE: + return buffer_margin_free.Ptr(); + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_AVAIL: + return buffer_margin_avail.Ptr(); + default: + // Trying in parent class. + return Indicator::GetSpecificValueStorage(_type); + } + } + + /** + * Checks whether indicator support given value storage type. + */ + bool HasSpecificValueStorage(ENUM_INDI_VS_TYPE _type) override { + switch (_type) { + case INDI_VS_TYPE_ACCOUNT_STATS_DATE_TIME: + case INDI_VS_TYPE_ACCOUNT_STATS_BALANCE: + case INDI_VS_TYPE_ACCOUNT_STATS_CREDIT: + case INDI_VS_TYPE_ACCOUNT_STATS_EQUITY: + case INDI_VS_TYPE_ACCOUNT_STATS_PROFIT: + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_USED: + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_FREE: + case INDI_VS_TYPE_ACCOUNT_STATS_MARGIN_AVAIL: + return true; + default: + // Trying in parent class. + return Indicator::HasSpecificValueStorage(_type); + } + } + + /** + * Called when data source emits new entry (historic or future one). + */ + virtual void OnDataSourceEntry(IndicatorDataEntry &entry, + ENUM_INDI_EMITTED_ENTRY_TYPE type = INDI_EMITTED_ENTRY_TYPE_PARENT) { + Indicator::OnDataSourceEntry(entry, type); + + if (type != INDI_EMITTED_ENTRY_TYPE_CANDLE) { + return; + } + + // Adding new account stats entry. + + Print("New candle: ", entry.ToString()); + } +}; diff --git a/Indicators/Account/tests/Indi_AccountStats.test.mq4 b/Indicators/Account/tests/Indi_AccountStats.test.mq4 new file mode 100644 index 000000000..8356c50b5 --- /dev/null +++ b/Indicators/Account/tests/Indi_AccountStats.test.mq4 @@ -0,0 +1,27 @@ +//+------------------------------------------------------------------+ +//| EA31337 framework | +//| Copyright 2016-2021, EA31337 Ltd | +//| https://github.com/EA31337 | +//+------------------------------------------------------------------+ + +/* + * This file is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * (at your option) any later version. + + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + */ + +/** + * @file + * Test functionality of Indi_AC indicator class. + */ + +#include "Indi_AC.test.mq5" diff --git a/Indicators/Account/tests/Indi_AccountStats.test.mq5 b/Indicators/Account/tests/Indi_AccountStats.test.mq5 new file mode 100644 index 000000000..c81e0da79 --- /dev/null +++ b/Indicators/Account/tests/Indi_AccountStats.test.mq5 @@ -0,0 +1,61 @@ +//+------------------------------------------------------------------+ +//| EA31337 framework | +//| Copyright 2016-2021, EA31337 Ltd | +//| https://github.com/EA31337 | +//+------------------------------------------------------------------+ + +/* + * This file is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * (at your option) any later version. + + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + */ + +// Includes. +#include "../../../Account/AccountMt.h" +#include "../../../Platform.h" +#include "../../../Test.mqh" +#include "../Indi_AccountStats.mqh" + +/** + * @file + * Test functionality of Indi_AccountStats indicator class. + */ + +Ref indi_account_mt; + +int OnInit() { + Ref account_mt = new AccountMt(); + Indi_AccountStats_Params indi_params(account_mt.Ptr()); + indi_account_mt = new Indi_AccountStats(indi_params); + + Platform::Init(); + + Platform::AddWithDefaultBindings(indi_account_mt.Ptr()); + + bool _result = true; + assertTrueOrFail(indi_account_mt REF_DEREF IsValid(), "Error on IsValid!"); + return (_result && _LastError == ERR_NO_ERROR ? INIT_SUCCEEDED : INIT_FAILED); +} + +void OnTick() { + Platform::Tick(); + if (Platform::IsNewHour()) { + IndicatorDataEntry _entry = indi_account_mt REF_DEREF GetEntry(); + bool _is_ready = indi_account_mt REF_DEREF Get(STRUCT_ENUM(IndicatorState, INDICATOR_STATE_PROP_IS_READY)); + bool _is_valid = _entry.IsValid(); + Print(indi_account_mt REF_DEREF ToString(), _is_ready ? "" : " (Not yet ready)"); + if (_is_ready && !_is_valid) { + Print(indi_account_mt REF_DEREF ToString(), " (Invalid entry!)"); + assertTrueOrExit(_entry.IsValid(), "Invalid entry!"); + } + } +} \ No newline at end of file diff --git a/Indicators/Indi_AMA.mqh b/Indicators/Indi_AMA.mqh index ec615c8ac..4fde34ef7 100644 --- a/Indicators/Indi_AMA.mqh +++ b/Indicators/Indi_AMA.mqh @@ -256,14 +256,6 @@ class Indi_AMA : public Indicator { return _value; } - /** - * Called when data source emits new entry (historic or future one). - */ - void OnDataSourceEntry(IndicatorDataEntry &entry) override { - // Just to be able to make a breakpoint here. - int x = 4; - }; - /** * Called if data source is requested, but wasn't yet set. May be used to initialize indicators that must operate on * some data source. diff --git a/Indicators/Tick/Indi_TickMt.mqh b/Indicators/Tick/Indi_TickMt.mqh index 726004fff..b60ee2d04 100644 --- a/Indicators/Tick/Indi_TickMt.mqh +++ b/Indicators/Tick/Indi_TickMt.mqh @@ -159,7 +159,7 @@ class Indi_TickMt : public IndicatorTick { _tmp_ticks[i].ask, ", ", _tmp_ticks[i].bid); #endif - EmitEntry(TickToEntry(_tmp_ticks[i].time, _tick)); + EmitEntry(TickToEntry(_tmp_ticks[i].time, _tick), INDI_EMITTED_ENTRY_TYPE_TICK); if (_num_yet_to_copy <= 0) { break; @@ -205,7 +205,7 @@ class Indi_TickMt : public IndicatorTick { // DebugBreak(); // Just emitting zeroes in case of error. TickAB _tick(0, 0); - EmitEntry(TickToEntry(TimeCurrent(), _tick)); + EmitEntry(TickToEntry(TimeCurrent(), _tick), INDI_EMITTED_ENTRY_TYPE_TICK); return; } @@ -224,6 +224,6 @@ class Indi_TickMt : public IndicatorTick { TickAB _tick(_ask, _bid); IndicatorDataEntry _entry(TickToEntry(_time, _tick)); StoreEntry(_entry); - EmitEntry(_entry); + EmitEntry(_entry, INDI_EMITTED_ENTRY_TYPE_TICK); } }; diff --git a/Storage/IValueStorage.h b/Storage/IValueStorage.h index 290c2552d..6f4897bf6 100644 --- a/Storage/IValueStorage.h +++ b/Storage/IValueStorage.h @@ -72,6 +72,11 @@ class IValueStorage : public Dynamic { DebugBreak(); return false; } + + /** + * Returns real array index for this given shift. + **/ + virtual int GetRealIndex(int _shift) { return IsSeries() ? (Size() - 1 - _shift) : _shift; } }; /** diff --git a/Storage/ValueStorage.h b/Storage/ValueStorage.h index 6c7979ba0..865b2d651 100644 --- a/Storage/ValueStorage.h +++ b/Storage/ValueStorage.h @@ -181,6 +181,15 @@ class ValueStorage : public IValueStorage { DebugBreak(); } + /** + * Inserts new value at the end of the buffer. If buffer works as As-Series, + * then new value will act as the one at index 0. + */ + virtual void Append(C _value) { + Alert(__FUNCSIG__, " does not implement Append()!"); + DebugBreak(); + } + /** * Sets buffer drawing attributes. Currently does nothing. */ diff --git a/Storage/ValueStorage.native.h b/Storage/ValueStorage.native.h index 46ca12565..0c6180881 100644 --- a/Storage/ValueStorage.native.h +++ b/Storage/ValueStorage.native.h @@ -60,19 +60,28 @@ class NativeValueStorage : public ValueStorage { * Fetches value from a given shift. Takes into consideration as-series flag. */ C Fetch(int _shift) override { - if (_shift < 0 || _shift >= ArraySize(_values)) { + if (_shift < 0 || _shift >= Size()) { return (C)EMPTY_VALUE; - // Print("Invalid buffer data index: ", _shift, ". Buffer size: ", ArraySize(_values)); - // DebugBreak(); } - return _values[_shift]; + int _index = GetRealIndex(_shift); + + return _values[_index]; } /** * Stores value at a given shift. Takes into consideration as-series flag. */ - void Store(int _shift, C _value) override { Array::ArrayStore(_values, _shift, _value, 4096); } + void Store(int _shift, C _value) override { + if (_shift < 0 || _shift >= Size()) { + Alert("Error: NativeValueStorage: Invalid buffer data index: ", _shift, ". Buffer size: ", Size()); + DebugBreak(); + } + + int _index = GetRealIndex(_shift); + + Array::ArrayStore(_values, _index, _value, 4096); + } /** * Returns number of values available to fetch (size of the values buffer). From 4e763ae798519f18c02e1983dbdcd2647e05c49e Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Wed, 27 Mar 2024 15:13:43 +0100 Subject: [PATCH 3/8] WIP. Refs EA31337/EA31337-indicators-stats#2. Indi_AccountStats closer to be finished. Awaiting to fix MT5 compiler bug. --- IndicatorData.mqh | 48 ++++++++++--------- Indicators/Account/Indi_AccountStats.mqh | 26 ++++++++-- .../Account/tests/Indi_AccountStats.test.mq5 | 2 +- Storage/ValueStorage.native.h | 9 ++++ 4 files changed, 58 insertions(+), 27 deletions(-) diff --git a/IndicatorData.mqh b/IndicatorData.mqh index bc0c498b1..990c2e1f3 100644 --- a/IndicatorData.mqh +++ b/IndicatorData.mqh @@ -68,7 +68,7 @@ class IndicatorData : public IndicatorBase { case IDATA_ICUSTOM: break; case IDATA_INDICATOR: - if (indi_src.IsSet() == NULL) { + if (indi_src.IsSet()) { // Indi_Price* _indi_price = Indi_Price::GetCached(GetSymbol(), GetTf(), iparams.GetShift()); // SetDataSource(_indi_price, true, PRICE_OPEN); } @@ -109,7 +109,7 @@ class IndicatorData : public IndicatorBase { /** * Class constructor. */ - IndicatorData(const IndicatorDataParams& _idparams, IndicatorBase* _indi_src = NULL, int _indi_mode = 0) + IndicatorData(const IndicatorDataParams& _idparams, IndicatorBase* _indi_src = nullptr, int _indi_mode = 0) : do_draw(false), idparams(_idparams), indi_src(_indi_src) { Init(); } @@ -572,9 +572,9 @@ class IndicatorData : public IndicatorBase { * Returns currently selected data source doing validation. */ IndicatorData* GetDataSource(bool _validate = true) { - IndicatorData* _result = NULL; + IndicatorData* _result = nullptr; - if (GetDataSourceRaw() != NULL) { + if (GetDataSourceRaw() != nullptr) { _result = GetDataSourceRaw(); } else if (Get(STRUCT_ENUM(IndicatorDataParams, IDATA_PARAM_SRC_ID)) != -1) { int _source_id = Get(STRUCT_ENUM(IndicatorDataParams, IDATA_PARAM_SRC_ID)); @@ -603,7 +603,7 @@ class IndicatorData : public IndicatorBase { // Requesting potential data source. _result = OnDataSourceRequest(); - if (_result != NULL) { + if (_result != nullptr) { // Initializing with new data source. SetDataSource(_result); Set(STRUCT_ENUM(IndicatorDataParams, IDATA_PARAM_IDSTYPE), IDATA_INDICATOR); @@ -621,7 +621,7 @@ class IndicatorData : public IndicatorBase { * Returns given data source type. Used by i*OnIndicator methods if indicator's Calculate() uses other indicators. */ IndicatorData* GetDataSource(ENUM_INDICATOR_TYPE _type) { - IndicatorData* _result = NULL; + IndicatorData* _result = nullptr; if (indicators.KeyExists((int)_type)) { _result = indicators[(int)_type].Ptr(); } else { @@ -697,11 +697,15 @@ class IndicatorData : public IndicatorBase { } if (Get(STRUCT_ENUM(IndicatorDataParams, IDATA_PARAM_IDSTYPE)) == IDATA_INDICATOR && - GetDataSourceRaw() == NULL && _try_initialize) { + GetDataSourceRaw() == nullptr && _try_initialize) { SetDataSource(OnDataSourceRequest()); } - return GetDataSourceRaw() != NULL; + IndicatorData* _ptr = GetDataSourceRaw(); + + bool _result = _ptr != nullptr; + + return _result; } /** @@ -802,7 +806,7 @@ class IndicatorData : public IndicatorBase { indi_src.Ptr().RemoveListener(THIS_PTR); } indi_src = _indi; - if (_indi != NULL) { + if (_indi != nullptr) { indi_src.Ptr().AddListener(THIS_PTR); Set(STRUCT_ENUM(IndicatorDataParams, IDATA_PARAM_SRC_ID), -1); Set(STRUCT_ENUM(IndicatorDataParams, IDATA_PARAM_SRC_MODE), _input_mode); @@ -862,7 +866,7 @@ class IndicatorData : public IndicatorBase { last_tick_time = _current_time; // Checking and potentially initializing new data source. - if (HasDataSource(true) != NULL) { + if (HasDataSource(true)) { // Ticking data source if not yet ticked. GetDataSource().Tick(); } @@ -905,13 +909,13 @@ class IndicatorData : public IndicatorBase { * Loads and validates built-in indicators whose can be used as data source. */ void ValidateDataSource(IndicatorData* _target, IndicatorData* _source) { - if (_target == NULL) { + if (_target == nullptr) { Alert("Internal Error! _target is NULL in ", __FUNCTION_LINE__, "."); DebugBreak(); return; } - if (_source == NULL) { + if (_source == nullptr) { Alert("Error! You have to select source indicator's via SetDataSource()."); DebugBreak(); return; @@ -1129,7 +1133,7 @@ class IndicatorData : public IndicatorBase { _originator PTR_DEREF GetFullName(), "!"); DebugBreak(); } - return NULL; + return nullptr; } } @@ -1146,7 +1150,7 @@ class IndicatorData : public IndicatorBase { /** * Returns the indicator's struct value via index. */ - virtual IndicatorDataEntry GetEntry(long _index = 0) = NULL; + virtual IndicatorDataEntry GetEntry(long _index = 0) = 0; /** * Returns the indicator's struct value via timestamp. @@ -1171,7 +1175,7 @@ class IndicatorData : public IndicatorBase { /** * Returns the indicator's entry value. */ - virtual IndicatorDataEntryValue GetEntryValue(int _mode = 0, int _shift = 0) = NULL; + virtual IndicatorDataEntryValue GetEntryValue(int _mode = 0, int _shift = 0) = 0; /** * Returns the shift of the maximum value over a specific number of periods depending on type. @@ -1238,7 +1242,7 @@ class IndicatorData : public IndicatorBase { * * When indicator values are not valid, returns empty signals. */ - virtual IndicatorSignal GetSignals(int _count = 3, int _shift = 0, int _mode1 = 0, int _mode2 = 0) = NULL; + virtual IndicatorSignal GetSignals(int _count = 3, int _shift = 0, int _mode1 = 0, int _mode2 = 0) = 0; /** * Returns spread for the bar. @@ -1331,7 +1335,7 @@ class IndicatorData : public IndicatorBase { ", only PRICE_(OPEN|HIGH|LOW|CLOSE|MEDIAN|TYPICAL|WEIGHTED) are currently supported by " "IndicatorBase::GetSpecificAppliedPriceValueStorage()!"); DebugBreak(); - return NULL; + return nullptr; } } @@ -1374,7 +1378,7 @@ class IndicatorData : public IndicatorBase { "Volume) in the hierarchy!"); DebugBreak(); } - return NULL; + return nullptr; } /** @@ -1390,7 +1394,7 @@ class IndicatorData : public IndicatorBase { virtual IValueStorage* GetSpecificValueStorage(ENUM_INDI_VS_TYPE _type) { Print("Error: ", GetFullName(), " indicator has no storage type ", EnumToString(_type), "!"); DebugBreak(); - return NULL; + return nullptr; } /** @@ -1632,7 +1636,7 @@ class IndicatorData : public IndicatorBase { /** * Provides built-in indicators whose can be used as data source. */ - virtual IndicatorData* FetchDataSource(ENUM_INDICATOR_TYPE _id) { return NULL; } + virtual IndicatorData* FetchDataSource(ENUM_INDICATOR_TYPE _id) { return nullptr; } /** * Checks whether indicator support given value storage type. @@ -1700,7 +1704,7 @@ class IndicatorData : public IndicatorBase { " without explicitly selecting an indicator, ", GetFullName(), " must override OnDataSourceRequest() method and return new instance of data source to be used by default."); DebugBreak(); - return NULL; + return nullptr; } /** @@ -1708,7 +1712,7 @@ class IndicatorData : public IndicatorBase { */ virtual IndicatorData* DataSourceRequestReturnDefault(int _applied_price) { DebugBreak(); - return NULL; + return nullptr; } /** diff --git a/Indicators/Account/Indi_AccountStats.mqh b/Indicators/Account/Indi_AccountStats.mqh index da45fef7d..975064d96 100644 --- a/Indicators/Account/Indi_AccountStats.mqh +++ b/Indicators/Account/Indi_AccountStats.mqh @@ -145,7 +145,6 @@ class Indi_AccountStats : public Indicator { * Returns value storage of given kind. */ IValueStorage *GetSpecificValueStorage(ENUM_INDI_VS_TYPE _type) override { - // Returning Price indicator which provides applied price in the only buffer #0. switch (_type) { case INDI_VS_TYPE_ACCOUNT_STATS_DATE_TIME: return buffer_date_time.Ptr(); @@ -200,8 +199,27 @@ class Indi_AccountStats : public Indicator { return; } - // Adding new account stats entry. - - Print("New candle: ", entry.ToString()); + // New candle means that account stats for current index 0 will be that we + // will now extract and store in the buffers. + + // Extracting current account stats. + datetime stats_date_time = iparams.GetAccount() PTR_DEREF GetDateTime(); + float stats_balance = iparams.GetAccount() PTR_DEREF GetBalance(); + float stats_credit = iparams.GetAccount() PTR_DEREF GetCredit(); + float stats_profit = iparams.GetAccount() PTR_DEREF GetProfit(); + float stats_equity = iparams.GetAccount() PTR_DEREF GetEquity(); + float stats_margin_used = iparams.GetAccount() PTR_DEREF GetMarginUsed(); + float stats_margin_free = iparams.GetAccount() PTR_DEREF GetMarginFree(); + float stats_margin_avail = iparams.GetAccount() PTR_DEREF GetMarginAvail(); + + // Appending account stats into buffers. + buffer_date_time REF_DEREF Append(stats_date_time); + buffer_balance REF_DEREF Append(stats_balance); + buffer_credit REF_DEREF Append(stats_credit); + buffer_profit REF_DEREF Append(stats_profit); + buffer_equity REF_DEREF Append(stats_equity); + buffer_margin_used REF_DEREF Append(stats_margin_used); + buffer_margin_free REF_DEREF Append(stats_margin_free); + buffer_margin_avail REF_DEREF Append(stats_margin_avail); } }; diff --git a/Indicators/Account/tests/Indi_AccountStats.test.mq5 b/Indicators/Account/tests/Indi_AccountStats.test.mq5 index c81e0da79..9f72f4356 100644 --- a/Indicators/Account/tests/Indi_AccountStats.test.mq5 +++ b/Indicators/Account/tests/Indi_AccountStats.test.mq5 @@ -48,7 +48,7 @@ int OnInit() { void OnTick() { Platform::Tick(); - if (Platform::IsNewHour()) { + if (Platform::IsNewMinute()) { IndicatorDataEntry _entry = indi_account_mt REF_DEREF GetEntry(); bool _is_ready = indi_account_mt REF_DEREF Get(STRUCT_ENUM(IndicatorState, INDICATOR_STATE_PROP_IS_READY)); bool _is_valid = _entry.IsValid(); diff --git a/Storage/ValueStorage.native.h b/Storage/ValueStorage.native.h index 0c6180881..7fb8219fc 100644 --- a/Storage/ValueStorage.native.h +++ b/Storage/ValueStorage.native.h @@ -83,6 +83,15 @@ class NativeValueStorage : public ValueStorage { Array::ArrayStore(_values, _index, _value, 4096); } + /** + * Inserts new value at the end of the buffer. If buffer works as As-Series, + * then new value will act as the one at index 0. + */ + void Append(C _value) override { + Resize(Size() + 1, 4096); + Store(Size() - 1, _value); + } + /** * Returns number of values available to fetch (size of the values buffer). */ From d6a3f579c325db308c287a907838c38c1bde4475 Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Fri, 12 Apr 2024 19:06:42 +0200 Subject: [PATCH 4/8] Refs #738. WIP. OpenCL wrapper class and OpenCL integration for Matrix class (not yet complete). --- Matrix.mqh | 101 ++++++++- OpenCL.h | 523 +++++++++++++++++++++++++++++++++++++++++++ tests/OpenCLTest.mq5 | 55 +++++ 3 files changed, 677 insertions(+), 2 deletions(-) create mode 100644 OpenCL.h create mode 100644 tests/OpenCLTest.mq5 diff --git a/Matrix.mqh b/Matrix.mqh index 59f3eef39..a2f38ce7a 100644 --- a/Matrix.mqh +++ b/Matrix.mqh @@ -24,6 +24,8 @@ #ifndef MATRIX_MQH #define MATRIX_MQH +#define MATRIX_USE_OPENCL + #ifdef USE_MQL_MATH_STAT #ifdef __MQL5__ #include @@ -31,6 +33,7 @@ #endif #include "Math.h" +#include "OpenCL.h" #define MATRIX_DIMENSIONS 6 #define MATRIX_VALUES_ARRAY_INCREMENT 500 @@ -719,6 +722,40 @@ class MatrixDimension { } }; +enum ENUM_MATRIX_FLAGS { MATRIX_FLAGS_NONE, MATRIX_FLAGS_USE_OPENCL }; + +/** + * Buffer used for CL operations. + */ +template +struct MatrixOpenCLBuffer : Dynamic { + // Flattened matrix data. + ARRAY(X, data); + + // Current version of the data. + long version; + + // CL buffer. + Ref buffer; + + public: + /** + * Constructor. + */ + MatrixBuffer() { version = 0; } + + /** + * Prepares buffer to be used by CL. Copies flattened data from the given matrix into buffer. + */ + void FillData(const Matrix& src) { + src.GetRawArray(data); + + if (!buffer.IsSet()) { + buffer = OpenCL::Alloc(ArraySize(data), true); + } + } +}; + /** * Matrix class. */ @@ -728,6 +765,11 @@ class Matrix { // First/root dimension. MatrixDimension* ptr_first_dimension; + // Map of data size -> CL buffer to be used e.g., by CL-based MatMul method. + DictStruct> _cl_buffers_in_0; + DictStruct> _cl_buffers_in_1; + DictStruct> _cl_buffers_out; + // Array with declaration of items per matrix's dimension. int dimensions[MATRIX_DIMENSIONS]; @@ -737,10 +779,18 @@ class Matrix { // Number of matrix dimensions. int num_dimensions; + // Flags. + int flags; + /** * Constructor. */ - Matrix(string _data) { FromString(_data); } + Matrix(string _data) { + FromString(_data); +#ifdef MATRIX_USE_OPENCL + InitializeOpenCL(); +#endif + } /** * Constructor. @@ -748,12 +798,20 @@ class Matrix { Matrix(const int num_1d = 0, const int num_2d = 0, const int num_3d = 0, const int num_4d = 0, const int num_5d = 0) { ptr_first_dimension = NULL; SetShape(num_1d, num_2d, num_3d, num_4d, num_5d); +#ifdef MATRIX_USE_OPENCL + InitializeOpenCL(); +#endif } /** * Constructor. */ - Matrix(MatrixDimension* _dimension) : ptr_first_dimension(NULL) { Initialize(_dimension); } + Matrix(MatrixDimension* _dimension) : ptr_first_dimension(NULL) { + Initialize(_dimension); +#ifdef MATRIX_USE_OPENCL + InitializeOpenCL(); +#endif + } /** * Copy constructor. @@ -764,6 +822,9 @@ class Matrix { } Initialize(_right.ptr_first_dimension.Clone()); +#ifdef MATRIX_USE_OPENCL + InitializeOpenCL(); +#endif } /** @@ -773,7 +834,23 @@ class Matrix { private: Matrix(const Matrix* _right) {} +#ifdef MATRIX_USE_OPENCL + + /** + * + */ + void InitializeOpenCL() {} + +#endif + public: + /** + * Returns/allocs and returns buffer of the given size to be used in CL operations as first input parameter. + * + * Buffer will be read only. + */ + MatrixOpenCLBuffer* GetCLBufferInArg0(int size) { if (_) } + /** * Matrix initializer. */ @@ -1247,6 +1324,11 @@ class Matrix { } static void MatMul(Matrix& source, Matrix& target, Matrix& output) { +#ifdef MATRIX_USE_OPENCL + MatMulCL(source, target, output); + return; +#endif + if (source.GetSize() != target.GetRange(1)) { Alert("Inconsistent size of matrices!"); } @@ -1264,6 +1346,21 @@ class Matrix { } } + /** + * Performs matrix multiplication via OpenCL. Note that MATRIX_USE_OPENCL must be defined in order matrix to use this + * method. + */ + static void MatMulCL(Matrix& source, Matrix& target, Matrix& output) { + if (source.GetSize() != target.GetRange(1)) { + Alert("Inconsistent size of matrices!"); + } + + int num_outputs = target.GetRange(0); + int num_inputs = target.GetRange(1); + + output.SetShape(num_outputs); + } + /** * Performs matrix multiplication. */ diff --git a/OpenCL.h b/OpenCL.h new file mode 100644 index 000000000..156c4ff65 --- /dev/null +++ b/OpenCL.h @@ -0,0 +1,523 @@ +#ifndef __MQL__ +#pragma once +#endif + +#include "DictStruct.mqh" + +// Defines. +#define OPENCL_PROGRAM_MAX_ARGS 8 + +// Forward declarations; +class OpenCLProgram; + +/** + * Memory buffer. + */ +class OpenCLBuffer : public Dynamic { + // Handle to memory buffer. + int buffer_handle; + + // Allocated buffer size. + int buffer_size; + + // Whether buffer is global or local to the kernel otherwise. + bool is_global; + + // Buffer version. Should be incremented after each change. + long version; + + public: + /** + * Constructor. + */ + OpenCLBuffer(int _size, bool _is_global, unsigned int _flags = CL_MEM_READ_WRITE); + + /** + * Writes/uploads data into buffer if needed. + */ + void Write(const ARRAY_REF(double, _arr), int _arr_version = -1) { + if (ArraySize(_arr) > buffer_size) { + Alert("Array passed is too large for the allocated buffer. Tries to pass ", ArraySize(_arr), + " elements into buffer of size ", buffer_size, "."); + DebugBreak(); + return; + } + + // Do we need to reupload data into GPU? + if (_arr_version != -1 && _arr_version <= version) { + // Buffer has already up-to-date data. + return; + } + + CLBufferWrite(buffer_handle, _arr); + + version = (_arr_version != -1) ? _arr_version : (version + 1); + } + + /** + * Reads data from buffer. + */ + void Read(ARRAY_REF(double, _arr)) { + if (!ArrayIsDynamic(_arr) && ArraySize(_arr) < buffer_size) { + Alert("Array passed is too small to be the target. Buffer has size ", buffer_size, + " and you tried to read it into buffer of size ", ArraySize(_arr), "."); + DebugBreak(); + return; + } + CLBufferRead(buffer_handle, _arr); + } + + /** + * Whether buffer is global or local to the kernel otherwise. + */ + bool IsGlobal() { return is_global; } + + /** + * Returns data version. + */ + long GetVersion() { return version; } + + /** + * Returns handle to buffer. + */ + int GetHandle() { return buffer_handle; } + + /** + * Destructor. + */ + ~OpenCLBuffer() { + if (buffer_handle != INVALID_HANDLE) { + CLBufferFree(buffer_handle); + } + } +}; + +/** + * Single program (code) + kernel (function name) to be invoked. + */ +class OpenCLProgram : public Dynamic { + // Handle to program. + int program_handle; + + // Handle to kernel. + int kernel_handle; + + // Buffer handles previously passed as arguments. Used to check if buffer needs to be reuploaded. + int arg_handles[OPENCL_PROGRAM_MAX_ARGS]; + + // Version of argument data. Used to check if buffer needs to be reuploaded. + long arg_versions[OPENCL_PROGRAM_MAX_ARGS]; + + public: + /** + * Constructor. + */ + OpenCLProgram() : program_handle(INVALID_HANDLE), kernel_handle(INVALID_HANDLE) { + for (int i = 0; i < OPENCL_PROGRAM_MAX_ARGS; ++i) { + arg_handles[i] = INVALID_HANDLE; + arg_versions[i] = -1; + } + } + + /** + * Destructor. + */ + ~OpenCLProgram() { + if (kernel_handle != INVALID_HANDLE) { + CLKernelFree(kernel_handle); + } + + if (program_handle != INVALID_HANDLE) { + CLProgramFree(program_handle); + } + } + + /** + * Passes argument to the kernel. Will not set kernel argument if not needed. + * + * Note that buffer reuploading is to be done via OpenCLBuffer::Write() in + * which you can pass version of your data, so no reupload will take place if + * your version isn't greater that the one already set in the buffer. + */ + void SetArg(int _index, OpenCLBuffer* _buffer) { + if (_buffer PTR_DEREF GetHandle() == arg_handles[_index] && + _buffer PTR_DEREF GetVersion() >= arg_versions[_index]) { + // Already uploaded recent version. + return; + } + + if (_buffer PTR_DEREF IsGlobal()) { + CLSetKernelArgMem(kernel_handle, _index, _buffer PTR_DEREF GetHandle()); + } else { + CLSetKernelArgMemLocal(kernel_handle, _index, _buffer PTR_DEREF GetHandle()); + } + + // Buffer will occupy argument slot. + arg_handles[_index] = _buffer PTR_DEREF GetHandle(); + + // Storing buffer version in the argument slot. + arg_versions[_index] = _buffer PTR_DEREF GetVersion(); + } + + /** + * Executes a single kernel. + */ + bool Run() { + if (!CLExecute(kernel_handle)) { + Alert("OpenCL error occured when tried to run kernel: ", GetLastError(), "!"); + return false; + } + + return true; + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a) { + SetArg(0, a); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b) { + SetArg(0, a); + SetArg(1, b); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b, C c) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b, C c, D d) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b, C c, D d, E e) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b, C c, D d, E e, F f) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + SetArg(5, f); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b, C c, D d, E e, F f, G g) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + SetArg(5, f); + SetArg(6, g); + return Run(); + } + + /** + * Executes a single kernel. Allows passing arugments to kernel. + */ + template + bool Run(A a, B b, C c, D d, E e, F f, G g, H h) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + SetArg(5, f); + SetArg(6, g); + SetArg(7, h); + return Run(); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. + */ + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size) { + if (!CLExecute(kernel_handle)) { + Alert("OpenCL error occured when tried to run kernel: ", GetLastError(), "!"); + return false; + } + + return true; + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a) { + SetArg(0, a); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b) { + SetArg(0, a); + SetArg(1, b); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, + C c) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, + D d) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, + D d, E e) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, + D d, E e, F f) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + SetArg(5, f); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, + D d, E e, F f, G g) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + SetArg(5, f); + SetArg(6, g); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. + */ + template + bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, + D d, E e, F f, G g, H h) { + SetArg(0, a); + SetArg(1, b); + SetArg(2, c); + SetArg(3, d); + SetArg(4, e); + SetArg(5, f); + SetArg(6, g); + SetArg(7, h); + return RunMany(_dimension, _work_offset, _work_size); + } + + /** + * Returns handle to OpenCL program. + */ + int GetProgramHandle() { return program_handle; } + + /** + * Sets handle to OpenCL program. + */ + void SetProgramHandle(int _handle) { + if (program_handle != -1) { + Alert("Cannot change program handle!"); + DebugBreak(); + return; + } + program_handle = _handle; + } + + /** + * Returns handle to OpenCL kernel. + */ + int GetKernelHandle() { return kernel_handle; } + + /** + * Sets handle to OpenCL kernel. + */ + void SetKernelHandle(int _handle) { + if (kernel_handle != -1) { + Alert("Cannot change kernel handle!"); + DebugBreak(); + return; + } + kernel_handle = _handle; + } +}; + +/** + * Wrapper for OpenCL. + */ +class OpenCL { + // OpenCL handles. + static int context_handle; + + // OpenCL global memory handles. + static int cl_mem_0, cl_mem_1, cl_mem_2; + + DictStruct> programs; + + public: + /** + * Initializes CL contexts. Called automatically by OpenCLLifetimeManager. + */ + static void Initialize() { + context_handle = CLContextCreate(); + if (context_handle == INVALID_HANDLE) { + Alert("Could not create OpenCL context. Error code: ", GetLastError(), "."); + DebugBreak(); + return; + } + } + + /** + * Frees CL contexts. Called automatically by OpenCLLifetimeManager. + */ + static void Deinitialize() { CLContextFree(context_handle); } + + /** + * Allocates memory to be later passed to OpenCLProgram. + */ + static OpenCLBuffer* Alloc(int _size, bool _is_global = true) { return new OpenCLBuffer(_size, _is_global); } + + /** + * Compiles given program and returns its id or -1 in case of error. + */ + static OpenCLProgram* Compile(string _source, string _fn_name) { + OpenCLProgram* _program = new OpenCLProgram(); + + // Log of CLProgramCreate(). + string _compilation_log; + + _program PTR_DEREF SetProgramHandle(CLProgramCreate(context_handle, _source, _compilation_log)); + + if (_program PTR_DEREF GetProgramHandle() == INVALID_HANDLE) { + Alert("Could not create OpenCL program. Error code: ", GetLastError(), ". Compilation log: ", _compilation_log, + "."); + DebugBreak(); + return nullptr; + } + + _program PTR_DEREF SetKernelHandle(CLKernelCreate(_program PTR_DEREF GetProgramHandle(), _fn_name)); + + if (_program PTR_DEREF GetKernelHandle() == INVALID_HANDLE) { + Alert("Could not create OpenCL kernel. Error code: ", GetLastError(), "."); + DebugBreak(); + return nullptr; + } + + return _program; + } + + /** + * Returns handle to OpenCL context. + */ + static int GetContextHandle() { return context_handle; } +}; + +static int OpenCL::context_handle; + +/** + * Manages initialization and deinitialization of static variables for OpenCL class. + */ +class OpenCLLifetimeManager { + public: + OpenCLLifetimeManager() { OpenCL::Initialize(); } + + ~OpenCLLifetimeManager() { OpenCL::Deinitialize(); } +}; + +OpenCLLifetimeManager __opencl_lifetime_manager; + +/** + * OpenCLBuffer constructor. + */ +OpenCLBuffer::OpenCLBuffer(int _size, bool _is_global, unsigned int _flags) { + buffer_handle = CLBufferCreate(OpenCL::GetContextHandle(), _size * sizeof(double), _flags); + if (buffer_handle == INVALID_HANDLE) { + Alert("Could not create OpenCL buffer. Error code: ", GetLastError(), "."); + DebugBreak(); + } + buffer_size = _size; + is_global = _is_global; + version = 0; +} diff --git a/tests/OpenCLTest.mq5 b/tests/OpenCLTest.mq5 new file mode 100644 index 000000000..8f13e2a30 --- /dev/null +++ b/tests/OpenCLTest.mq5 @@ -0,0 +1,55 @@ +//+------------------------------------------------------------------+ +//| EA31337 framework | +//| Copyright 2016-2021, EA31337 Ltd | +//| https://github.com/EA31337 | +//+------------------------------------------------------------------+ + +/* + * This file is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * (at your option) any later version. + + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + + * You should have received a copy of the GNU General Public License + * along with this program. If not, see . + */ + +/** + * @file + * Test functionality of Ticker class. + */ + +// Includes. +#include "../Matrix.mqh" +#include "../OpenCL.h" + +/** + * Implements initialization function. + */ +int OnInit() { + Ref program = + OpenCL::Compile("#pragma OPENCL EXTENSION cl_khr_fp64 : enable" NL + "__kernel void test(__global double *data) {" NL " data[0] = 5;" NL "}" NL, + + "test"); + + double result[] = {0}; + + Ref buffer = OpenCL::Alloc(1 /* 1 double */); + + if (!program REF_DEREF Run(buffer.Ptr())) { + Alert("Error running program!"); + } + + buffer REF_DEREF Read(result); + + Print("Output: ", result[0]); + + ExpertRemove(); + return (INIT_SUCCEEDED); +} From 2f2fce7b44abe6001cef897ea0aa09fe5691c663 Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Fri, 19 Apr 2024 21:01:01 +0200 Subject: [PATCH 5/8] Refs #738. WIP. OpenCL integration for Matrix class needs rework as previous implementation we based on was buggy. --- Array.mqh | 2 +- Indicator.struct.h | 2 +- Matrix.mqh | 274 +++++++++++++++++++++++++++++++++++++++---- OpenCL.h | 101 +++++++++------- tests/OpenCLTest.mq5 | 12 ++ 5 files changed, 328 insertions(+), 63 deletions(-) diff --git a/Array.mqh b/Array.mqh index 65d8c529a..7ee00dbd2 100644 --- a/Array.mqh +++ b/Array.mqh @@ -772,7 +772,7 @@ static int GetLowestArrDoubleValue(double& arr[][], int key) { template void ArrayPush(ARRAY_REF(X, array), X value) { - ArrayResize(ArraySize(array) + 1); + ArrayResize(array, ArraySize(array) + 1); array[ArraySize(array) - 1] = value; } template diff --git a/Indicator.struct.h b/Indicator.struct.h index 1819695f3..78a128274 100644 --- a/Indicator.struct.h +++ b/Indicator.struct.h @@ -120,7 +120,7 @@ struct IndicatorParams { } void SetMaxParams(int _value) { max_params = _value; - ArrayResize(input_params, max_params); + ArrayResize(input_params, (int)max_params); } void SetName(string _name) { name = _name; }; void SetShift(int _shift) { shift = _shift; } diff --git a/Matrix.mqh b/Matrix.mqh index a2f38ce7a..00cbea7af 100644 --- a/Matrix.mqh +++ b/Matrix.mqh @@ -160,6 +160,7 @@ struct MatrixDimensionAccessor { if (ptr_dimension.type != MATRIX_DIMENSION_TYPE_VALUES) { \ Print("Error: Trying to use matrix", ptr_matrix.Repr(), \ "'s value operator " #OP " in a dimension which doesn't contain values!"); \ + DebugBreak(); \ return; \ } \ \ @@ -177,6 +178,7 @@ struct MatrixDimensionAccessor { void operator=(X _value) { if (ptr_dimension.type != MATRIX_DIMENSION_TYPE_VALUES) { Print("Error: Trying to set matrix", ptr_matrix.Repr(), "'s value in a dimension which doesn't contain values!"); + DebugBreak(); return; } @@ -189,6 +191,7 @@ struct MatrixDimensionAccessor { X Val() { if (ptr_dimension.type != MATRIX_DIMENSION_TYPE_VALUES) { Print("Error: Trying to get value from matrix", ptr_matrix.Repr(), "'s dimension which doesn't contain values!"); + DebugBreak(); return (X)EMPTY_VALUE; } @@ -202,6 +205,7 @@ struct MatrixDimensionAccessor { X ValOrZero() { if (ptr_dimension.type != MATRIX_DIMENSION_TYPE_VALUES) { Print("Error: Trying to get value from matrix", ptr_matrix.Repr(), "'s dimension which doesn't contain values!"); + DebugBreak(); return (X)EMPTY_VALUE; } @@ -728,7 +732,7 @@ enum ENUM_MATRIX_FLAGS { MATRIX_FLAGS_NONE, MATRIX_FLAGS_USE_OPENCL }; * Buffer used for CL operations. */ template -struct MatrixOpenCLBuffer : Dynamic { +class MatrixOpenCLBuffer : public Dynamic { // Flattened matrix data. ARRAY(X, data); @@ -742,18 +746,24 @@ struct MatrixOpenCLBuffer : Dynamic { /** * Constructor. */ - MatrixBuffer() { version = 0; } + MatrixOpenCLBuffer(int _size) { + version = 0; + buffer = OpenCL::Alloc(_size); + ArrayResize(data, _size); + } /** * Prepares buffer to be used by CL. Copies flattened data from the given matrix into buffer. */ void FillData(const Matrix& src) { src.GetRawArray(data); - - if (!buffer.IsSet()) { - buffer = OpenCL::Alloc(ArraySize(data), true); - } + buffer REF_DEREF Write(data, version); } + + /** + * Returns pointer to the CL buffer. + */ + OpenCLBuffer* GetBuffer() { return buffer.Ptr(); } }; /** @@ -766,9 +776,9 @@ class Matrix { MatrixDimension* ptr_first_dimension; // Map of data size -> CL buffer to be used e.g., by CL-based MatMul method. - DictStruct> _cl_buffers_in_0; - DictStruct> _cl_buffers_in_1; - DictStruct> _cl_buffers_out; + static DictStruct>> cl_buffers_in_0; + static DictStruct>> cl_buffers_in_1; + static DictStruct>> cl_buffers_out; // Array with declaration of items per matrix's dimension. int dimensions[MATRIX_DIMENSIONS]; @@ -782,6 +792,16 @@ class Matrix { // Flags. int flags; + // Unique id of the matrix. + int uuid; + + // Static counter, so each matrix will have its own uuid. + static int uuid_counter; + + // OpenCL program. + static Ref cl_program_matmul; + static Ref cl_program_matmul_single; + /** * Constructor. */ @@ -790,6 +810,7 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif + uuid = uuid_counter++; } /** @@ -801,6 +822,7 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif + uuid = uuid_counter++; } /** @@ -811,6 +833,7 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif + uuid = uuid_counter++; } /** @@ -825,6 +848,9 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif + + // We mark new matrix as unique one, even though we clone another matrix. + uuid = uuid_counter++; } /** @@ -839,17 +865,83 @@ class Matrix { /** * */ - void InitializeOpenCL() {} + void InitializeOpenCL() { + if (cl_program_matmul.IsSet()) { + // Already initialized. + return; + } + + cl_program_matmul = OpenCL::Compile( + "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" NL "__kernel void matmul(" NL " const int Mdim," NL + " const int Ndim," NL " const int Pdim," NL " __global double* A," NL " __global double* B," NL + " __global double* C," NL " __local double* Bwrk" NL ")" NL "{" NL " int k, j;" NL + " int i = get_global_id(0);" NL " int iloc = get_local_id(0);" NL + " int nloc = get_local_size(0);" NL " double Awrk[1000];" NL " double tmp;" NL "" NL + " for (k = 0; k < Pdim; k++)" NL " Awrk[k] = A[i * Ndim + k];" NL "" NL + " for (j = 0; j < Mdim; j++)" NL " {" NL " for (k = iloc; k < Pdim; k = k + nloc)" NL + " Bwrk[k] = B[k * Pdim + j];" NL "" NL " barrier(CLK_LOCAL_MEM_FENCE);" NL "" NL + " tmp = 0.0f;" NL "" NL " for (k = 0; k < Pdim; k++)" NL + " tmp += Awrk[k] * Bwrk[k];" NL " C[i * Ndim + j] += tmp;" NL " }" NL "}" NL, + "matmul"); + + cl_program_matmul_single = OpenCL::Compile( + "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" NL "__kernel void matmul(" NL "const int Mdim," NL + "const int Ndim," NL "const int Pdim," NL "__global float *A," NL "__global float *B," NL "__global float *C" NL + ")" NL "{" NL "int i, j, k;" NL "for (i=0; i* GetCLBufferInArg0(int _size) { + Ref> _buffer; + + _buffer = cl_buffers_in_0.GetByKey(_size, _buffer); + + if (!_buffer.IsSet()) { + _buffer = new MatrixOpenCLBuffer(_size); + cl_buffers_in_0.Set(_size, _buffer); + } + + return _buffer.Ptr(); + } + + /** + * Returns/allocs and returns buffer of the given size to be used in CL operations as second input parameter. + */ + static MatrixOpenCLBuffer* GetCLBufferInArg1(int _size) { + Ref> _buffer; + + _buffer = cl_buffers_in_1.GetByKey(_size, _buffer); + + if (!_buffer.IsSet()) { + _buffer = new MatrixOpenCLBuffer(_size); + cl_buffers_in_1.Set(_size, _buffer); + } + + return _buffer.Ptr(); + } + + /** + * Returns/allocs and returns buffer of the given size to be used in CL operations as output parameter. + */ + static MatrixOpenCLBuffer* GetCLBufferOutArg(int _size) { + Ref> _buffer; + + _buffer = cl_buffers_out.GetByKey(_size, _buffer); + + if (!_buffer.IsSet()) { + _buffer = new MatrixOpenCLBuffer(_size); + cl_buffers_out.Set(_size, _buffer); + } + + return _buffer.Ptr(); + } /** * Matrix initializer. @@ -876,6 +968,7 @@ class Matrix { break; } else { Print("Internal error: unknown dimension type!"); + DebugBreak(); } } @@ -986,7 +1079,7 @@ class Matrix { /** * Returns total number of values the matrix contain of. */ - int GetSize() { return size; } + int GetSize() const { return size; } /** * Returns number of matrix dimensions. @@ -1325,7 +1418,8 @@ class Matrix { static void MatMul(Matrix& source, Matrix& target, Matrix& output) { #ifdef MATRIX_USE_OPENCL - MatMulCL(source, target, output); + // MatMulCL(source, target, output); + MatMulCLSingle(source, target, output); return; #endif @@ -1350,15 +1444,116 @@ class Matrix { * Performs matrix multiplication via OpenCL. Note that MATRIX_USE_OPENCL must be defined in order matrix to use this * method. */ - static void MatMulCL(Matrix& source, Matrix& target, Matrix& output) { - if (source.GetSize() != target.GetRange(1)) { + static void MatMulCL(Matrix& _source, Matrix& _target, Matrix& _output) { + if (_source.GetSize() != _target.GetRange(1)) { Alert("Inconsistent size of matrices!"); } - int num_outputs = target.GetRange(0); - int num_inputs = target.GetRange(1); + unsigned int _m_dim = _source.GetRange(0); + unsigned int _n_dim = _source.GetRange(1); + unsigned int _p_dim = _target.GetRange(1); - output.SetShape(num_outputs); + OpenCLBuffer* _in_1 = GetCLBufferInArg0(_source.GetSize()) PTR_DEREF GetBuffer(); + OpenCLBuffer* _in_2 = GetCLBufferInArg1(_target.GetSize()) PTR_DEREF GetBuffer(); + OpenCLBuffer* _out = GetCLBufferOutArg(_source.GetRange(0)) PTR_DEREF GetBuffer(); + + double _in_1_data[]; + double _in_2_data[]; + double _out_data[]; + + _source.GetRawArray(_in_1_data); + _target.GetRawArray(_in_2_data); + + cl_program_matmul REF_DEREF SetArg(0, _m_dim); + cl_program_matmul REF_DEREF SetArg(1, _n_dim); + cl_program_matmul REF_DEREF SetArg(2, _p_dim); + cl_program_matmul REF_DEREF SetArg(3, _in_1); + cl_program_matmul REF_DEREF SetArg(4, _in_2); + cl_program_matmul REF_DEREF SetArg(5, _out); + cl_program_matmul REF_DEREF SetArgLocalMem(6, _p_dim * sizeof(X)); + ARRAY(unsigned int, _global_work_offset); + ARRAY(unsigned int, _global_work_size); + ARRAY(unsigned int, _local_work_size); + ArrayPush(_global_work_offset, 0U); + ArrayPush(_global_work_size, _n_dim); + ArrayPush(_global_work_size, _m_dim); + ArrayPush(_local_work_size, _m_dim); + + if (!cl_program_matmul REF_DEREF RunMany(2, _global_work_offset, _global_work_size, _local_work_size)) { + Alert("Errpr: Could not run Matrix::MatMulCL() over OpenCL!"); + DebugBreak(); + } + + _out PTR_DEREF Read(_out_data); + + // _output.SetShape(num_outputs); + } + + /** + * Performs matrix multiplication via OpenCL. Note that MATRIX_USE_OPENCL must be defined in order matrix to use this + * method. + */ + static void MatMulCLSingle(Matrix& _source, Matrix& _target, Matrix& _output) { + if (_source.GetRange(1) != _target.GetRange(0)) { + Alert("Inconsistent size of matrices!"); + } + + unsigned int _m_dim = _source.GetRange(0); + unsigned int _n_dim = _source.GetRange(1); + unsigned int _p_dim = _target.GetRange(1); + + OpenCLBuffer* _in_1 = GetCLBufferInArg0(_n_dim * _p_dim) PTR_DEREF GetBuffer(); + OpenCLBuffer* _in_2 = GetCLBufferInArg1(_p_dim * _m_dim) PTR_DEREF GetBuffer(); + OpenCLBuffer* _out = GetCLBufferOutArg(_m_dim * _p_dim) PTR_DEREF GetBuffer(); + + double _in_1_data[]; + double _in_2_data[]; + double _out_data[]; + + ArrayResize(_out_data, _out PTR_DEREF GetSizeItems()); + + _source.GetRawArray(_in_1_data); + _target.GetRawArray(_in_2_data); + + MatMulCL_CPU(_m_dim, _n_dim, _p_dim, _in_1_data, _in_2_data, _out_data); + + cl_program_matmul REF_DEREF SetArg(0, _n_dim); + cl_program_matmul REF_DEREF SetArg(1, _m_dim); + cl_program_matmul REF_DEREF SetArg(2, _p_dim); + cl_program_matmul REF_DEREF SetArg(3, _in_1); + cl_program_matmul REF_DEREF SetArg(4, _in_2); + cl_program_matmul REF_DEREF SetArg(5, _out); + + if (!cl_program_matmul_single REF_DEREF Run()) { + Alert("Errpr: Could not run Matrix::MatMulCL() over OpenCL!"); + DebugBreak(); + } + + _out PTR_DEREF Read(_out_data); + + // _output.SetShape(num_outputs); + } + + static void MatMulCL_CPU(int Mdim, int Ndim, int Pdim, double& A[], double& B[], double& C[]) { + /* + for (int i = 0; i < Mdim; ++i) { + for (int j = 0; j < Pdim; ++j) { + C[i * Pdim + j] = 0.0; + for (int k = 0; k < Ndim; ++k) { + C[i * Pdim + j] += A[i * Ndim + k] * B[k * Pdim + j]; + } + } + } + */ + + int i, j, k; + for (i = 0; i < Ndim; i++) { + for (j = 0; j < Mdim; j++) { + for (k = 0; k < Pdim; k++) { // C(i,j) = sum(over k) A(i,k) * B(k,j) + C[i * Pdim + j] += A[i * Ndim + k] * B[k * Pdim + j]; + } + } + } } /** @@ -1370,6 +1565,15 @@ class Matrix { return output; } + /** + * Performs matrix multiplication. + */ + Matrix* MatMul(Matrix* target) { + Matrix* output = new Matrix(); + MatMul(this, target, output); + return output; + } + /** * Performs matrix multiplication. */ @@ -1475,7 +1679,7 @@ class Matrix { /** * Fills array with all values from the matrix. */ - void GetRawArray(X& array[]) { + void GetRawArray(X& array[]) const { ArrayResize(array, GetSize()); int offset = 0; ptr_first_dimension.FillArray(array, offset); @@ -2359,4 +2563,32 @@ class Matrix { } }; +#ifdef __MQL__ +template +static int Matrix::uuid_counter = 0; +template +static Ref Matrix::cl_program_matmul; +template +static Ref Matrix::cl_program_matmul_single; +template +static DictStruct>> Matrix::cl_buffers_in_0; +template +static DictStruct>> Matrix::cl_buffers_in_1; +template +static DictStruct>> Matrix::cl_buffers_out; +#else +template +static int Matrix::uuid_counter = 0; +template +static Ref Matrix::cl_program_matmul; +template +static Ref Matrix::cl_program_matmul_single; +template +static DictStruct>> Matrix::cl_buffers_in_0; +template +static DictStruct>> Matrix::cl_buffers_in_1; +template +static DictStruct>> Matrix::cl_buffers_out; +#endif + #endif diff --git a/OpenCL.h b/OpenCL.h index 156c4ff65..ffeb1314f 100644 --- a/OpenCL.h +++ b/OpenCL.h @@ -20,9 +20,6 @@ class OpenCLBuffer : public Dynamic { // Allocated buffer size. int buffer_size; - // Whether buffer is global or local to the kernel otherwise. - bool is_global; - // Buffer version. Should be incremented after each change. long version; @@ -30,12 +27,12 @@ class OpenCLBuffer : public Dynamic { /** * Constructor. */ - OpenCLBuffer(int _size, bool _is_global, unsigned int _flags = CL_MEM_READ_WRITE); + OpenCLBuffer(int _size, unsigned int _flags = CL_MEM_READ_WRITE); /** * Writes/uploads data into buffer if needed. */ - void Write(const ARRAY_REF(double, _arr), int _arr_version = -1) { + void Write(const ARRAY_REF(double, _arr), long _arr_version = -1) { if (ArraySize(_arr) > buffer_size) { Alert("Array passed is too large for the allocated buffer. Tries to pass ", ArraySize(_arr), " elements into buffer of size ", buffer_size, "."); @@ -68,9 +65,14 @@ class OpenCLBuffer : public Dynamic { } /** - * Whether buffer is global or local to the kernel otherwise. + * Returns buffer size in bytes. */ - bool IsGlobal() { return is_global; } + int GetSizeBytes() { return buffer_size * sizeof(double); } + + /** + * Returns buffer size in items. + */ + int GetSizeItems() { return buffer_size; } /** * Returns data version. @@ -132,6 +134,20 @@ class OpenCLProgram : public Dynamic { } } + /** + * Passes local memory size argument to the kernel. + */ + void SetArgLocalMem(int _index, unsigned long _mem_size) { CLSetKernelArgMemLocal(kernel_handle, _index, _mem_size); } + + /** + * Passes argument to the kernel. Will not set kernel argument if not needed. + * + * Note that buffer reuploading is to be done via OpenCLBuffer::Write() in + * which you can pass version of your data, so no reupload will take place if + * your version isn't greater that the one already set in the buffer. + */ + void SetArg(int _index, double value) { CLSetKernelArg(kernel_handle, _index, value); } + /** * Passes argument to the kernel. Will not set kernel argument if not needed. * @@ -146,11 +162,7 @@ class OpenCLProgram : public Dynamic { return; } - if (_buffer PTR_DEREF IsGlobal()) { - CLSetKernelArgMem(kernel_handle, _index, _buffer PTR_DEREF GetHandle()); - } else { - CLSetKernelArgMemLocal(kernel_handle, _index, _buffer PTR_DEREF GetHandle()); - } + CLSetKernelArgMem(kernel_handle, _index, _buffer PTR_DEREF GetHandle()); // Buffer will occupy argument slot. arg_handles[_index] = _buffer PTR_DEREF GetHandle(); @@ -274,8 +286,9 @@ class OpenCLProgram : public Dynamic { /** * Executes multiple kernels where work is subdivided among kernels. */ - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size) { - if (!CLExecute(kernel_handle)) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size)) { + if (!CLExecute(kernel_handle, _dimension, _global_work_offset, _global_work_size, _local_work_size)) { Alert("OpenCL error occured when tried to run kernel: ", GetLastError(), "!"); return false; } @@ -287,81 +300,89 @@ class OpenCLProgram : public Dynamic { * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a) { SetArg(0, a); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b) { SetArg(0, a); SetArg(1, b); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, - C c) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b, C c) { SetArg(0, a); SetArg(1, b); SetArg(2, c); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, - D d) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b, C c, D d) { SetArg(0, a); SetArg(1, b); SetArg(2, c); SetArg(3, d); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, - D d, E e) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b, C c, D d, E e) { SetArg(0, a); SetArg(1, b); SetArg(2, c); SetArg(3, d); SetArg(4, e); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, - D d, E e, F f) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b, C c, D d, E e, F f) { SetArg(0, a); SetArg(1, b); SetArg(2, c); SetArg(3, d); SetArg(4, e); SetArg(5, f); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, - D d, E e, F f, G g) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b, C c, D d, E e, F f, G g) { SetArg(0, a); SetArg(1, b); SetArg(2, c); @@ -369,15 +390,16 @@ class OpenCLProgram : public Dynamic { SetArg(4, e); SetArg(5, f); SetArg(6, g); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. */ template - bool RunMany(unsigned int _dimension, const unsigned int& _work_offset, const unsigned int& _work_size, A a, B b, C c, - D d, E e, F f, G g, H h) { + bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), + const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, + B b, C c, D d, E e, F f, G g, H h) { SetArg(0, a); SetArg(1, b); SetArg(2, c); @@ -386,7 +408,7 @@ class OpenCLProgram : public Dynamic { SetArg(5, f); SetArg(6, g); SetArg(7, h); - return RunMany(_dimension, _work_offset, _work_size); + return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); } /** @@ -431,7 +453,7 @@ class OpenCL { // OpenCL handles. static int context_handle; - // OpenCL global memory handles. + // OpenCL memory handles. static int cl_mem_0, cl_mem_1, cl_mem_2; DictStruct> programs; @@ -457,7 +479,7 @@ class OpenCL { /** * Allocates memory to be later passed to OpenCLProgram. */ - static OpenCLBuffer* Alloc(int _size, bool _is_global = true) { return new OpenCLBuffer(_size, _is_global); } + static OpenCLBuffer* Alloc(int _size) { return new OpenCLBuffer(_size); } /** * Compiles given program and returns its id or -1 in case of error. @@ -511,13 +533,12 @@ OpenCLLifetimeManager __opencl_lifetime_manager; /** * OpenCLBuffer constructor. */ -OpenCLBuffer::OpenCLBuffer(int _size, bool _is_global, unsigned int _flags) { +OpenCLBuffer::OpenCLBuffer(int _size, unsigned int _flags) { buffer_handle = CLBufferCreate(OpenCL::GetContextHandle(), _size * sizeof(double), _flags); if (buffer_handle == INVALID_HANDLE) { Alert("Could not create OpenCL buffer. Error code: ", GetLastError(), "."); DebugBreak(); } buffer_size = _size; - is_global = _is_global; version = 0; } diff --git a/tests/OpenCLTest.mq5 b/tests/OpenCLTest.mq5 index 8f13e2a30..800eb08ed 100644 --- a/tests/OpenCLTest.mq5 +++ b/tests/OpenCLTest.mq5 @@ -50,6 +50,18 @@ int OnInit() { Print("Output: ", result[0]); + Matrix *in1, *in2, *out; + in1 = Matrix::CreateFromString("[[1,2,3], [4,5,6]]"); // 2 x 3 + Print("in1 shape: ", in1 PTR_DEREF GetRange(0), " x ", in1 PTR_DEREF GetRange(1)); + in2 = Matrix::CreateFromString("[[7,8,9,10], [11,12,13,14], [15,16,17,18]]"); // 3 x 4 + Print("in2 shape: ", in2 PTR_DEREF GetRange(0), " x ", in2 PTR_DEREF GetRange(1)); + out = in1 PTR_DEREF MatMul(in2); + Print("out shape: ", out PTR_DEREF GetRange(0), " x ", out PTR_DEREF GetRange(1)); + + delete in1; + delete in2; + delete out; + ExpertRemove(); return (INIT_SUCCEEDED); } From b612dac1c4b9c57c3a57132ad035b1938d8e103b Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Wed, 24 Apr 2024 19:20:02 +0200 Subject: [PATCH 6/8] Refs #738. WIP. Partially working Matrix multiplication via OpenCL. Now we need some improvements and Matrix::Deflatten(). --- Matrix.matmul.cl | 115 +++++++++++++++++++++++++++++++++++++++++ Matrix.matmul.naive.cl | 23 +++++++++ Matrix.matmul.test.cl | 19 +++++++ Matrix.mqh | 97 +++++++++++++++++----------------- OpenCL.h | 4 +- tests/OpenCLTest.mq5 | 2 +- 6 files changed, 210 insertions(+), 50 deletions(-) create mode 100644 Matrix.matmul.cl create mode 100644 Matrix.matmul.naive.cl create mode 100644 Matrix.matmul.test.cl diff --git a/Matrix.matmul.cl b/Matrix.matmul.cl new file mode 100644 index 000000000..df718efcc --- /dev/null +++ b/Matrix.matmul.cl @@ -0,0 +1,115 @@ +#define WIDTH 1 +#define TSM 128 // The tile-size in dimension M +#define TSN 128 // The tile-size in dimension N +#define TSK 16 // The tile-size in dimension K +#define WPTM 8 // The work-per-thread in dimension M +#define WPTN 8 // The work-per-thread in dimension N +#define RTSM (TSM/WPTM) // The reduced tile-size in dimension M +#define RTSN (TSN/WPTN) // The reduced tile-size in dimension N +#define LPTA ((TSK*TSM)/(RTSM*RTSN)) // Loads-per-thread for A +#define LPTB ((TSK*TSN)/(RTSM*RTSN)) // Loads-per-thread for B +__kernel void matmul(const int M, const int N, const int K, + const __global double* A, + const __global double* B, + __global float* C) { + + // Thread identifiers + const int tidm = get_local_id(0); // Local row ID (max: TSM/WPTM) + const int tidn = get_local_id(1); // Local col ID (max: TSN/WPTN) + const int offsetM = TSM*get_group_id(0); // Work-group offset + const int offsetN = TSN*get_group_id(1); // Work-group offset + + // Local memory to fit a tile of A and B + __local float Asub[TSK][TSM]; + __local float Bsub[TSK][TSN]; + + // Allocate register space + float Areg; + float Breg[WPTN]; + float acc[WPTM][WPTN]; + + // Initialise the accumulation registers + for (int wm=0; wm(_size); + _buffer = new MatrixOpenCLBuffer(_size, CL_MEM_READ_ONLY); cl_buffers_in_0.Set(_size, _buffer); } @@ -920,7 +913,7 @@ class Matrix { _buffer = cl_buffers_in_1.GetByKey(_size, _buffer); if (!_buffer.IsSet()) { - _buffer = new MatrixOpenCLBuffer(_size); + _buffer = new MatrixOpenCLBuffer(_size, CL_MEM_READ_ONLY); cl_buffers_in_1.Set(_size, _buffer); } @@ -936,7 +929,7 @@ class Matrix { _buffer = cl_buffers_out.GetByKey(_size, _buffer); if (!_buffer.IsSet()) { - _buffer = new MatrixOpenCLBuffer(_size); + _buffer = new MatrixOpenCLBuffer(_size, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR); cl_buffers_out.Set(_size, _buffer); } @@ -1419,7 +1412,7 @@ class Matrix { static void MatMul(Matrix& source, Matrix& target, Matrix& output) { #ifdef MATRIX_USE_OPENCL // MatMulCL(source, target, output); - MatMulCLSingle(source, target, output); + MatMulCL(source, target, output); return; #endif @@ -1445,48 +1438,56 @@ class Matrix { * method. */ static void MatMulCL(Matrix& _source, Matrix& _target, Matrix& _output) { - if (_source.GetSize() != _target.GetRange(1)) { + if (_source.GetRange(1) != _target.GetRange(0)) { Alert("Inconsistent size of matrices!"); } - unsigned int _m_dim = _source.GetRange(0); - unsigned int _n_dim = _source.GetRange(1); - unsigned int _p_dim = _target.GetRange(1); + unsigned int _rows_a = _source.GetRange(0); + unsigned int _cols_a = _source.GetRange(1); + unsigned int _cols_b = _target.GetRange(1); - OpenCLBuffer* _in_1 = GetCLBufferInArg0(_source.GetSize()) PTR_DEREF GetBuffer(); - OpenCLBuffer* _in_2 = GetCLBufferInArg1(_target.GetSize()) PTR_DEREF GetBuffer(); - OpenCLBuffer* _out = GetCLBufferOutArg(_source.GetRange(0)) PTR_DEREF GetBuffer(); + OpenCLBuffer* _in_1 = GetCLBufferInArg0(_rows_a * _cols_a) PTR_DEREF GetBuffer(); + OpenCLBuffer* _in_2 = GetCLBufferInArg1(_cols_a * _cols_b) PTR_DEREF GetBuffer(); + OpenCLBuffer* _out = GetCLBufferOutArg(_rows_a * _cols_b) PTR_DEREF GetBuffer(); double _in_1_data[]; double _in_2_data[]; double _out_data[]; + // ArrayResize(_out_data, _out PTR_DEREF GetSizeItems()); + _source.GetRawArray(_in_1_data); _target.GetRawArray(_in_2_data); - cl_program_matmul REF_DEREF SetArg(0, _m_dim); - cl_program_matmul REF_DEREF SetArg(1, _n_dim); - cl_program_matmul REF_DEREF SetArg(2, _p_dim); - cl_program_matmul REF_DEREF SetArg(3, _in_1); - cl_program_matmul REF_DEREF SetArg(4, _in_2); - cl_program_matmul REF_DEREF SetArg(5, _out); - cl_program_matmul REF_DEREF SetArgLocalMem(6, _p_dim * sizeof(X)); + _in_1 PTR_DEREF Write(_in_1_data); + _in_2 PTR_DEREF Write(_in_2_data); + + cl_program_matmul REF_DEREF SetArg(0, _in_1); + cl_program_matmul REF_DEREF SetArg(1, _in_2); + cl_program_matmul REF_DEREF SetArg(2, _out); + cl_program_matmul REF_DEREF SetArg(3, (int)_rows_a); + cl_program_matmul REF_DEREF SetArg(4, (int)_cols_a); + cl_program_matmul REF_DEREF SetArg(5, (int)_cols_b); + ARRAY(unsigned int, _global_work_offset); ARRAY(unsigned int, _global_work_size); ARRAY(unsigned int, _local_work_size); + ArrayPush(_global_work_offset, 0U); - ArrayPush(_global_work_size, _n_dim); - ArrayPush(_global_work_size, _m_dim); - ArrayPush(_local_work_size, _m_dim); + ArrayPush(_global_work_offset, 0U); + ArrayPush(_global_work_size, (unsigned int)_rows_a); + ArrayPush(_global_work_size, (unsigned int)_cols_b); + ArrayPush(_local_work_size, 1U); + ArrayPush(_local_work_size, 1U); if (!cl_program_matmul REF_DEREF RunMany(2, _global_work_offset, _global_work_size, _local_work_size)) { - Alert("Errpr: Could not run Matrix::MatMulCL() over OpenCL!"); + Alert("Error: Could not run Matrix::MatMulCL() over OpenCL!"); DebugBreak(); } _out PTR_DEREF Read(_out_data); - // _output.SetShape(num_outputs); + _output.SetShape(_rows_a, _cols_b); } /** @@ -1498,13 +1499,13 @@ class Matrix { Alert("Inconsistent size of matrices!"); } - unsigned int _m_dim = _source.GetRange(0); - unsigned int _n_dim = _source.GetRange(1); - unsigned int _p_dim = _target.GetRange(1); + unsigned int _cols_a = _source.GetRange(0); + unsigned int _rows_a = _source.GetRange(1); + unsigned int _cols_b = _target.GetRange(1); - OpenCLBuffer* _in_1 = GetCLBufferInArg0(_n_dim * _p_dim) PTR_DEREF GetBuffer(); - OpenCLBuffer* _in_2 = GetCLBufferInArg1(_p_dim * _m_dim) PTR_DEREF GetBuffer(); - OpenCLBuffer* _out = GetCLBufferOutArg(_m_dim * _p_dim) PTR_DEREF GetBuffer(); + OpenCLBuffer* _in_1 = GetCLBufferInArg0(_rows_a * _cols_b) PTR_DEREF GetBuffer(); + OpenCLBuffer* _in_2 = GetCLBufferInArg1(_cols_b * _cols_a) PTR_DEREF GetBuffer(); + OpenCLBuffer* _out = GetCLBufferOutArg(_cols_a * _cols_b) PTR_DEREF GetBuffer(); double _in_1_data[]; double _in_2_data[]; @@ -1515,16 +1516,16 @@ class Matrix { _source.GetRawArray(_in_1_data); _target.GetRawArray(_in_2_data); - MatMulCL_CPU(_m_dim, _n_dim, _p_dim, _in_1_data, _in_2_data, _out_data); + MatMulCL_CPU(_cols_a, _rows_a, _cols_b, _in_1_data, _in_2_data, _out_data); - cl_program_matmul REF_DEREF SetArg(0, _n_dim); - cl_program_matmul REF_DEREF SetArg(1, _m_dim); - cl_program_matmul REF_DEREF SetArg(2, _p_dim); + cl_program_matmul REF_DEREF SetArg(0, (int)_rows_a); + cl_program_matmul REF_DEREF SetArg(1, (int)_cols_a); + cl_program_matmul REF_DEREF SetArg(2, (int)_cols_b); cl_program_matmul REF_DEREF SetArg(3, _in_1); cl_program_matmul REF_DEREF SetArg(4, _in_2); cl_program_matmul REF_DEREF SetArg(5, _out); - if (!cl_program_matmul_single REF_DEREF Run()) { + if (!cl_program_matmul REF_DEREF Run()) { Alert("Errpr: Could not run Matrix::MatMulCL() over OpenCL!"); DebugBreak(); } diff --git a/OpenCL.h b/OpenCL.h index ffeb1314f..b6fe7bc54 100644 --- a/OpenCL.h +++ b/OpenCL.h @@ -61,6 +61,7 @@ class OpenCLBuffer : public Dynamic { DebugBreak(); return; } + ArrayResize(_arr, buffer_size); CLBufferRead(buffer_handle, _arr); } @@ -147,6 +148,7 @@ class OpenCLProgram : public Dynamic { * your version isn't greater that the one already set in the buffer. */ void SetArg(int _index, double value) { CLSetKernelArg(kernel_handle, _index, value); } + void SetArg(int _index, int value) { CLSetKernelArg(kernel_handle, _index, value); } /** * Passes argument to the kernel. Will not set kernel argument if not needed. @@ -479,7 +481,7 @@ class OpenCL { /** * Allocates memory to be later passed to OpenCLProgram. */ - static OpenCLBuffer* Alloc(int _size) { return new OpenCLBuffer(_size); } + static OpenCLBuffer* Alloc(int _size, unsigned int _flags) { return new OpenCLBuffer(_size, _flags); } /** * Compiles given program and returns its id or -1 in case of error. diff --git a/tests/OpenCLTest.mq5 b/tests/OpenCLTest.mq5 index 800eb08ed..c26cf3153 100644 --- a/tests/OpenCLTest.mq5 +++ b/tests/OpenCLTest.mq5 @@ -40,7 +40,7 @@ int OnInit() { double result[] = {0}; - Ref buffer = OpenCL::Alloc(1 /* 1 double */); + Ref buffer = OpenCL::Alloc(1 /* 1 double */, CL_MEM_READ_WRITE); if (!program REF_DEREF Run(buffer.Ptr())) { Alert("Error running program!"); From dce2921eaa9431bba219e0675f925124df43a05d Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Fri, 26 Apr 2024 21:31:19 +0200 Subject: [PATCH 7/8] Refs #738. WIP. Matrix multiplication via OpenCL. Working on OpenCL buffers caching when passing matrices into kernel arguments. Got rid of templated version of Run() and RunMany(). --- Matrix.mqh | 207 +++++++++++++----------------- OpenCL.h | 298 ++++++++++--------------------------------- tests/OpenCLTest.mq5 | 5 +- 3 files changed, 159 insertions(+), 351 deletions(-) diff --git a/Matrix.mqh b/Matrix.mqh index eea71b0a3..45fbf5b6b 100644 --- a/Matrix.mqh +++ b/Matrix.mqh @@ -643,7 +643,7 @@ class MatrixDimension { } /** - * Extracts dimensions's values to the given array. Used internally. + * Extracts dimensions' values to the given array. Used internally. */ void FillArray(X& array[], int& offset) { int i; @@ -732,44 +732,6 @@ class MatrixDimension { enum ENUM_MATRIX_FLAGS { MATRIX_FLAGS_NONE, MATRIX_FLAGS_USE_OPENCL }; -/** - * Buffer used for CL operations. - */ -template -class MatrixOpenCLBuffer : public Dynamic { - // Flattened matrix data. - ARRAY(X, data); - - // Current version of the data. - long version; - - // CL buffer. - Ref buffer; - - public: - /** - * Constructor. - */ - MatrixOpenCLBuffer(int _size, unsigned int _flags) { - version = 0; - buffer = OpenCL::Alloc(_size, _flags); - ArrayResize(data, _size); - } - - /** - * Prepares buffer to be used by CL. Copies flattened data from the given matrix into buffer. - */ - void FillData(const Matrix& src) { - src.GetRawArray(data); - buffer REF_DEREF Write(data, version); - } - - /** - * Returns pointer to the CL buffer. - */ - OpenCLBuffer* GetBuffer() { return buffer.Ptr(); } -}; - /** * Matrix class. */ @@ -779,10 +741,14 @@ class Matrix { // First/root dimension. MatrixDimension* ptr_first_dimension; +#ifdef MATRIX_USE_OPENCL + // Map of data size -> CL buffer to be used e.g., by CL-based MatMul method. - static DictStruct>> cl_buffers_in_0; - static DictStruct>> cl_buffers_in_1; - static DictStruct>> cl_buffers_out; + static DictStruct> cl_buffers_in_0; + static DictStruct> cl_buffers_in_1; + static DictStruct> cl_buffers_out; + +#endif // Array with declaration of items per matrix's dimension. int dimensions[MATRIX_DIMENSIONS]; @@ -796,14 +762,25 @@ class Matrix { // Flags. int flags; - // Unique id of the matrix. - int uuid; + // Static counter, so each new matrix will have its own version. For new + // matrices new 32-bit range of versions are given and it should be more + // that enough. + static unsigned long version_counter; + + // Cache of previously flattened data. + ARRAY(X, flattened_cache); - // Static counter, so each matrix will have its own uuid. - static int uuid_counter; + // Version of the data that was flattened. + unsigned long flattened_cache_version; - // OpenCL program. + // Version of the data stored in dimensions arrays. Incremented after each + // change to this matrix. + unsigned long version; + + // OpenCL program for multi-core MatMul. static Ref cl_program_matmul; + + // OpenCL program for single-core MatMul. static Ref cl_program_matmul_single; /** @@ -814,7 +791,7 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif - uuid = uuid_counter++; + Initialize(); } /** @@ -826,7 +803,7 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif - uuid = uuid_counter++; + Initialize(); } /** @@ -837,7 +814,7 @@ class Matrix { #ifdef MATRIX_USE_OPENCL InitializeOpenCL(); #endif - uuid = uuid_counter++; + Initialize(); } /** @@ -853,8 +830,9 @@ class Matrix { InitializeOpenCL(); #endif - // We mark new matrix as unique one, even though we clone another matrix. - uuid = uuid_counter++; + // Note that we mark new matrix as unique one, even though we clone another + // matrix. + Initialize(); } /** @@ -864,6 +842,18 @@ class Matrix { private: Matrix(const Matrix* _right) {} + /** + * Initializes new or copy of another matrix. + */ + void Initialize() { + // Cache will have version lower that the data so matrix will be flattened in the first occasion. + flattened_cache_version = version_counter; + version = version_counter + 1; + + // Each new matrix will have its own 32-bit range of versions. + version_counter += UINT_MAX; + } + #ifdef MATRIX_USE_OPENCL /** @@ -891,13 +881,13 @@ class Matrix { /** * Returns/allocs and returns buffer of the given size to be used in CL operations as first input parameter. */ - static MatrixOpenCLBuffer* GetCLBufferInArg0(int _size) { - Ref> _buffer; + static OpenCLBuffer* GetCLBufferInArg0(int _size) { + Ref _buffer; _buffer = cl_buffers_in_0.GetByKey(_size, _buffer); if (!_buffer.IsSet()) { - _buffer = new MatrixOpenCLBuffer(_size, CL_MEM_READ_ONLY); + _buffer = new OpenCLBuffer(_size, CL_MEM_READ_ONLY); cl_buffers_in_0.Set(_size, _buffer); } @@ -907,13 +897,13 @@ class Matrix { /** * Returns/allocs and returns buffer of the given size to be used in CL operations as second input parameter. */ - static MatrixOpenCLBuffer* GetCLBufferInArg1(int _size) { - Ref> _buffer; + static OpenCLBuffer* GetCLBufferInArg1(int _size) { + Ref _buffer; _buffer = cl_buffers_in_1.GetByKey(_size, _buffer); if (!_buffer.IsSet()) { - _buffer = new MatrixOpenCLBuffer(_size, CL_MEM_READ_ONLY); + _buffer = new OpenCLBuffer(_size, CL_MEM_READ_ONLY); cl_buffers_in_1.Set(_size, _buffer); } @@ -923,13 +913,13 @@ class Matrix { /** * Returns/allocs and returns buffer of the given size to be used in CL operations as output parameter. */ - static MatrixOpenCLBuffer* GetCLBufferOutArg(int _size) { - Ref> _buffer; + static OpenCLBuffer* GetCLBufferOutArg(int _size) { + Ref _buffer; _buffer = cl_buffers_out.GetByKey(_size, _buffer); if (!_buffer.IsSet()) { - _buffer = new MatrixOpenCLBuffer(_size, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR); + _buffer = new OpenCLBuffer(_size, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR); cl_buffers_out.Set(_size, _buffer); } @@ -1411,7 +1401,6 @@ class Matrix { static void MatMul(Matrix& source, Matrix& target, Matrix& output) { #ifdef MATRIX_USE_OPENCL - // MatMulCL(source, target, output); MatMulCL(source, target, output); return; #endif @@ -1433,6 +1422,8 @@ class Matrix { } } +#ifdef MATRIX_USE_OPENCL + /** * Performs matrix multiplication via OpenCL. Note that MATRIX_USE_OPENCL must be defined in order matrix to use this * method. @@ -1446,48 +1437,32 @@ class Matrix { unsigned int _cols_a = _source.GetRange(1); unsigned int _cols_b = _target.GetRange(1); - OpenCLBuffer* _in_1 = GetCLBufferInArg0(_rows_a * _cols_a) PTR_DEREF GetBuffer(); - OpenCLBuffer* _in_2 = GetCLBufferInArg1(_cols_a * _cols_b) PTR_DEREF GetBuffer(); - OpenCLBuffer* _out = GetCLBufferOutArg(_rows_a * _cols_b) PTR_DEREF GetBuffer(); + // Reusable 0-offset. + static ARRAY(unsigned int, _global_work_offset) = {0U, 0U}; - double _in_1_data[]; - double _in_2_data[]; - double _out_data[]; + ARRAY(unsigned int, _global_work_size) = {(unsigned int)_rows_a, (unsigned int)_cols_b}; - // ArrayResize(_out_data, _out PTR_DEREF GetSizeItems()); + // @todo Make local work size adapt to output matrix size. + ARRAY(unsigned int, _local_work_size) = {1U, 1U}; - _source.GetRawArray(_in_1_data); - _target.GetRawArray(_in_2_data); - - _in_1 PTR_DEREF Write(_in_1_data); - _in_2 PTR_DEREF Write(_in_2_data); - - cl_program_matmul REF_DEREF SetArg(0, _in_1); - cl_program_matmul REF_DEREF SetArg(1, _in_2); - cl_program_matmul REF_DEREF SetArg(2, _out); + cl_program_matmul REF_DEREF SetArg(0, _source, OPENCL_MATRIX_ARG_IN_1); + cl_program_matmul REF_DEREF SetArg(1, _target, OPENCL_MATRIX_ARG_IN_2); + cl_program_matmul REF_DEREF SetArg(2, _output, OPENCL_MATRIX_ARG_OUT); cl_program_matmul REF_DEREF SetArg(3, (int)_rows_a); cl_program_matmul REF_DEREF SetArg(4, (int)_cols_a); cl_program_matmul REF_DEREF SetArg(5, (int)_cols_b); - ARRAY(unsigned int, _global_work_offset); - ARRAY(unsigned int, _global_work_size); - ARRAY(unsigned int, _local_work_size); - - ArrayPush(_global_work_offset, 0U); - ArrayPush(_global_work_offset, 0U); - ArrayPush(_global_work_size, (unsigned int)_rows_a); - ArrayPush(_global_work_size, (unsigned int)_cols_b); - ArrayPush(_local_work_size, 1U); - ArrayPush(_local_work_size, 1U); - - if (!cl_program_matmul REF_DEREF RunMany(2, _global_work_offset, _global_work_size, _local_work_size)) { + if (!cl_program_matmul REF_DEREF RunMany(2U, _global_work_offset, _global_work_size, _local_work_size)) { Alert("Error: Could not run Matrix::MatMulCL() over OpenCL!"); DebugBreak(); } - _out PTR_DEREF Read(_out_data); - + // Extracting data from + ARRAY(X, _out_data); + ArrayResize(_out_data, _rows_a * _cols_b); + _output.GetBuffer() PTR_DEREF Read(_out_data); _output.SetShape(_rows_a, _cols_b); + _output.FillFromArray(_out_data); } /** @@ -1535,27 +1510,7 @@ class Matrix { // _output.SetShape(num_outputs); } - static void MatMulCL_CPU(int Mdim, int Ndim, int Pdim, double& A[], double& B[], double& C[]) { - /* - for (int i = 0; i < Mdim; ++i) { - for (int j = 0; j < Pdim; ++j) { - C[i * Pdim + j] = 0.0; - for (int k = 0; k < Ndim; ++k) { - C[i * Pdim + j] += A[i * Ndim + k] * B[k * Pdim + j]; - } - } - } - */ - - int i, j, k; - for (i = 0; i < Ndim; i++) { - for (j = 0; j < Mdim; j++) { - for (k = 0; k < Pdim; k++) { // C(i,j) = sum(over k) A(i,k) * B(k,j) - C[i * Pdim + j] += A[i * Ndim + k] * B[k * Pdim + j]; - } - } - } - } +#endif /** * Performs matrix multiplication. @@ -1703,6 +1658,12 @@ class Matrix { return result; } + /** + * Fills matrix from flattened data. Shape of the array must be initialized + * before deflatenning. + */ + void Deflatten(const ARRAY_REF(X, array)) {} + /** * Initializer that generates tensors with a uniform distribution. */ @@ -2564,32 +2525,36 @@ class Matrix { } }; +#ifdef MATRIX_USE_OPENCL + #ifdef __MQL__ template -static int Matrix::uuid_counter = 0; +static unsigned long Matrix::version_counter = 0UL; template static Ref Matrix::cl_program_matmul; template static Ref Matrix::cl_program_matmul_single; template -static DictStruct>> Matrix::cl_buffers_in_0; +static DictStruct> Matrix::cl_buffers_in_0; template -static DictStruct>> Matrix::cl_buffers_in_1; +static DictStruct> Matrix::cl_buffers_in_1; template -static DictStruct>> Matrix::cl_buffers_out; +static DictStruct> Matrix::cl_buffers_out; #else template -static int Matrix::uuid_counter = 0; +static unsigned long Matrix::version_counter = 0UL; template static Ref Matrix::cl_program_matmul; template static Ref Matrix::cl_program_matmul_single; template -static DictStruct>> Matrix::cl_buffers_in_0; +static DictStruct> Matrix::cl_buffers_in_0; template -static DictStruct>> Matrix::cl_buffers_in_1; +static DictStruct> Matrix::cl_buffers_in_1; template -static DictStruct>> Matrix::cl_buffers_out; +static DictStruct> Matrix::cl_buffers_out; #endif #endif + +#endif \ No newline at end of file diff --git a/OpenCL.h b/OpenCL.h index b6fe7bc54..0ce41a1f1 100644 --- a/OpenCL.h +++ b/OpenCL.h @@ -9,6 +9,11 @@ // Forward declarations; class OpenCLProgram; +template +class Matrix; + +// Type of the matrix passed as argument to the OpenCLProgram. +enum ENUM_OPENCL_MATRIX_ARG { OPENCL_MATRIX_ARG_IN_1, OPENCL_MATRIX_ARG_IN_2, OPENCL_MATRIX_ARG_OUT }; /** * Memory buffer. @@ -20,8 +25,8 @@ class OpenCLBuffer : public Dynamic { // Allocated buffer size. int buffer_size; - // Buffer version. Should be incremented after each change. - long version; + // Version of the data. The same one that was passed to the Write() method. + unsigned long version; public: /** @@ -29,10 +34,15 @@ class OpenCLBuffer : public Dynamic { */ OpenCLBuffer(int _size, unsigned int _flags = CL_MEM_READ_WRITE); + /** + * Checks whether stored data version differs from the passed version. + */ + bool RequiresReupload(unsigned long _data_version) { return _data_version == ULONG_MAX || version != _data_version; } + /** * Writes/uploads data into buffer if needed. */ - void Write(const ARRAY_REF(double, _arr), long _arr_version = -1) { + void Write(const ARRAY_REF(double, _arr), unsigned long _data_version = ULONG_MAX) { if (ArraySize(_arr) > buffer_size) { Alert("Array passed is too large for the allocated buffer. Tries to pass ", ArraySize(_arr), " elements into buffer of size ", buffer_size, "."); @@ -40,15 +50,13 @@ class OpenCLBuffer : public Dynamic { return; } - // Do we need to reupload data into GPU? - if (_arr_version != -1 && _arr_version <= version) { - // Buffer has already up-to-date data. + if (!RequiresReupload(_data_version)) { return; } CLBufferWrite(buffer_handle, _arr); - version = (_arr_version != -1) ? _arr_version : (version + 1); + version = _data_version; } /** @@ -78,7 +86,7 @@ class OpenCLBuffer : public Dynamic { /** * Returns data version. */ - long GetVersion() { return version; } + unsigned long GetVersion() { return version; } /** * Returns handle to buffer. @@ -109,7 +117,7 @@ class OpenCLProgram : public Dynamic { int arg_handles[OPENCL_PROGRAM_MAX_ARGS]; // Version of argument data. Used to check if buffer needs to be reuploaded. - long arg_versions[OPENCL_PROGRAM_MAX_ARGS]; + unsigned long arg_versions[OPENCL_PROGRAM_MAX_ARGS]; public: /** @@ -140,6 +148,13 @@ class OpenCLProgram : public Dynamic { */ void SetArgLocalMem(int _index, unsigned long _mem_size) { CLSetKernelArgMemLocal(kernel_handle, _index, _mem_size); } + /** + * Checks whether given argument requires reupload of the buffer into GPU. + */ + bool RequiresReupload(int _index, OpenCLBuffer* _buffer, unsigned long _data_version) { + return _buffer PTR_DEREF GetHandle() != arg_handles[_index] || _data_version != arg_versions[_index]; + } + /** * Passes argument to the kernel. Will not set kernel argument if not needed. * @@ -157,10 +172,10 @@ class OpenCLProgram : public Dynamic { * which you can pass version of your data, so no reupload will take place if * your version isn't greater that the one already set in the buffer. */ - void SetArg(int _index, OpenCLBuffer* _buffer) { - if (_buffer PTR_DEREF GetHandle() == arg_handles[_index] && - _buffer PTR_DEREF GetVersion() >= arg_versions[_index]) { - // Already uploaded recent version. + void SetArg(int _index, OpenCLBuffer* _buffer, unsigned long _data_version) { + if (!RequiresReupload(_index, _buffer, _data_version)) { + // Buffer is already set via CLSetKernelArgMem() and argument's version + // is the same as _data_version. return; } @@ -174,115 +189,53 @@ class OpenCLProgram : public Dynamic { } /** - * Executes a single kernel. - */ - bool Run() { - if (!CLExecute(kernel_handle)) { - Alert("OpenCL error occured when tried to run kernel: ", GetLastError(), "!"); - return false; + * Passes matrix argument to the kernel. Will not upload data if not needed. + * + * The idea is to retrieve existing buffer that matches matrix size and its + * purpose. If such buffer already exists in the same version in the argument + * slot then no reupload will take place. + */ + template + void SetArg(int _index, Matrix& _matrix, ENUM_OPENCL_MATRIX_ARG _matrix_type) { + unsigned long _matrix_data_version = _matrix.GetVersion(); + OpenCLBuffer* _buffer = nullptr; + + switch (_matrix_type) { + case OPENCL_MATRIX_ARG_IN_1: + _buffer = GetCLBufferInArg0(_matrix.GetSize()); + break; + + case OPENCL_MATRIX_ARG_IN_2: + _buffer = GetCLBufferInArg1(_matrix.GetSize()); + break; + + case OPENCL_MATRIX_ARG_OUT: + _buffer = GetCLBufferOutArg(_matrix.GetSize()); + break; } - return true; - } + if (RequiresReupload(_index, _buffer, _matrix_data_version)) { + // Flattening matrix data in order to upload it into GPU. + double _flattened_data[]; + _matrix.GetRawArray(_flattened_data); - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a) { - SetArg(0, a); - return Run(); - } + _buffer PTR_DEREF Write(_flattened_data) - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a, B b) { - SetArg(0, a); - SetArg(1, b); - return Run(); - } - - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a, B b, C c) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - return Run(); - } - - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a, B b, C c, D d) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - return Run(); - } - - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a, B b, C c, D d, E e) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - return Run(); - } - - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a, B b, C c, D d, E e, F f) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - SetArg(5, f); - return Run(); + // Do we need to reupload the data? + SetArg(_index, _buffer, _matrix_data_version); + } } /** - * Executes a single kernel. Allows passing arugments to kernel. + * Executes a single kernel. */ - template - bool Run(A a, B b, C c, D d, E e, F f, G g) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - SetArg(5, f); - SetArg(6, g); - return Run(); - } + bool Run() { + if (!CLExecute(kernel_handle)) { + Alert("OpenCL error occured when tried to run kernel: ", GetLastError(), "!"); + return false; + } - /** - * Executes a single kernel. Allows passing arugments to kernel. - */ - template - bool Run(A a, B b, C c, D d, E e, F f, G g, H h) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - SetArg(5, f); - SetArg(6, g); - SetArg(7, h); - return Run(); + return true; } /** @@ -298,121 +251,6 @@ class OpenCLProgram : public Dynamic { return true; } - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a) { - SetArg(0, a); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b) { - SetArg(0, a); - SetArg(1, b); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b, C c) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b, C c, D d) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b, C c, D d, E e) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b, C c, D d, E e, F f) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - SetArg(5, f); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b, C c, D d, E e, F f, G g) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - SetArg(5, f); - SetArg(6, g); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - - /** - * Executes multiple kernels where work is subdivided among kernels. Allows passing arugments to kernels. - */ - template - bool RunMany(unsigned int _dimension, const ARRAY_REF(unsigned int, _global_work_offset), - const ARRAY_REF(unsigned int, _global_work_size), const ARRAY_REF(unsigned int, _local_work_size), A a, - B b, C c, D d, E e, F f, G g, H h) { - SetArg(0, a); - SetArg(1, b); - SetArg(2, c); - SetArg(3, d); - SetArg(4, e); - SetArg(5, f); - SetArg(6, g); - SetArg(7, h); - return RunMany(_dimension, _global_work_offset, _global_work_size, _local_work_size); - } - /** * Returns handle to OpenCL program. */ @@ -542,5 +380,7 @@ OpenCLBuffer::OpenCLBuffer(int _size, unsigned int _flags) { DebugBreak(); } buffer_size = _size; - version = 0; + // Ensuring there won't be initial version clash when checking if buffer data + // need to be reuploaded. + version = ULONG_MAX; } diff --git a/tests/OpenCLTest.mq5 b/tests/OpenCLTest.mq5 index c26cf3153..0e0d2ef50 100644 --- a/tests/OpenCLTest.mq5 +++ b/tests/OpenCLTest.mq5 @@ -42,7 +42,9 @@ int OnInit() { Ref buffer = OpenCL::Alloc(1 /* 1 double */, CL_MEM_READ_WRITE); - if (!program REF_DEREF Run(buffer.Ptr())) { + program REF_DEREF SetArg(0, buffer.Ptr(), ULONG_MAX); + + if (!program REF_DEREF Run()) { Alert("Error running program!"); } @@ -57,6 +59,7 @@ int OnInit() { Print("in2 shape: ", in2 PTR_DEREF GetRange(0), " x ", in2 PTR_DEREF GetRange(1)); out = in1 PTR_DEREF MatMul(in2); Print("out shape: ", out PTR_DEREF GetRange(0), " x ", out PTR_DEREF GetRange(1)); + Print("out data: ", out PTR_DEREF ToString()); delete in1; delete in2; From ad24e5b786147864e696c5d29380d8171baa2302 Mon Sep 17 00:00:00 2001 From: Adrian Kierzkowski Date: Wed, 1 May 2024 20:07:17 +0200 Subject: [PATCH 8/8] Refs #738. Matrix multiplication via OpenCL. Ready for testing. --- Matrix.matmul.naive.cl | 29 +++---- Matrix.mqh | 166 ++++++++++++++++++++++++++--------------- OpenCL.h | 34 +++++++-- tests/MatrixTest.mq5 | 15 ++-- tests/OpenCLTest.mq5 | 2 +- 5 files changed, 148 insertions(+), 98 deletions(-) diff --git a/Matrix.matmul.naive.cl b/Matrix.matmul.naive.cl index 612c2eb97..7ce5b64a1 100644 --- a/Matrix.matmul.naive.cl +++ b/Matrix.matmul.naive.cl @@ -1,23 +1,14 @@ #pragma OPENCL EXTENSION cl_khr_fp64 : enable -__kernel void matmul( - __global double* A, - __global double* B, - __global double* C, - int rowsA, - int colsA, - int colsB - ) -{ - int row = get_global_id(0); - int col = get_global_id(1); - - double sum = 0.0; - - for(int k = 0; k < colsA; ++k) { - sum += A[row * colsA + k] * B[k * colsB + col]; - //sum += col; +__kernel void matmul(__global double* A, __global double* B, __global double* C, int rowsA, int colsA, int colsB) { + int i = get_global_id(0); + int j = get_global_id(1); + + if (i < rowsA && j < colsB) { + float sum = 0.0f; + for (int k = 0; k < colsA; ++k) { + sum += A[i * colsA + k] * B[k * colsB + j]; + } + C[i * colsB + j] = sum; } - - C[row * colsB + col] = sum; } \ No newline at end of file diff --git a/Matrix.mqh b/Matrix.mqh index 45fbf5b6b..53d161160 100644 --- a/Matrix.mqh +++ b/Matrix.mqh @@ -878,6 +878,21 @@ class Matrix { #endif public: + /** + * Returns matrix's data version. + */ + unsigned long GetVersion() { return version; } + + /** + * Returns matrix flattened data version. + */ + unsigned long GetFlattenedCacheVersion() { return flattened_cache_version; } + + /** + * Acknowledges matrix that it's data has been changed. + */ + void Modified() { ++version; } + /** * Returns/allocs and returns buffer of the given size to be used in CL operations as first input parameter. */ @@ -958,6 +973,8 @@ class Matrix { num_dimensions = i; RecalculateSize(); + + Modified(); } void RecalculateSize() { @@ -1045,12 +1062,14 @@ class Matrix { size *= dimensions[i]; } } + + Modified(); } /** * Returns length of the given dimension. */ - int GetRange(int _dimension) { + int GetRangeRaw(int _dimension) { if (_dimension >= MATRIX_DIMENSIONS) { Print("Matrix::GetRange(): Dimension should be between 0 and ", MATRIX_DIMENSIONS - 1, ". Got ", _dimension, "!"); return -1; @@ -1059,6 +1078,22 @@ class Matrix { return dimensions[_dimension]; } + /** + * Returns length of the given dimension. In case that e.g., dimension 0 has values, dimension 1 will return range 1. + */ + int GetRange(int _dimension) { + if (_dimension < num_dimensions) { + return GetRangeRaw(_dimension); + } + + if (_dimension == num_dimensions) { + // Last dimension always contain values. + return 1; + } + + return 0; + } + /** * Returns total number of values the matrix contain of. */ @@ -1081,6 +1116,7 @@ class Matrix { int initial_container_size = ptr_first_dimension.DuplicateDimension(_level, _num); dimensions[_level] += _num * initial_container_size; RecalculateSize(); + Modified(); } /** @@ -1209,6 +1245,8 @@ class Matrix { } accessor = _value; + + Modified(); } /** @@ -1231,6 +1269,7 @@ class Matrix { void Add(X value) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_ADD, value); + Modified(); } } @@ -1245,6 +1284,7 @@ class Matrix { void Sub(X value) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_SUBTRACT, value); + Modified(); } } @@ -1259,6 +1299,7 @@ class Matrix { void Mul(X value) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_MULTIPLY, value); + Modified(); } } @@ -1273,6 +1314,7 @@ class Matrix { void Div(X value) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_DIVIDE, value); + Modified(); } } @@ -1282,6 +1324,7 @@ class Matrix { void Fill(X value) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_FILL, value); + Modified(); } } @@ -1291,6 +1334,7 @@ class Matrix { void FillRandom(int _seed = -1) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_FILL_RANDOM, _seed); + Modified(); } } @@ -1300,6 +1344,7 @@ class Matrix { void FillRandom(X _start, X _end, int _seed = -1) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_FILL_RANDOM_RANGE, _start, _end, _seed); + Modified(); } } @@ -1309,6 +1354,7 @@ class Matrix { void FillPosAdd() { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_FILL_POS_ADD); + Modified(); } } @@ -1318,11 +1364,12 @@ class Matrix { void FillPosMul() { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_FILL_POS_MUL); + Modified(); } } /** - * Replaces existing matrix's values by random value from a given range. + * Calculates sum fo all matrix's values. */ X Sum() { X _out1 = 0, _out2; @@ -1370,9 +1417,13 @@ class Matrix { return MinOf((X)0); } + /** + * Performs power operation on all matrix's values. + */ void Power(X value) { if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_POWER, value); + Modified(); } } @@ -1420,6 +1471,8 @@ class Matrix { output[output_idx] += source[input_idx].Val() * target[output_idx][input_idx].Val(); } } + + output.Modified(); } #ifdef MATRIX_USE_OPENCL @@ -1440,10 +1493,16 @@ class Matrix { // Reusable 0-offset. static ARRAY(unsigned int, _global_work_offset) = {0U, 0U}; - ARRAY(unsigned int, _global_work_size) = {(unsigned int)_rows_a, (unsigned int)_cols_b}; + // Local work size (up to 8 x 8). + static ARRAY(unsigned int, _local_work_size) = {MathMin(_rows_a, 8U), MathMin(_cols_a, 8U)}; - // @todo Make local work size adapt to output matrix size. - ARRAY(unsigned int, _local_work_size) = {1U, 1U}; + // Our global size could be greater that data. + ARRAY(unsigned int, _global_work_size) = { + (unsigned int)MathCeil((double)_rows_a / _local_work_size[0]) * _local_work_size[0], + (unsigned int)MathCeil((double)_cols_a / _local_work_size[1]) * _local_work_size[1], + }; + + _output.SetShape(_rows_a, _cols_b); cl_program_matmul REF_DEREF SetArg(0, _source, OPENCL_MATRIX_ARG_IN_1); cl_program_matmul REF_DEREF SetArg(1, _target, OPENCL_MATRIX_ARG_IN_2); @@ -1460,56 +1519,11 @@ class Matrix { // Extracting data from ARRAY(X, _out_data); ArrayResize(_out_data, _rows_a * _cols_b); - _output.GetBuffer() PTR_DEREF Read(_out_data); + cl_program_matmul REF_DEREF GetArgBuffer(2) PTR_DEREF Read(_out_data); _output.SetShape(_rows_a, _cols_b); _output.FillFromArray(_out_data); } - /** - * Performs matrix multiplication via OpenCL. Note that MATRIX_USE_OPENCL must be defined in order matrix to use this - * method. - */ - static void MatMulCLSingle(Matrix& _source, Matrix& _target, Matrix& _output) { - if (_source.GetRange(1) != _target.GetRange(0)) { - Alert("Inconsistent size of matrices!"); - } - - unsigned int _cols_a = _source.GetRange(0); - unsigned int _rows_a = _source.GetRange(1); - unsigned int _cols_b = _target.GetRange(1); - - OpenCLBuffer* _in_1 = GetCLBufferInArg0(_rows_a * _cols_b) PTR_DEREF GetBuffer(); - OpenCLBuffer* _in_2 = GetCLBufferInArg1(_cols_b * _cols_a) PTR_DEREF GetBuffer(); - OpenCLBuffer* _out = GetCLBufferOutArg(_cols_a * _cols_b) PTR_DEREF GetBuffer(); - - double _in_1_data[]; - double _in_2_data[]; - double _out_data[]; - - ArrayResize(_out_data, _out PTR_DEREF GetSizeItems()); - - _source.GetRawArray(_in_1_data); - _target.GetRawArray(_in_2_data); - - MatMulCL_CPU(_cols_a, _rows_a, _cols_b, _in_1_data, _in_2_data, _out_data); - - cl_program_matmul REF_DEREF SetArg(0, (int)_rows_a); - cl_program_matmul REF_DEREF SetArg(1, (int)_cols_a); - cl_program_matmul REF_DEREF SetArg(2, (int)_cols_b); - cl_program_matmul REF_DEREF SetArg(3, _in_1); - cl_program_matmul REF_DEREF SetArg(4, _in_2); - cl_program_matmul REF_DEREF SetArg(5, _out); - - if (!cl_program_matmul REF_DEREF Run()) { - Alert("Errpr: Could not run Matrix::MatMulCL() over OpenCL!"); - DebugBreak(); - } - - _out PTR_DEREF Read(_out_data); - - // _output.SetShape(num_outputs); - } - #endif /** @@ -1563,6 +1577,7 @@ class Matrix { void operator+=(const Matrix& r) { if (ptr_first_dimension && r.ptr_first_dimension) { ptr_first_dimension.Op(r.ptr_first_dimension, MATRIX_OPERATION_ADD); + Modified(); } } @@ -1585,6 +1600,7 @@ class Matrix { void operator-=(const Matrix& r) { if (ptr_first_dimension && r.ptr_first_dimension) { ptr_first_dimension.Op(r.ptr_first_dimension, MATRIX_OPERATION_SUBTRACT); + Modified(); } } @@ -1607,6 +1623,7 @@ class Matrix { void operator*=(const Matrix& r) { if (ptr_first_dimension && r.ptr_first_dimension) { ptr_first_dimension.Op(r.ptr_first_dimension, MATRIX_OPERATION_MULTIPLY); + Modified(); } } @@ -1629,16 +1646,29 @@ class Matrix { void operator/=(const Matrix& r) { if (ptr_first_dimension && r.ptr_first_dimension) { ptr_first_dimension.Op(r.ptr_first_dimension, MATRIX_OPERATION_DIVIDE); + Modified(); } } /** * Fills array with all values from the matrix. */ - void GetRawArray(X& array[]) const { - ArrayResize(array, GetSize()); + void GetRawArray(X& array[]) { + if (flattened_cache_version == version) { + // No need to flatten again as our cache is up to date. + ArrayCopy(array, flattened_cache); + return; + } + + // Filling target array with flattened matrix data. int offset = 0; + ArrayResize(array, GetSize()); ptr_first_dimension.FillArray(array, offset); + + // Copying target array into our flattened data cache. + ArrayResize(flattened_cache, 0, 1024); + ArrayCopy(flattened_cache, array); + flattened_cache_version = version; } /** @@ -1658,12 +1688,6 @@ class Matrix { return result; } - /** - * Fills matrix from flattened data. Shape of the array must be initialized - * before deflatenning. - */ - void Deflatten(const ARRAY_REF(X, array)) {} - /** * Initializer that generates tensors with a uniform distribution. */ @@ -1696,6 +1720,8 @@ class Matrix { int offset = 0; ptr_first_dimension.FromArray(_array, offset); + + Modified(); } /** @@ -1703,17 +1729,24 @@ class Matrix { */ void FillTruncatedNormal(X _mean, X _stddev, int _seeds = -1) { Print("Matrix::FillTruncatedNormal() is not yet implemented!"); + Modified(); } /** * The Glorot normal initializer, also called Xavier normal initializer. */ - void FillGlorotNormal(int _seed = -1) { Print("Matrix::FillGlorotNormal() is not yet implemented!"); } + void FillGlorotNormal(int _seed = -1) { + Print("Matrix::FillGlorotNormal() is not yet implemented!"); + Modified(); + } /** * The Glorot uniform initializer, also called Xavier uniform initializer. */ - void FillGlorotUniform(int _seed = -1) { Print("Matrix::FillGlorotUniform() is not yet implemented!"); } + void FillGlorotUniform(int _seed = -1) { + Print("Matrix::FillGlorotUniform() is not yet implemented!"); + Modified(); + } /** * Initializer that generates the identity matrix. @@ -1736,7 +1769,10 @@ class Matrix { /** * Initializer that generates an orthogonal matrix. */ - void FillOrthogonal(X _gain, int _seed = -1) { Print("Matrix::FillOrthogonal() is not yet implemented!"); } + void FillOrthogonal(X _gain, int _seed = -1) { + Print("Matrix::FillOrthogonal() is not yet implemented!"); + Modified(); + } /** * Calculates absolute difference between this tensor and given one using optional weights tensor. @@ -1801,6 +1837,7 @@ class Matrix { void ReduceSimple(bool _only_last_dimension = true, ENUM_MATRIX_OPERATION _reduce_op = MATRIX_OPERATION_SUM) { if (ptr_first_dimension != NULL) { ptr_first_dimension.ReduceSimple(_only_last_dimension ? GetDimensions() - 1 : 0, _reduce_op); + Modified(); } } @@ -1816,6 +1853,7 @@ class Matrix { } RecalculateSize(); + Modified(); } /** @@ -2056,6 +2094,7 @@ class Matrix { int _out3; if (ptr_first_dimension) { ptr_first_dimension.Op(MATRIX_OPERATION_RELU, 0, 0, 0, _out1, _out2, _out3); + Modified(); } } @@ -2086,6 +2125,7 @@ class Matrix { } else { this[_1d][_2d][_3d][_4d][_5d] = value; } + Modified(); } Matrix* GetConv2d(int _in_channels, int _out_channels, int _krn_1d, int _krn_2d, @@ -2370,6 +2410,8 @@ class Matrix { Print("Matrix::ChunkOp(): Invalid operation ", EnumToString(_op), "!"); } + // Note: ChuckOp() does not modify the matrix. + return 0; } diff --git a/OpenCL.h b/OpenCL.h index 0ce41a1f1..b6aa15303 100644 --- a/OpenCL.h +++ b/OpenCL.h @@ -119,6 +119,10 @@ class OpenCLProgram : public Dynamic { // Version of argument data. Used to check if buffer needs to be reuploaded. unsigned long arg_versions[OPENCL_PROGRAM_MAX_ARGS]; + // Buffers passed as arguments. Note that we store weak references, so we + // need to check if buffer still exists. + WeakRef arg_buffers[OPENCL_PROGRAM_MAX_ARGS]; + public: /** * Constructor. @@ -127,6 +131,7 @@ class OpenCLProgram : public Dynamic { for (int i = 0; i < OPENCL_PROGRAM_MAX_ARGS; ++i) { arg_handles[i] = INVALID_HANDLE; arg_versions[i] = -1; + arg_buffers[i] = nullptr; } } @@ -186,6 +191,22 @@ class OpenCLProgram : public Dynamic { // Storing buffer version in the argument slot. arg_versions[_index] = _buffer PTR_DEREF GetVersion(); + + // Storing pointer to the buffer as weak reference. + arg_buffers[_index] = _buffer; + } + + /** + * Returns buffer passed as argument. + */ + OpenCLBuffer* GetArgBuffer(int _index) { + if (!arg_buffers[_index].ObjectExists()) { + Alert("Error: Trying to retrieve buffer at argument index ", _index, + ", but the buffer was either not set or deleted."); + DebugBreak(); + } + + return arg_buffers[_index].Ptr(); } /** @@ -202,15 +223,15 @@ class OpenCLProgram : public Dynamic { switch (_matrix_type) { case OPENCL_MATRIX_ARG_IN_1: - _buffer = GetCLBufferInArg0(_matrix.GetSize()); + _buffer = _matrix.GetCLBufferInArg0(_matrix.GetSize()); break; case OPENCL_MATRIX_ARG_IN_2: - _buffer = GetCLBufferInArg1(_matrix.GetSize()); + _buffer = _matrix.GetCLBufferInArg1(_matrix.GetSize()); break; case OPENCL_MATRIX_ARG_OUT: - _buffer = GetCLBufferOutArg(_matrix.GetSize()); + _buffer = _matrix.GetCLBufferOutArg(_matrix.GetSize()); break; } @@ -218,11 +239,8 @@ class OpenCLProgram : public Dynamic { // Flattening matrix data in order to upload it into GPU. double _flattened_data[]; _matrix.GetRawArray(_flattened_data); - - _buffer PTR_DEREF Write(_flattened_data) - - // Do we need to reupload the data? - SetArg(_index, _buffer, _matrix_data_version); + _buffer PTR_DEREF Write(_flattened_data); + SetArg(_index, _buffer, _matrix_data_version); } } diff --git a/tests/MatrixTest.mq5 b/tests/MatrixTest.mq5 index f1d3d679e..8f0aa09b2 100644 --- a/tests/MatrixTest.mq5 +++ b/tests/MatrixTest.mq5 @@ -595,30 +595,29 @@ int OnInit() { assertTrueOrFail(matrix_24_relu.ToString(false, 4) == "[0.0000,0.0000,0.0000,1.0000,2.0000]", "Matrix::Relu(): Invalid result!"); - Matrix matrix_26("[1, 2, 3]"); + Matrix matrix_26("[[1, 2, 3]]"); Matrix matrix_27( "[[1, 1, 1]" " [0, 1, 0]" " [0, 0, 2]]"); Matrix* ptr_matrix_27_matmul = matrix_26 ^ matrix_27; - - assertTrueOrFail(ptr_matrix_27_matmul.ToString(false, 2) == "[6.00,2.00,6.00]", + assertTrueOrFail(ptr_matrix_27_matmul.ToString(false, 2) == "[[1.00,3.00,7.00]]", "Matrix::operator=(MatrixDimension): Invalid result!"); - Matrix matrix_27_bias("[1, 0, 5]"); + Matrix matrix_27_bias("[[1, 0, 5]]"); Matrix* ptr_matrix_27_add = ptr_matrix_27_matmul + matrix_27_bias; - assertTrueOrFail(ptr_matrix_27_add.ToString(false, 0) == "[7,2,11]", "Matrix::operator+(Matrix): Invalid result!"); + assertTrueOrFail(ptr_matrix_27_add.ToString(false, 0) == "[[2,3,12]]", "Matrix::operator+(Matrix): Invalid result!"); Matrix* ptr_matrix_27_sub = ptr_matrix_27_matmul - matrix_27_bias; - assertTrueOrFail(ptr_matrix_27_sub.ToString(false, 0) == "[5,2,1]", "Matrix::operator+(Matrix): Invalid result!"); + assertTrueOrFail(ptr_matrix_27_sub.ToString(false, 0) == "[[0,3,2]]", "Matrix::operator+(Matrix): Invalid result!"); Matrix* ptr_matrix_27_mul = ptr_matrix_27_matmul * matrix_27_bias; - assertTrueOrFail(ptr_matrix_27_mul.ToString(false, 0) == "[6,0,30]", "Matrix::operator+(Matrix): Invalid result!"); + assertTrueOrFail(ptr_matrix_27_mul.ToString(false, 0) == "[[1,0,35]]", "Matrix::operator+(Matrix): Invalid result!"); - delete ptr_matrix_27_matmul; delete ptr_matrix_27_add; + delete ptr_matrix_27_matmul; delete ptr_matrix_27_sub; delete ptr_matrix_27_mul; diff --git a/tests/OpenCLTest.mq5 b/tests/OpenCLTest.mq5 index 0e0d2ef50..65eb84a19 100644 --- a/tests/OpenCLTest.mq5 +++ b/tests/OpenCLTest.mq5 @@ -55,7 +55,7 @@ int OnInit() { Matrix *in1, *in2, *out; in1 = Matrix::CreateFromString("[[1,2,3], [4,5,6]]"); // 2 x 3 Print("in1 shape: ", in1 PTR_DEREF GetRange(0), " x ", in1 PTR_DEREF GetRange(1)); - in2 = Matrix::CreateFromString("[[7,8,9,10], [11,12,13,14], [15,16,17,18]]"); // 3 x 4 + in2 = Matrix::CreateFromString("[[7,8,9,10,11], [11,12,13,14,16], [15,16,17,18,19]]"); // 3 x 5 Print("in2 shape: ", in2 PTR_DEREF GetRange(0), " x ", in2 PTR_DEREF GetRange(1)); out = in1 PTR_DEREF MatMul(in2); Print("out shape: ", out PTR_DEREF GetRange(0), " x ", out PTR_DEREF GetRange(1));