@@ -164,7 +164,7 @@ Main::Main(QWidget *parent) :
statusBar()->addPermanentWidget(ui->chainStatus);
statusBar()->addPermanentWidget(ui->blockCount);
- ui->blockCount->setText(QString("PV%2 D%3 H%4 v%5").arg(eth::c_protocolVersion).arg(c_databaseVersion).arg(c_ethashVersion).arg(dev::Version));
+ ui->blockCount->setText(QString("PV%2 D%3 %4-%5 v%6").arg(eth::c_protocolVersion).arg(c_databaseVersion).arg(QString::fromStdString(ProofOfWork::name())).arg(ProofOfWork::revision()).arg(dev::Version));
connect(ui->ourAccounts->model(), SIGNAL(rowsMoved(const QModelIndex &, int, int, const QModelIndex &, int)), SLOT(ourAccountsRowsMoved()));
@@ -952,7 +952,7 @@ void Main::on_preview_triggered()
void Main::refreshMining()
{
- MineProgress p = ethereum()->miningProgress();
+ MiningProgress p = ethereum()->miningProgress();
ui->mineStatus->setText(ethereum()->isMining() ? QString("%1s @ %2kH/s").arg(p.ms / 1000).arg(p.ms ? p.hashes / p.ms : 0) : "Not mining");
if (!ui->miningView->isVisible())
return;
@@ -1481,7 +1481,7 @@ void Main::on_blocks_currentItemChanged()
s << "Difficulty: " << info.difficulty << "" << "
";
if (info.number)
{
- auto e = Ethasher::eval(info);
+ auto e = EthashAux::eval(info);
s << "Proof-of-Work: " << e.value << " <= " << (h256)u256((bigint(1) << 256) / info.difficulty) << " (mixhash: " << e.mixHash.abridged() << ")" << "
";
s << "Parent: " << info.parentHash << "" << "
";
}
@@ -1510,7 +1510,7 @@ void Main::on_blocks_currentItemChanged()
s << line << "Nonce: " << uncle.nonce << "" << "";
s << line << "Hash w/o nonce: " << uncle.headerHash(WithoutNonce) << "" << "";
s << line << "Difficulty: " << uncle.difficulty << "" << "";
- auto e = Ethasher::eval(uncle);
+ auto e = EthashAux::eval(uncle);
s << line << "Proof-of-Work: " << e.value << " <= " << (h256)u256((bigint(1) << 256) / uncle.difficulty) << " (mixhash: " << e.mixHash.abridged() << ")" << "";
}
if (info.parentHash)
@@ -1752,6 +1752,11 @@ void Main::on_clearPending_triggered()
refreshAll();
}
+void Main::on_retryUnknown_triggered()
+{
+ ethereum()->retryUnkonwn();
+}
+
void Main::on_killBlockchain_triggered()
{
writeSettings();
diff --git a/alethzero/MainWin.h b/alethzero/MainWin.h
index a5c74eeaa..a8579ed01 100644
--- a/alethzero/MainWin.h
+++ b/alethzero/MainWin.h
@@ -163,6 +163,7 @@ private slots:
void on_usePrivate_triggered();
void on_turboMining_triggered();
void on_jitvm_triggered();
+ void on_retryUnknown_triggered();
// Debugger
void on_debugCurrent_triggered();
diff --git a/alethzero/MiningView.cpp b/alethzero/MiningView.cpp
index 63d1fcf99..e020408ea 100644
--- a/alethzero/MiningView.cpp
+++ b/alethzero/MiningView.cpp
@@ -36,7 +36,7 @@ using namespace dev::eth;
// types
using dev::eth::MineInfo;
-using dev::eth::MineProgress;
+using dev::eth::MiningProgress;
// functions
using dev::toString;
@@ -50,12 +50,13 @@ MiningView::MiningView(QWidget* _p): QWidget(_p)
{
}
-void MiningView::appendStats(list const& _i, MineProgress const& _p)
+void MiningView::appendStats(list const& _i, MiningProgress const& _p)
{
+ (void)_p;
if (_i.empty())
return;
- unsigned o = m_values.size();
+/* unsigned o = m_values.size();
for (MineInfo const& i: _i)
{
m_values.push_back(i.best);
@@ -91,7 +92,7 @@ void MiningView::appendStats(list const& _i, MineProgress const& _p)
m_completes.erase(remove_if(m_completes.begin(), m_completes.end(), [](int i){return i < 0;}), m_completes.end());
m_progress = _p;
- update();
+ update();*/
}
void MiningView::resetStats()
@@ -101,6 +102,7 @@ void MiningView::resetStats()
void MiningView::paintEvent(QPaintEvent*)
{
+ /*
Grapher g;
QPainter p(this);
@@ -115,4 +117,5 @@ void MiningView::paintEvent(QPaintEvent*)
g.ruleY(r - 1, QColor(128, 128, 128));
for (auto r: m_completes)
g.ruleY(r, QColor(192, 64, 64));
+ */
}
diff --git a/alethzero/MiningView.h b/alethzero/MiningView.h
index 8f3135f75..65b9f2ec9 100644
--- a/alethzero/MiningView.h
+++ b/alethzero/MiningView.h
@@ -42,14 +42,14 @@ class MiningView: public QWidget
public:
MiningView(QWidget* _p = nullptr);
- void appendStats(std::list const& _l, dev::eth::MineProgress const& _p);
+ void appendStats(std::list const& _l, dev::eth::MiningProgress const& _p);
void resetStats();
protected:
virtual void paintEvent(QPaintEvent*);
private:
- dev::eth::MineProgress m_progress;
+ dev::eth::MiningProgress m_progress;
unsigned m_duration = 300;
std::vector m_values;
std::vector m_bests;
diff --git a/alethzero/Transact.cpp b/alethzero/Transact.cpp
index f1f7477fe..1ebdf9e23 100644
--- a/alethzero/Transact.cpp
+++ b/alethzero/Transact.cpp
@@ -37,7 +37,7 @@
#include
#include
#include
-#ifndef _MSC_VER
+#if ETH_SERPENT
#include
#include
#endif
@@ -220,7 +220,7 @@ static tuple, bytes, string> userInputToCode(string const& _user,
errors.push_back("Solidity: Uncaught exception");
}
}
-#ifndef _MSC_VER
+#if ETH_SERPENT
else if (sourceIsSerpent(_user))
{
try
diff --git a/eth/main.cpp b/eth/main.cpp
index e9af192f9..db6acbbca 100644
--- a/eth/main.cpp
+++ b/eth/main.cpp
@@ -32,6 +32,8 @@
#include
#include
#include
+#include
+#include
#include
#include
#include
@@ -44,7 +46,6 @@
#include
#include
#endif
-#include
#include "BuildInfo.h"
using namespace std;
using namespace dev;
@@ -111,6 +112,7 @@ void help()
<< " -b,--bootstrap Connect to the default Ethereum peerserver." << endl
<< " -B,--block-fees Set the block fee profit in the reference unit e.g. ยข (Default: 15)." << endl
<< " -c,--client-name Add a name to your client's version string (default: blank)." << endl
+ << " -C,--check-pow Check PoW credentials for validity." << endl
<< " -d,--db-path Load database from path (default: ~/.ethereum " << endl
<< " /Etherum or Library/Application Support/Ethereum)." << endl
<< " -D,--create-dag Create the DAG in preparation for mining on given block and exit." << endl
@@ -126,26 +128,26 @@ void help()
#if ETH_JSONRPC
<< " -j,--json-rpc Enable JSON-RPC server (default: off)." << endl
<< " --json-rpc-port Specify JSON-RPC server port (implies '-j', default: " << SensibleHttpPort << ")." << endl
+#endif
+#if ETH_EVMJIT
+ << " -J,--jit Enable EVM JIT (default: off)." << endl
#endif
<< " -K,--kill First kill the blockchain." << endl
<< " --listen-ip Listen on the given port for incoming connections (default: 30303)." << endl
<< " -l,--listen Listen on the given IP for incoming connections (default: 0.0.0.0)." << endl
<< " -u,--public-ip Force public ip to given (default: auto)." << endl
<< " -m,--mining Enable mining, optionally for a specified number of blocks (Default: off)" << endl
- << " -n,--upnp Use upnp for NAT (default: on)." << endl
+ << " -n,-u,--upnp Use upnp for NAT (default: on)." << endl
<< " -o,--mode Start a full node or a peer node (Default: full)." << endl
<< " -p,--port Connect to remote port (default: 30303)." << endl
<< " -P,--priority <0 - 100> Default % priority of a transaction (default: 50)." << endl
<< " -R,--rebuild First rebuild the blockchain from the existing database." << endl
<< " -r,--remote Connect to remote host (default: none)." << endl
<< " -s,--secret Set the secret key for use with send command (default: auto)." << endl
- << " -t,--miners Number of mining threads to start (Default: " << thread::hardware_concurrency() << ")" << endl
+ << " -S,--temporary-secret Set the secret key for use with send command, for this session only." << endl
<< " -v,--verbosity <0 - 9> Set the log verbosity from 0 to 9 (Default: 8)." << endl
<< " -x,--peers Attempt to connect to given number of peers (Default: 5)." << endl
<< " -V,--version Show the version and exit." << endl
-#if ETH_EVMJIT
- << " --jit Use EVM JIT (default: off)." << endl
-#endif
;
exit(0);
}
@@ -210,7 +212,7 @@ void doInitDAG(unsigned _n)
BlockInfo bi;
bi.number = _n;
cout << "Initializing DAG for epoch beginning #" << (bi.number / 30000 * 30000) << " (seedhash " << bi.seedHash().abridged() << "). This will take a while." << endl;
- Ethasher::get()->full(bi);
+ Ethash::prep(bi);
exit(0);
}
@@ -271,10 +273,10 @@ int main(int argc, char** argv)
/// Mining params
unsigned mining = ~(unsigned)0;
- int miners = -1;
bool forceMining = false;
- KeyPair us = KeyPair::create();
- Address coinbase = us.address();
+ KeyPair sigKey = KeyPair::create();
+ Secret sessionSecret;
+ Address coinbase = sigKey.address();
/// Structured logging params
bool structuredLogging = false;
@@ -290,7 +292,7 @@ int main(int argc, char** argv)
if (b.size())
{
RLP config(b);
- us = KeyPair(config[0].toHash());
+ sigKey = KeyPair(config[0].toHash());
coinbase = config[1].toHash();
}
@@ -338,7 +340,7 @@ int main(int argc, char** argv)
exportFrom = argv[++i];
else if (arg == "--only" && i + 1 < argc)
exportTo = exportFrom = argv[++i];
- else if ((arg == "-n" || arg == "--upnp") && i + 1 < argc)
+ else if ((arg == "-n" || arg == "-u" || arg == "--upnp") && i + 1 < argc)
{
string m = argv[++i];
if (isTrue(m))
@@ -362,7 +364,7 @@ int main(int argc, char** argv)
{
coinbase = h160(fromHex(argv[++i], WhenError::Throw));
}
- catch (BadHexCharacter& _e)
+ catch (BadHexCharacter&)
{
cerr << "Bad hex in " << arg << " option: " << argv[i] << endl;
return -1;
@@ -373,7 +375,9 @@ int main(int argc, char** argv)
return -1;
}
else if ((arg == "-s" || arg == "--secret") && i + 1 < argc)
- us = KeyPair(h256(fromHex(argv[++i])));
+ sigKey = KeyPair(h256(fromHex(argv[++i])));
+ else if ((arg == "-S" || arg == "--session-secret") && i + 1 < argc)
+ sessionSecret = h256(fromHex(argv[++i]));
else if (arg == "--structured-logging-format" && i + 1 < argc)
structuredLoggingFormat = string(argv[++i]);
else if (arg == "--structured-logging")
@@ -399,6 +403,43 @@ int main(int argc, char** argv)
return -1;
}
}
+ else if ((arg == "-C" || arg == "--check-pow") && i + 4 < argc)
+ {
+ string m;
+ try
+ {
+ BlockInfo bi;
+ m = boost::to_lower_copy(string(argv[++i]));
+ h256 powHash(m);
+ m = boost::to_lower_copy(string(argv[++i]));
+ h256 seedHash;
+ if (m.size() == 64 || m.size() == 66)
+ seedHash = h256(m);
+ else
+ seedHash = EthashAux::seedHash(stol(m));
+ m = boost::to_lower_copy(string(argv[++i]));
+ bi.difficulty = u256(m);
+ auto boundary = bi.boundary();
+ m = boost::to_lower_copy(string(argv[++i]));
+ bi.nonce = h64(m);
+ auto r = EthashAux::eval(seedHash, powHash, bi.nonce);
+ bool valid = r.value < boundary;
+ cout << (valid ? "VALID :-)" : "INVALID :-(") << endl;
+ cout << r.value << (valid ? " < " : " >= ") << boundary << endl;
+ cout << " where " << boundary << " = 2^256 / " << bi.difficulty << endl;
+ cout << " and " << r.value << " = ethash(" << powHash << ", " << bi.nonce << ")" << endl;
+ cout << " with seed as " << seedHash << endl;
+ if (valid)
+ cout << "(mixHash = " << r.mixHash << ")" << endl;
+ cout << "SHA3( light(seed) ) = " << sha3(bytesConstRef((byte const*)EthashAux::light(seedHash), EthashAux::params(seedHash).cache_size)) << endl;
+ exit(0);
+ }
+ catch (...)
+ {
+ cerr << "Bad " << arg << " option: " << m << endl;
+ return -1;
+ }
+ }
else if ((arg == "-B" || arg == "--block-fees") && i + 1 < argc)
{
try
@@ -477,8 +518,6 @@ int main(int argc, char** argv)
g_logVerbosity = atoi(argv[++i]);
else if ((arg == "-x" || arg == "--peers") && i + 1 < argc)
peers = atoi(argv[++i]);
- else if ((arg == "-t" || arg == "--miners") && i + 1 < argc)
- miners = atoi(argv[++i]);
else if ((arg == "-o" || arg == "--mode") && i + 1 < argc)
{
string m = argv[++i];
@@ -492,15 +531,12 @@ int main(int argc, char** argv)
return -1;
}
}
- else if (arg == "--jit")
- {
#if ETH_EVMJIT
+ else if (arg == "-J" || arg == "--jit")
+ {
jit = true;
-#else
- cerr << "EVM JIT not enabled" << endl;
- return -1;
-#endif
}
+#endif
else if (arg == "-h" || arg == "--help")
help();
else if (arg == "-V" || arg == "--version")
@@ -514,10 +550,13 @@ int main(int argc, char** argv)
{
RLPStream config(2);
- config << us.secret() << coinbase;
+ config << sigKey.secret() << coinbase;
writeFile(configFile, config.out());
}
+ if (sessionSecret)
+ sigKey = KeyPair(sessionSecret);
+
// Two codepaths is necessary since named block require database, but numbered
// blocks are superuseful to have when database is already open in another process.
if (mode == OperationMode::DAGInit && !(initDAG == LatestBlock || initDAG == PendingBlock))
@@ -537,9 +576,7 @@ int main(int argc, char** argv)
killChain,
nodeMode == NodeMode::Full ? set{"eth", "shh"} : set(),
netPrefs,
- &nodesState,
- miners
- );
+ &nodesState);
if (mode == OperationMode::DAGInit)
doInitDAG(web3.ethereum()->blockChain().number() + (initDAG == PendingBlock ? 30000 : 0));
@@ -632,7 +669,7 @@ int main(int argc, char** argv)
c->setAddress(coinbase);
}
- cout << "Transaction Signer: " << us.address() << endl;
+ cout << "Transaction Signer: " << sigKey.address() << endl;
cout << "Mining Benefactor: " << coinbase << endl;
web3.startNetwork();
@@ -647,8 +684,7 @@ int main(int argc, char** argv)
if (jsonrpc > -1)
{
jsonrpcConnector = unique_ptr(new jsonrpc::HttpServer(jsonrpc, "", "", SensibleHttpThreads));
- jsonrpcServer = shared_ptr(new WebThreeStubServer(*jsonrpcConnector.get(), web3, vector({us})));
- jsonrpcServer->setIdentities({us});
+ jsonrpcServer = shared_ptr(new WebThreeStubServer(*jsonrpcConnector.get(), web3, vector({sigKey})));
jsonrpcServer->StartListening();
}
#endif
@@ -772,8 +808,7 @@ int main(int argc, char** argv)
if (jsonrpc < 0)
jsonrpc = SensibleHttpPort;
jsonrpcConnector = unique_ptr(new jsonrpc::HttpServer(jsonrpc, "", "", SensibleHttpThreads));
- jsonrpcServer = shared_ptr(new WebThreeStubServer(*jsonrpcConnector.get(), web3, vector({us})));
- jsonrpcServer->setIdentities({us});
+ jsonrpcServer = shared_ptr(new WebThreeStubServer(*jsonrpcConnector.get(), web3, vector({sigKey})));
jsonrpcServer->StartListening();
}
else if (cmd == "jsonstop")
@@ -785,12 +820,11 @@ int main(int argc, char** argv)
#endif
else if (cmd == "address")
{
- cout << "Current address:" << endl
- << toHex(us.address().asArray()) << endl;
+ cout << "Current address:" << endl << sigKey.address() << endl;
}
else if (cmd == "secret")
{
- cout << "Secret Key: " << toHex(us.secret().asArray()) << endl;
+ cout << "Secret Key: " << sigKey.secret() << endl;
}
else if (c && cmd == "block")
{
@@ -805,7 +839,7 @@ int main(int argc, char** argv)
}
else if (c && cmd == "balance")
{
- cout << "Current balance: " << formatBalance( c->balanceAt(us.address())) << " = " <balanceAt(us.address()) << " wei" << endl;
+ cout << "Current balance: " << formatBalance( c->balanceAt(sigKey.address())) << " = " <balanceAt(sigKey.address()) << " wei" << endl;
}
else if (c && cmd == "transact")
{
@@ -921,7 +955,7 @@ int main(int argc, char** argv)
try
{
Address dest = h160(fromHex(hexAddr, WhenError::Throw));
- c->submitTransaction(us.secret(), amount, dest, bytes(), minGas);
+ c->submitTransaction(sigKey.secret(), amount, dest, bytes(), minGas);
}
catch (BadHexCharacter& _e)
{
@@ -990,7 +1024,7 @@ int main(int argc, char** argv)
else if (gas < minGas)
cwarn << "Minimum gas amount is" << minGas;
else
- c->submitTransaction(us.secret(), endowment, init, gas, gasPrice);
+ c->submitTransaction(sigKey.secret(), endowment, init, gas, gasPrice);
}
else
cwarn << "Require parameters: contract ENDOWMENT GASPRICE GAS CODEHEX";
@@ -1107,7 +1141,7 @@ int main(int argc, char** argv)
{
string hexSec;
iss >> hexSec;
- us = KeyPair(h256(fromHex(hexSec)));
+ sigKey = KeyPair(h256(fromHex(hexSec)));
}
else
cwarn << "Require parameter: setSecret HEXSECRETKEY";
@@ -1147,7 +1181,7 @@ int main(int argc, char** argv)
string path;
iss >> path;
RLPStream config(2);
- config << us.secret() << coinbase;
+ config << sigKey.secret() << coinbase;
writeFile(path, config.out());
}
else
@@ -1163,7 +1197,7 @@ int main(int argc, char** argv)
if (b.size())
{
RLP config(b);
- us = KeyPair(config[0].toHash());
+ sigKey = KeyPair(config[0].toHash());
coinbase = config[1].toHash();
}
else
diff --git a/exp/main.cpp b/exp/main.cpp
index 48562f80e..20f287f43 100644
--- a/exp/main.cpp
+++ b/exp/main.cpp
@@ -25,6 +25,7 @@
#include "libethash-cl/cl.hpp"
#endif
#include
+#include
#include
#include
#include
@@ -34,11 +35,12 @@
#include
#include
#include
-#include
#include
#include
+#include
#include
#include
+#include
#include
#include
#include
@@ -106,21 +108,144 @@ int main()
cnote << "State after transaction: " << s;
cnote << before.diff(s);
}
-#else
+#elif 0
int main()
{
-#if ETH_ETHASHCL
- EthashCL ecl;
+ GenericFarm f;
BlockInfo genesis = CanonBlockChain::genesis();
genesis.difficulty = 1 << 18;
- cdebug << (h256)u256((bigint(1) << 256) / genesis.difficulty);
- std::pair r;
- while (!r.first.completed)
- r = ecl.mine(genesis, 1000);
- cdebug << r.second.mixHash << r.second.nonce;
- EthashCL::assignResult(r.second, genesis);
- assert(EthashCPU::verify(genesis));
-#endif
+ cdebug << genesis.boundary();
+
+ auto mine = [](GenericFarm& f, BlockInfo const& g, unsigned timeout) {
+ BlockInfo bi = g;
+ bool completed = false;
+ f.onSolutionFound([&](ProofOfWork::Solution sol)
+ {
+ ProofOfWork::assignResult(sol, bi);
+ return completed = true;
+ });
+ f.setWork(bi);
+ for (unsigned i = 0; !completed && i < timeout * 10; ++i, cout << f.miningProgress() << "\r" << flush)
+ this_thread::sleep_for(chrono::milliseconds(100));
+ cout << endl << flush;
+ cdebug << bi.mixHash << bi.nonce << (Ethash::verify(bi) ? "GOOD" : "bad");
+ };
+
+ Ethash::prep(genesis);
+
+ genesis.difficulty = u256(1) << 40;
+ genesis.noteDirty();
+ f.startCPU();
+ mine(f, genesis, 10);
+
+ f.startGPU();
+
+ cdebug << "Good:";
+ genesis.difficulty = 1 << 18;
+ genesis.noteDirty();
+ mine(f, genesis, 30);
+
+ cdebug << "Bad:";
+ genesis.difficulty = (u256(1) << 40);
+ genesis.noteDirty();
+ mine(f, genesis, 30);
+
+ f.stop();
+
+ return 0;
+}
+#elif 0
+
+void mine(State& s, BlockChain const& _bc)
+{
+ s.commitToMine(_bc);
+ GenericFarm f;
+ bool completed = false;
+ f.onSolutionFound([&](ProofOfWork::Solution sol)
+ {
+ return completed = s.completeMine(sol);
+ });
+ f.setWork(s.info());
+ f.startCPU();
+ while (!completed)
+ this_thread::sleep_for(chrono::milliseconds(20));
+}
+#elif 0
+int main()
+{
+ cnote << "Testing State...";
+
+ KeyPair me = sha3("Gav Wood");
+ KeyPair myMiner = sha3("Gav's Miner");
+// KeyPair you = sha3("123");
+
+ Defaults::setDBPath(boost::filesystem::temp_directory_path().string() + "/" + toString(chrono::system_clock::now().time_since_epoch().count()));
+
+ OverlayDB stateDB = State::openDB();
+ CanonBlockChain bc;
+ cout << bc;
+
+ State s(stateDB, BaseState::CanonGenesis, myMiner.address());
+ cout << s;
+
+ // Sync up - this won't do much until we use the last state.
+ s.sync(bc);
+
+ cout << s;
+
+ // Mine to get some ether!
+ mine(s, bc);
+
+ bc.attemptImport(s.blockData(), stateDB);
+
+ cout << bc;
+
+ s.sync(bc);
+
+ cout << s;
+
+ // Inject a transaction to transfer funds from miner to me.
+ Transaction t(1000, 10000, 30000, me.address(), bytes(), s.transactionsFrom(myMiner.address()), myMiner.secret());
+ assert(t.sender() == myMiner.address());
+ s.execute(bc.lastHashes(), t);
+
+ cout << s;
+
+ // Mine to get some ether and set in stone.
+ s.commitToMine(bc);
+ s.commitToMine(bc);
+ mine(s, bc);
+ bc.attemptImport(s.blockData(), stateDB);
+
+ cout << bc;
+
+ s.sync(bc);
+
+ cout << s;
+
+ return 0;
+}
+#else
+int main()
+{
+ string tempDir = boost::filesystem::temp_directory_path().string() + "/" + toString(chrono::system_clock::now().time_since_epoch().count());
+
+ KeyPair myMiner = sha3("Gav's Miner");
+
+ p2p::Host net("Test");
+ cdebug << "Path:" << tempDir;
+ Client c(&net, tempDir);
+
+ c.setAddress(myMiner.address());
+
+ this_thread::sleep_for(chrono::milliseconds(1000));
+
+ c.startMining();
+
+ this_thread::sleep_for(chrono::milliseconds(6000));
+
+ c.stopMining();
+
return 0;
}
#endif
diff --git a/libdevcore/Common.cpp b/libdevcore/Common.cpp
index b6e8e7f93..78b3d9c30 100644
--- a/libdevcore/Common.cpp
+++ b/libdevcore/Common.cpp
@@ -27,7 +27,7 @@ using namespace dev;
namespace dev
{
-char const* Version = "0.9.7";
+char const* Version = "0.9.8";
}
diff --git a/libdevcore/CommonData.h b/libdevcore/CommonData.h
index 38ccd71f0..93bad71a3 100644
--- a/libdevcore/CommonData.h
+++ b/libdevcore/CommonData.h
@@ -116,9 +116,9 @@ inline void toBigEndian(_T _val, _Out& o_out)
template
inline _T fromBigEndian(_In const& _bytes)
{
- _T ret = 0;
+ _T ret = (_T)0;
for (auto i: _bytes)
- ret = (ret << 8) | (byte)(typename std::make_unsigned::type)i;
+ ret = (_T)((ret << 8) | (byte)(typename std::make_unsigned::type)i);
return ret;
}
diff --git a/libdevcore/Guards.h b/libdevcore/Guards.h
index f5c64b041..4229428ce 100644
--- a/libdevcore/Guards.h
+++ b/libdevcore/Guards.h
@@ -38,4 +38,75 @@ using UpgradableGuard = boost::upgrade_lock;
using UpgradeGuard = boost::upgrade_to_unique_lock;
using WriteGuard = boost::unique_lock;
+template
+struct GenericGuardBool: GuardType
+{
+ GenericGuardBool(MutexType& _m): GuardType(_m) {}
+ bool b = true;
+};
+template
+struct GenericUnguardBool
+{
+ GenericUnguardBool(MutexType& _m): m(_m) { m.unlock(); }
+ ~GenericUnguardBool() { m.lock(); }
+ bool b = true;
+ MutexType& m;
+};
+template
+struct GenericUnguardSharedBool
+{
+ GenericUnguardSharedBool(MutexType& _m): m(_m) { m.unlock_shared(); }
+ ~GenericUnguardSharedBool() { m.lock_shared(); }
+ bool b = true;
+ MutexType& m;
+};
+
+/** @brief Simple block guard.
+ * The expression/block following is guarded though the given mutex.
+ * Usage:
+ * @code
+ * Mutex m;
+ * unsigned d;
+ * ...
+ * ETH_GUARDED(m) d = 1;
+ * ...
+ * ETH_GUARDED(m) { for (auto d = 10; d > 0; --d) foo(d); d = 0; }
+ * @endcode
+ *
+ * There are several variants of this basic mechanism for different Mutex types and Guards.
+ *
+ * There is also the UNGUARD variant which allows an unguarded expression/block to exist within a
+ * guarded expression. eg:
+ *
+ * @code
+ * Mutex m;
+ * int d;
+ * ...
+ * ETH_GUARDED(m)
+ * {
+ * for (auto d = 50; d > 25; --d)
+ * foo(d);
+ * ETH_UNGUARDED(m)
+ * bar();
+ * for (; d > 0; --d)
+ * foo(d);
+ * }
+ * @endcode
+ */
+
+#define ETH_GUARDED(MUTEX) \
+ for (GenericGuardBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+#define ETH_READ_GUARDED(MUTEX) \
+ for (GenericGuardBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+#define ETH_WRITE_GUARDED(MUTEX) \
+ for (GenericGuardBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+#define ETH_RECURSIVE_GUARDED(MUTEX) \
+ for (GenericGuardBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+#define ETH_UNGUARDED(MUTEX) \
+ for (GenericUnguardBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+#define ETH_READ_UNGUARDED(MUTEX) \
+ for (GenericUnguardSharedBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+#define ETH_WRITE_UNGUARDED(MUTEX) \
+ for (GenericUnguardBool __eth_l(MUTEX); __eth_l.b; __eth_l.b = false)
+
}
diff --git a/libdevcore/Worker.cpp b/libdevcore/Worker.cpp
index 175323620..8c1fbb9c7 100644
--- a/libdevcore/Worker.cpp
+++ b/libdevcore/Worker.cpp
@@ -27,24 +27,28 @@
using namespace std;
using namespace dev;
-void Worker::startWorking()
+void Worker::startWorking(IfRunning _ir)
{
- cnote << "startWorking for thread" << m_name;
+// cnote << "startWorking for thread" << m_name;
Guard l(x_work);
- if (m_work)
- return;
+
+ if (m_work && m_work->joinable())
+ try {
+ if (_ir == IfRunning::Detach)
+ m_work->detach();
+ else if (_ir == IfRunning::Join)
+ m_work->join();
+ else
+ return;
+ } catch (...) {}
cnote << "Spawning" << m_name;
m_stop = false;
m_work.reset(new thread([&]()
{
setThreadName(m_name.c_str());
startedWorking();
- while (!m_stop)
- {
- if (m_idleWaitMs)
- this_thread::sleep_for(chrono::milliseconds(m_idleWaitMs));
- doWork();
- }
+ workLoop();
+ m_work->detach();
cnote << "Finishing up worker thread";
doneWorking();
}));
@@ -52,14 +56,26 @@ void Worker::startWorking()
void Worker::stopWorking()
{
- cnote << "stopWorking for thread" << m_name;
+// cnote << "stopWorking for thread" << m_name;
Guard l(x_work);
- if (!m_work)
+ if (!m_work || !m_work->joinable())
return;
cnote << "Stopping" << m_name;
m_stop = true;
- m_work->join();
+ try {
+ m_work->join();
+ }
+ catch (...) {}
m_work.reset();
cnote << "Stopped" << m_name;
}
+void Worker::workLoop()
+{
+ while (!m_stop)
+ {
+ if (m_idleWaitMs)
+ this_thread::sleep_for(chrono::milliseconds(m_idleWaitMs));
+ doWork();
+ }
+}
diff --git a/libdevcore/Worker.h b/libdevcore/Worker.h
index 40bc118aa..287ff6d6f 100644
--- a/libdevcore/Worker.h
+++ b/libdevcore/Worker.h
@@ -23,11 +23,19 @@
#include
#include
+#include
#include "Guards.h"
namespace dev
{
+enum class IfRunning
+{
+ Fail,
+ Join,
+ Detach
+};
+
class Worker
{
protected:
@@ -45,7 +53,7 @@ protected:
void setName(std::string _n) { if (!isWorking()) m_name = _n; }
/// Starts worker thread; causes startedWorking() to be called.
- void startWorking();
+ void startWorking(IfRunning _ir = IfRunning::Fail);
/// Stop worker thread; causes call to stopWorking().
void stopWorking();
@@ -57,11 +65,18 @@ protected:
virtual void startedWorking() {}
/// Called continuously following sleep for m_idleWaitMs.
- virtual void doWork() = 0;
+ virtual void doWork() {}
+
+ /// Overrides doWork(); should call shouldStop() often and exit when true.
+ virtual void workLoop();
+ bool shouldStop() const { return m_stop; }
/// Called when is to be stopped, just prior to thread being joined.
virtual void doneWorking() {}
+ /// Blocks caller into worker thread has finished.
+ void join() const { Guard l(x_work); try { if (m_work) m_work->join(); } catch (...) {} }
+
private:
std::string m_name;
unsigned m_idleWaitMs = 0;
diff --git a/libethash-cl/CMakeLists.txt b/libethash-cl/CMakeLists.txt
new file mode 100644
index 000000000..7b00a22bd
--- /dev/null
+++ b/libethash-cl/CMakeLists.txt
@@ -0,0 +1,47 @@
+cmake_minimum_required(VERSION 2.8)
+
+set(LIBRARY ethash-cl)
+#set(CMAKE_BUILD_TYPE Release)
+
+include(bin2h.cmake)
+bin2h(SOURCE_FILE ethash_cl_miner_kernel.cl
+ VARIABLE_NAME ethash_cl_miner_kernel
+ HEADER_FILE ${CMAKE_CURRENT_BINARY_DIR}/ethash_cl_miner_kernel.h)
+
+if (NOT MSVC)
+ # Initialize CXXFLAGS for c++11
+ set(CMAKE_CXX_FLAGS "-Wall -std=c++11")
+ set(CMAKE_CXX_FLAGS_DEBUG "-O0 -g")
+ set(CMAKE_CXX_FLAGS_MINSIZEREL "-Os -DNDEBUG")
+ set(CMAKE_CXX_FLAGS_RELEASE "-O4 -DNDEBUG")
+ set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g")
+
+ # Compiler-specific C++11 activation.
+ if ("${CMAKE_CXX_COMPILER_ID}" MATCHES "GNU")
+ execute_process(
+ COMMAND ${CMAKE_CXX_COMPILER} -dumpversion OUTPUT_VARIABLE GCC_VERSION)
+ if (NOT (GCC_VERSION VERSION_GREATER 4.7 OR GCC_VERSION VERSION_EQUAL 4.7))
+ message(FATAL_ERROR "${PROJECT_NAME} requires g++ 4.7 or greater.")
+ endif ()
+ elseif ("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -stdlib=libc++")
+ else ()
+ message(FATAL_ERROR "Your C++ compiler does not support C++11.")
+ endif ()
+endif()
+
+set(OpenCL_FOUND TRUE)
+set(OpenCL_INCLUDE_DIRS /usr/include/CL)
+set(OpenCL_LIBRARIES -lOpenCL)
+
+if (NOT OpenCL_FOUND)
+ find_package(OpenCL)
+endif()
+
+if (OpenCL_FOUND)
+ set(CMAKE_CXX_FLAGS "-std=c++11 -Wall -Wno-unknown-pragmas -Wextra -Werror -pedantic -fPIC ${CMAKE_CXX_FLAGS}")
+ include_directories(${OpenCL_INCLUDE_DIRS} ${CMAKE_CURRENT_BINARY_DIR})
+ include_directories(..)
+ add_library(${LIBRARY} ethash_cl_miner.cpp ethash_cl_miner.h cl.hpp)
+ TARGET_LINK_LIBRARIES(${LIBRARY} ${OpenCL_LIBRARIES} ethash)
+endif()
diff --git a/libethash-cl/bin2h.cmake b/libethash-cl/bin2h.cmake
new file mode 100644
index 000000000..90ca9cc5b
--- /dev/null
+++ b/libethash-cl/bin2h.cmake
@@ -0,0 +1,86 @@
+# https://gist.github.com/sivachandran/3a0de157dccef822a230
+include(CMakeParseArguments)
+
+# Function to wrap a given string into multiple lines at the given column position.
+# Parameters:
+# VARIABLE - The name of the CMake variable holding the string.
+# AT_COLUMN - The column position at which string will be wrapped.
+function(WRAP_STRING)
+ set(oneValueArgs VARIABLE AT_COLUMN)
+ cmake_parse_arguments(WRAP_STRING "${options}" "${oneValueArgs}" "" ${ARGN})
+
+ string(LENGTH ${${WRAP_STRING_VARIABLE}} stringLength)
+ math(EXPR offset "0")
+
+ while(stringLength GREATER 0)
+
+ if(stringLength GREATER ${WRAP_STRING_AT_COLUMN})
+ math(EXPR length "${WRAP_STRING_AT_COLUMN}")
+ else()
+ math(EXPR length "${stringLength}")
+ endif()
+
+ string(SUBSTRING ${${WRAP_STRING_VARIABLE}} ${offset} ${length} line)
+ set(lines "${lines}\n${line}")
+
+ math(EXPR stringLength "${stringLength} - ${length}")
+ math(EXPR offset "${offset} + ${length}")
+ endwhile()
+
+ set(${WRAP_STRING_VARIABLE} "${lines}" PARENT_SCOPE)
+endfunction()
+
+# Function to embed contents of a file as byte array in C/C++ header file(.h). The header file
+# will contain a byte array and integer variable holding the size of the array.
+# Parameters
+# SOURCE_FILE - The path of source file whose contents will be embedded in the header file.
+# VARIABLE_NAME - The name of the variable for the byte array. The string "_SIZE" will be append
+# to this name and will be used a variable name for size variable.
+# HEADER_FILE - The path of header file.
+# APPEND - If specified appends to the header file instead of overwriting it
+# NULL_TERMINATE - If specified a null byte(zero) will be append to the byte array. This will be
+# useful if the source file is a text file and we want to use the file contents
+# as string. But the size variable holds size of the byte array without this
+# null byte.
+# Usage:
+# bin2h(SOURCE_FILE "Logo.png" HEADER_FILE "Logo.h" VARIABLE_NAME "LOGO_PNG")
+function(BIN2H)
+ set(options APPEND NULL_TERMINATE)
+ set(oneValueArgs SOURCE_FILE VARIABLE_NAME HEADER_FILE)
+ cmake_parse_arguments(BIN2H "${options}" "${oneValueArgs}" "" ${ARGN})
+
+ # reads source file contents as hex string
+ file(READ ${BIN2H_SOURCE_FILE} hexString HEX)
+ string(LENGTH ${hexString} hexStringLength)
+
+ # appends null byte if asked
+ if(BIN2H_NULL_TERMINATE)
+ set(hexString "${hexString}00")
+ endif()
+
+ # wraps the hex string into multiple lines at column 32(i.e. 16 bytes per line)
+ wrap_string(VARIABLE hexString AT_COLUMN 32)
+ math(EXPR arraySize "${hexStringLength} / 2")
+
+ # adds '0x' prefix and comma suffix before and after every byte respectively
+ string(REGEX REPLACE "([0-9a-f][0-9a-f])" "0x\\1, " arrayValues ${hexString})
+ # removes trailing comma
+ string(REGEX REPLACE ", $" "" arrayValues ${arrayValues})
+
+ # converts the variable name into proper C identifier
+ IF (${CMAKE_VERSION} GREATER 2.8.10) # fix for legacy cmake
+ string(MAKE_C_IDENTIFIER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME)
+ ENDIF()
+ string(TOUPPER "${BIN2H_VARIABLE_NAME}" BIN2H_VARIABLE_NAME)
+
+ # declares byte array and the length variables
+ set(arrayDefinition "const unsigned char ${BIN2H_VARIABLE_NAME}[] = { ${arrayValues} };")
+ set(arraySizeDefinition "const size_t ${BIN2H_VARIABLE_NAME}_SIZE = ${arraySize};")
+
+ set(declarations "${arrayDefinition}\n\n${arraySizeDefinition}\n\n")
+ if(BIN2H_APPEND)
+ file(APPEND ${BIN2H_HEADER_FILE} "${declarations}")
+ else()
+ file(WRITE ${BIN2H_HEADER_FILE} "${declarations}")
+ endif()
+endfunction()
diff --git a/libethash-cl/cl.hpp b/libethash-cl/cl.hpp
new file mode 100644
index 000000000..a38498762
--- /dev/null
+++ b/libethash-cl/cl.hpp
@@ -0,0 +1,4014 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2010 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
+
+/*! \file
+ *
+ * \brief C++ bindings for OpenCL 1.0 (rev 48) and OpenCL 1.1 (rev 33)
+ * \author Benedict R. Gaster and Laurent Morichetti
+ *
+ * Additions and fixes from Brian Cole, March 3rd 2010.
+ *
+ * \version 1.1
+ * \date June 2010
+ *
+ * Optional extension support
+ *
+ * cl
+ * cl_ext_device_fission
+ * #define USE_CL_DEVICE_FISSION
+ */
+
+/*! \mainpage
+ * \section intro Introduction
+ * For many large applications C++ is the language of choice and so it seems
+ * reasonable to define C++ bindings for OpenCL.
+ *
+ *
+ * The interface is contained with a single C++ header file \em cl.hpp and all
+ * definitions are contained within the namespace \em cl. There is no additional
+ * requirement to include \em cl.h and to use either the C++ or original C
+ * bindings it is enough to simply include \em cl.hpp.
+ *
+ * The bindings themselves are lightweight and correspond closely to the
+ * underlying C API. Using the C++ bindings introduces no additional execution
+ * overhead.
+ *
+ * For detail documentation on the bindings see:
+ *
+ * The OpenCL C++ Wrapper API 1.1 (revision 04)
+ * http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
+ *
+ * \section example Example
+ *
+ * The following example shows a general use case for the C++
+ * bindings, including support for the optional exception feature and
+ * also the supplied vector and string classes, see following sections for
+ * decriptions of these features.
+ *
+ * \code
+ * #define __CL_ENABLE_EXCEPTIONS
+ *
+ * #if defined(__APPLE__) || defined(__MACOSX)
+ * #include
+ * #else
+ * #include
+ * #endif
+ * #include
+ * #include
+ * #include
+ *
+ * const char * helloStr = "__kernel void "
+ * "hello(void) "
+ * "{ "
+ * " "
+ * "} ";
+ *
+ * int
+ * main(void)
+ * {
+ * cl_int err = CL_SUCCESS;
+ * try {
+ *
+ * std::vector platforms;
+ * cl::Platform::get(&platforms);
+ * if (platforms.size() == 0) {
+ * std::cout << "Platform size 0\n";
+ * return -1;
+ * }
+ *
+ * cl_context_properties properties[] =
+ * { CL_CONTEXT_PLATFORM, (cl_context_properties)(platforms[0])(), 0};
+ * cl::Context context(CL_DEVICE_TYPE_CPU, properties);
+ *
+ * std::vector devices = context.getInfo();
+ *
+ * cl::Program::Sources source(1,
+ * std::make_pair(helloStr,strlen(helloStr)));
+ * cl::Program program_ = cl::Program(context, source);
+ * program_.build(devices);
+ *
+ * cl::Kernel kernel(program_, "hello", &err);
+ *
+ * cl::Event event;
+ * cl::CommandQueue queue(context, devices[0], 0, &err);
+ * queue.enqueueNDRangeKernel(
+ * kernel,
+ * cl::NullRange,
+ * cl::NDRange(4,4),
+ * cl::NullRange,
+ * NULL,
+ * &event);
+ *
+ * event.wait();
+ * }
+ * catch (cl::Error err) {
+ * std::cerr
+ * << "ERROR: "
+ * << err.what()
+ * << "("
+ * << err.err()
+ * << ")"
+ * << std::endl;
+ * }
+ *
+ * return EXIT_SUCCESS;
+ * }
+ *
+ * \endcode
+ *
+ */
+#ifndef CL_HPP_
+#define CL_HPP_
+
+#ifdef _WIN32
+#include
+#include
+#if defined(USE_DX_INTEROP)
+#include
+#endif
+#endif // _WIN32
+
+//
+#if defined(USE_CL_DEVICE_FISSION)
+#include
+#endif
+
+#if defined(__APPLE__) || defined(__MACOSX)
+#include
+#include
+#else
+#include
+#include
+#endif // !__APPLE__
+
+#if !defined(CL_CALLBACK)
+#define CL_CALLBACK
+#endif //CL_CALLBACK
+
+#include
+
+#if !defined(__NO_STD_VECTOR)
+#include
+#endif
+
+#if !defined(__NO_STD_STRING)
+#include
+#endif
+
+#if defined(linux) || defined(__APPLE__) || defined(__MACOSX)
+# include
+#endif // linux
+
+#include
+
+/*! \namespace cl
+ *
+ * \brief The OpenCL C++ bindings are defined within this namespace.
+ *
+ */
+namespace cl {
+
+#define __INIT_CL_EXT_FCN_PTR(name) \
+ if(!pfn_##name) { \
+ pfn_##name = (PFN_##name) \
+ clGetExtensionFunctionAddress(#name); \
+ if(!pfn_##name) { \
+ } \
+ }
+
+class Program;
+class Device;
+class Context;
+class CommandQueue;
+class Memory;
+
+#if defined(__CL_ENABLE_EXCEPTIONS)
+#include
+/*! \class Error
+ * \brief Exception class
+ */
+class Error : public std::exception
+{
+private:
+ cl_int err_;
+ const char * errStr_;
+public:
+ /*! Create a new CL error exception for a given error code
+ * and corresponding message.
+ */
+ Error(cl_int err, const char * errStr = NULL) : err_(err), errStr_(errStr)
+ {}
+
+ ~Error() throw() {}
+
+ /*! \brief Get error string associated with exception
+ *
+ * \return A memory pointer to the error message string.
+ */
+ virtual const char * what() const throw ()
+ {
+ if (errStr_ == NULL) {
+ return "empty";
+ }
+ else {
+ return errStr_;
+ }
+ }
+
+ /*! \brief Get error code associated with exception
+ *
+ * \return The error code.
+ */
+ cl_int err(void) const { return err_; }
+};
+
+#define __ERR_STR(x) #x
+#else
+#define __ERR_STR(x) NULL
+#endif // __CL_ENABLE_EXCEPTIONS
+
+//! \cond DOXYGEN_DETAIL
+#if !defined(__CL_USER_OVERRIDE_ERROR_STRINGS)
+#define __GET_DEVICE_INFO_ERR __ERR_STR(clgetDeviceInfo)
+#define __GET_PLATFORM_INFO_ERR __ERR_STR(clGetPlatformInfo)
+#define __GET_DEVICE_IDS_ERR __ERR_STR(clGetDeviceIDs)
+#define __GET_PLATFORM_IDS_ERR __ERR_STR(clGetPlatformIDs)
+#define __GET_CONTEXT_INFO_ERR __ERR_STR(clGetContextInfo)
+#define __GET_EVENT_INFO_ERR __ERR_STR(clGetEventInfo)
+#define __GET_EVENT_PROFILE_INFO_ERR __ERR_STR(clGetEventProfileInfo)
+#define __GET_MEM_OBJECT_INFO_ERR __ERR_STR(clGetMemObjectInfo)
+#define __GET_IMAGE_INFO_ERR __ERR_STR(clGetImageInfo)
+#define __GET_SAMPLER_INFO_ERR __ERR_STR(clGetSamplerInfo)
+#define __GET_KERNEL_INFO_ERR __ERR_STR(clGetKernelInfo)
+#define __GET_KERNEL_WORK_GROUP_INFO_ERR __ERR_STR(clGetKernelWorkGroupInfo)
+#define __GET_PROGRAM_INFO_ERR __ERR_STR(clGetProgramInfo)
+#define __GET_PROGRAM_BUILD_INFO_ERR __ERR_STR(clGetProgramBuildInfo)
+#define __GET_COMMAND_QUEUE_INFO_ERR __ERR_STR(clGetCommandQueueInfo)
+
+#define __CREATE_CONTEXT_FROM_TYPE_ERR __ERR_STR(clCreateContextFromType)
+#define __GET_SUPPORTED_IMAGE_FORMATS_ERR __ERR_STR(clGetSupportedImageFormats)
+
+#define __CREATE_BUFFER_ERR __ERR_STR(clCreateBuffer)
+#define __CREATE_SUBBUFFER_ERR __ERR_STR(clCreateSubBuffer)
+#define __CREATE_GL_BUFFER_ERR __ERR_STR(clCreateFromGLBuffer)
+#define __GET_GL_OBJECT_INFO_ERR __ERR_STR(clGetGLObjectInfo)
+#define __CREATE_IMAGE2D_ERR __ERR_STR(clCreateImage2D)
+#define __CREATE_IMAGE3D_ERR __ERR_STR(clCreateImage3D)
+#define __CREATE_SAMPLER_ERR __ERR_STR(clCreateSampler)
+#define __SET_MEM_OBJECT_DESTRUCTOR_CALLBACK_ERR __ERR_STR(clSetMemObjectDestructorCallback)
+
+#define __CREATE_USER_EVENT_ERR __ERR_STR(clCreateUserEvent)
+#define __SET_USER_EVENT_STATUS_ERR __ERR_STR(clSetUserEventStatus)
+#define __SET_EVENT_CALLBACK_ERR __ERR_STR(clSetEventCallback)
+#define __WAIT_FOR_EVENTS_ERR __ERR_STR(clWaitForEvents)
+
+#define __CREATE_KERNEL_ERR __ERR_STR(clCreateKernel)
+#define __SET_KERNEL_ARGS_ERR __ERR_STR(clSetKernelArg)
+#define __CREATE_PROGRAM_WITH_SOURCE_ERR __ERR_STR(clCreateProgramWithSource)
+#define __CREATE_PROGRAM_WITH_BINARY_ERR __ERR_STR(clCreateProgramWithBinary)
+#define __BUILD_PROGRAM_ERR __ERR_STR(clBuildProgram)
+#define __CREATE_KERNELS_IN_PROGRAM_ERR __ERR_STR(clCreateKernelsInProgram)
+
+#define __CREATE_COMMAND_QUEUE_ERR __ERR_STR(clCreateCommandQueue)
+#define __SET_COMMAND_QUEUE_PROPERTY_ERR __ERR_STR(clSetCommandQueueProperty)
+#define __ENQUEUE_READ_BUFFER_ERR __ERR_STR(clEnqueueReadBuffer)
+#define __ENQUEUE_READ_BUFFER_RECT_ERR __ERR_STR(clEnqueueReadBufferRect)
+#define __ENQUEUE_WRITE_BUFFER_ERR __ERR_STR(clEnqueueWriteBuffer)
+#define __ENQUEUE_WRITE_BUFFER_RECT_ERR __ERR_STR(clEnqueueWriteBufferRect)
+#define __ENQEUE_COPY_BUFFER_ERR __ERR_STR(clEnqueueCopyBuffer)
+#define __ENQEUE_COPY_BUFFER_RECT_ERR __ERR_STR(clEnqueueCopyBufferRect)
+#define __ENQUEUE_READ_IMAGE_ERR __ERR_STR(clEnqueueReadImage)
+#define __ENQUEUE_WRITE_IMAGE_ERR __ERR_STR(clEnqueueWriteImage)
+#define __ENQUEUE_COPY_IMAGE_ERR __ERR_STR(clEnqueueCopyImage)
+#define __ENQUEUE_COPY_IMAGE_TO_BUFFER_ERR __ERR_STR(clEnqueueCopyImageToBuffer)
+#define __ENQUEUE_COPY_BUFFER_TO_IMAGE_ERR __ERR_STR(clEnqueueCopyBufferToImage)
+#define __ENQUEUE_MAP_BUFFER_ERR __ERR_STR(clEnqueueMapBuffer)
+#define __ENQUEUE_MAP_IMAGE_ERR __ERR_STR(clEnqueueMapImage)
+#define __ENQUEUE_UNMAP_MEM_OBJECT_ERR __ERR_STR(clEnqueueUnMapMemObject)
+#define __ENQUEUE_NDRANGE_KERNEL_ERR __ERR_STR(clEnqueueNDRangeKernel)
+#define __ENQUEUE_TASK_ERR __ERR_STR(clEnqueueTask)
+#define __ENQUEUE_NATIVE_KERNEL __ERR_STR(clEnqueueNativeKernel)
+#define __ENQUEUE_MARKER_ERR __ERR_STR(clEnqueueMarker)
+#define __ENQUEUE_WAIT_FOR_EVENTS_ERR __ERR_STR(clEnqueueWaitForEvents)
+#define __ENQUEUE_BARRIER_ERR __ERR_STR(clEnqueueBarrier)
+
+#define __ENQUEUE_ACQUIRE_GL_ERR __ERR_STR(clEnqueueAcquireGLObjects)
+#define __ENQUEUE_RELEASE_GL_ERR __ERR_STR(clEnqueueReleaseGLObjects)
+
+#define __UNLOAD_COMPILER_ERR __ERR_STR(clUnloadCompiler)
+
+#define __FLUSH_ERR __ERR_STR(clFlush)
+#define __FINISH_ERR __ERR_STR(clFinish)
+
+#define __CREATE_SUB_DEVICES __ERR_STR(clCreateSubDevicesEXT)
+#endif // __CL_USER_OVERRIDE_ERROR_STRINGS
+//! \endcond
+
+/*! \class string
+ * \brief Simple string class, that provides a limited subset of std::string
+ * functionality but avoids many of the issues that come with that class.
+ */
+class string
+{
+private:
+ ::size_t size_;
+ char * str_;
+public:
+ string(void) : size_(0), str_(NULL)
+ {
+ }
+
+ string(char * str, ::size_t size) :
+ size_(size),
+ str_(NULL)
+ {
+ str_ = new char[size_+1];
+ if (str_ != NULL) {
+ memcpy(str_, str, size_ * sizeof(char));
+ str_[size_] = '\0';
+ }
+ else {
+ size_ = 0;
+ }
+ }
+
+ string(char * str) :
+ str_(NULL)
+ {
+ size_= ::strlen(str);
+ str_ = new char[size_ + 1];
+ if (str_ != NULL) {
+ memcpy(str_, str, (size_ + 1) * sizeof(char));
+ }
+ else {
+ size_ = 0;
+ }
+ }
+
+ string& operator=(const string& rhs)
+ {
+ if (this == &rhs) {
+ return *this;
+ }
+
+ if (rhs.size_ == 0 || rhs.str_ == NULL) {
+ size_ = 0;
+ str_ = NULL;
+ }
+ else {
+ size_ = rhs.size_;
+ str_ = new char[size_ + 1];
+ if (str_ != NULL) {
+ memcpy(str_, rhs.str_, (size_ + 1) * sizeof(char));
+ }
+ else {
+ size_ = 0;
+ }
+ }
+
+ return *this;
+ }
+
+ string(const string& rhs)
+ {
+ *this = rhs;
+ }
+
+ ~string()
+ {
+ if (str_ != NULL) {
+ delete[] str_;
+ }
+ }
+
+ ::size_t size(void) const { return size_; }
+ ::size_t length(void) const { return size(); }
+
+ const char * c_str(void) const { return (str_) ? str_ : "";}
+};
+
+#if !defined(__USE_DEV_STRING) && !defined(__NO_STD_STRING)
+#include
+typedef std::string STRING_CLASS;
+#elif !defined(__USE_DEV_STRING)
+typedef cl::string STRING_CLASS;
+#endif
+
+#if !defined(__USE_DEV_VECTOR) && !defined(__NO_STD_VECTOR)
+#include
+#define VECTOR_CLASS std::vector
+#elif !defined(__USE_DEV_VECTOR)
+#define VECTOR_CLASS cl::vector
+#endif
+
+#if !defined(__MAX_DEFAULT_VECTOR_SIZE)
+#define __MAX_DEFAULT_VECTOR_SIZE 10
+#endif
+
+/*! \class vector
+ * \brief Fixed sized vector implementation that mirroring
+ * std::vector functionality.
+ */
+template
+class vector
+{
+private:
+ T data_[N];
+ unsigned int size_;
+ bool empty_;
+public:
+ vector() :
+ size_(-1),
+ empty_(true)
+ {}
+
+ ~vector() {}
+
+ unsigned int size(void) const
+ {
+ return size_ + 1;
+ }
+
+ void clear()
+ {
+ size_ = -1;
+ empty_ = true;
+ }
+
+ void push_back (const T& x)
+ {
+ if (size() < N) {
+ size_++;
+ data_[size_] = x;
+ empty_ = false;
+ }
+ }
+
+ void pop_back(void)
+ {
+ if (!empty_) {
+ data_[size_].~T();
+ size_--;
+ if (size_ == -1) {
+ empty_ = true;
+ }
+ }
+ }
+
+ vector(const vector& vec) :
+ size_(vec.size_),
+ empty_(vec.empty_)
+ {
+ if (!empty_) {
+ memcpy(&data_[0], &vec.data_[0], size() * sizeof(T));
+ }
+ }
+
+ vector(unsigned int size, const T& val = T()) :
+ size_(-1),
+ empty_(true)
+ {
+ for (unsigned int i = 0; i < size; i++) {
+ push_back(val);
+ }
+ }
+
+ vector& operator=(const vector& rhs)
+ {
+ if (this == &rhs) {
+ return *this;
+ }
+
+ size_ = rhs.size_;
+ empty_ = rhs.empty_;
+
+ if (!empty_) {
+ memcpy(&data_[0], &rhs.data_[0], size() * sizeof(T));
+ }
+
+ return *this;
+ }
+
+ bool operator==(vector &vec)
+ {
+ if (empty_ && vec.empty_) {
+ return true;
+ }
+
+ if (size() != vec.size()) {
+ return false;
+ }
+
+ return memcmp(&data_[0], &vec.data_[0], size() * sizeof(T)) == 0 ? true : false;
+ }
+
+ operator T* () { return data_; }
+ operator const T* () const { return data_; }
+
+ bool empty (void) const
+ {
+ return empty_;
+ }
+
+ unsigned int max_size (void) const
+ {
+ return N;
+ }
+
+ unsigned int capacity () const
+ {
+ return sizeof(T) * N;
+ }
+
+ T& operator[](int index)
+ {
+ return data_[index];
+ }
+
+ T operator[](int index) const
+ {
+ return data_[index];
+ }
+
+ template
+ void assign(I start, I end)
+ {
+ clear();
+ while(start < end) {
+ push_back(*start);
+ start++;
+ }
+ }
+
+ /*! \class iterator
+ * \brief Iterator class for vectors
+ */
+ class iterator
+ {
+ private:
+ vector vec_;
+ int index_;
+ bool initialized_;
+ public:
+ iterator(void) :
+ index_(-1),
+ initialized_(false)
+ {
+ index_ = -1;
+ initialized_ = false;
+ }
+
+ ~iterator(void) {}
+
+ static iterator begin(vector &vec)
+ {
+ iterator i;
+
+ if (!vec.empty()) {
+ i.index_ = 0;
+ }
+
+ i.vec_ = vec;
+ i.initialized_ = true;
+ return i;
+ }
+
+ static iterator end(vector &vec)
+ {
+ iterator i;
+
+ if (!vec.empty()) {
+ i.index_ = vec.size();
+ }
+ i.vec_ = vec;
+ i.initialized_ = true;
+ return i;
+ }
+
+ bool operator==(iterator i)
+ {
+ return ((vec_ == i.vec_) &&
+ (index_ == i.index_) &&
+ (initialized_ == i.initialized_));
+ }
+
+ bool operator!=(iterator i)
+ {
+ return (!(*this==i));
+ }
+
+ void operator++()
+ {
+ index_++;
+ }
+
+ void operator++(int x)
+ {
+ index_ += x;
+ }
+
+ void operator--()
+ {
+ index_--;
+ }
+
+ void operator--(int x)
+ {
+ index_ -= x;
+ }
+
+ T operator *()
+ {
+ return vec_[index_];
+ }
+ };
+
+ iterator begin(void)
+ {
+ return iterator::begin(*this);
+ }
+
+ iterator end(void)
+ {
+ return iterator::end(*this);
+ }
+
+ T& front(void)
+ {
+ return data_[0];
+ }
+
+ T& back(void)
+ {
+ return data_[size_];
+ }
+
+ const T& front(void) const
+ {
+ return data_[0];
+ }
+
+ const T& back(void) const
+ {
+ return data_[size_];
+ }
+};
+
+/*!
+ * \brief size_t class used to interface between C++ and
+ * OpenCL C calls that require arrays of size_t values, who's
+ * size is known statically.
+ */
+template
+struct size_t : public cl::vector< ::size_t, N> { };
+
+namespace detail {
+
+// GetInfo help struct
+template
+struct GetInfoHelper
+{
+ static cl_int
+ get(Functor f, cl_uint name, T* param)
+ {
+ return f(name, sizeof(T), param, NULL);
+ }
+};
+
+// Specialized GetInfoHelper for VECTOR_CLASS params
+template
+struct GetInfoHelper >
+{
+ static cl_int get(Func f, cl_uint name, VECTOR_CLASS* param)
+ {
+ ::size_t required;
+ cl_int err = f(name, 0, NULL, &required);
+ if (err != CL_SUCCESS) {
+ return err;
+ }
+
+ T* value = (T*) alloca(required);
+ err = f(name, required, value, NULL);
+ if (err != CL_SUCCESS) {
+ return err;
+ }
+
+ param->assign(&value[0], &value[required/sizeof(T)]);
+ return CL_SUCCESS;
+ }
+};
+
+// Specialized for getInfo
+template
+struct GetInfoHelper >
+{
+ static cl_int
+ get(Func f, cl_uint name, VECTOR_CLASS* param)
+ {
+ cl_uint err = f(name, param->size() * sizeof(char *), &(*param)[0], NULL);
+ if (err != CL_SUCCESS) {
+ return err;
+ }
+
+ return CL_SUCCESS;
+ }
+};
+
+// Specialized GetInfoHelper for STRING_CLASS params
+template
+struct GetInfoHelper
+{
+ static cl_int get(Func f, cl_uint name, STRING_CLASS* param)
+ {
+ ::size_t required;
+ cl_int err = f(name, 0, NULL, &required);
+ if (err != CL_SUCCESS) {
+ return err;
+ }
+
+ char* value = (char*) alloca(required);
+ err = f(name, required, value, NULL);
+ if (err != CL_SUCCESS) {
+ return err;
+ }
+
+ *param = value;
+ return CL_SUCCESS;
+ }
+};
+
+#define __GET_INFO_HELPER_WITH_RETAIN(CPP_TYPE) \
+namespace detail { \
+template \
+struct GetInfoHelper \
+{ \
+ static cl_int get(Func f, cl_uint name, CPP_TYPE* param) \
+ { \
+ cl_uint err = f(name, sizeof(CPP_TYPE), param, NULL); \
+ if (err != CL_SUCCESS) { \
+ return err; \
+ } \
+ \
+ return ReferenceHandler::retain((*param)()); \
+ } \
+}; \
+}
+
+
+#define __PARAM_NAME_INFO_1_0(F) \
+ F(cl_platform_info, CL_PLATFORM_PROFILE, STRING_CLASS) \
+ F(cl_platform_info, CL_PLATFORM_VERSION, STRING_CLASS) \
+ F(cl_platform_info, CL_PLATFORM_NAME, STRING_CLASS) \
+ F(cl_platform_info, CL_PLATFORM_VENDOR, STRING_CLASS) \
+ F(cl_platform_info, CL_PLATFORM_EXTENSIONS, STRING_CLASS) \
+ \
+ F(cl_device_info, CL_DEVICE_TYPE, cl_device_type) \
+ F(cl_device_info, CL_DEVICE_VENDOR_ID, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_COMPUTE_UNITS, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_WORK_GROUP_SIZE, ::size_t) \
+ F(cl_device_info, CL_DEVICE_MAX_WORK_ITEM_SIZES, VECTOR_CLASS< ::size_t>) \
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, cl_uint) \
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, cl_uint) \
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_CLOCK_FREQUENCY, cl_uint) \
+ F(cl_device_info, CL_DEVICE_ADDRESS_BITS, cl_bitfield) \
+ F(cl_device_info, CL_DEVICE_MAX_READ_IMAGE_ARGS, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_WRITE_IMAGE_ARGS, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_MEM_ALLOC_SIZE, cl_ulong) \
+ F(cl_device_info, CL_DEVICE_IMAGE2D_MAX_WIDTH, ::size_t) \
+ F(cl_device_info, CL_DEVICE_IMAGE2D_MAX_HEIGHT, ::size_t) \
+ F(cl_device_info, CL_DEVICE_IMAGE3D_MAX_WIDTH, ::size_t) \
+ F(cl_device_info, CL_DEVICE_IMAGE3D_MAX_HEIGHT, ::size_t) \
+ F(cl_device_info, CL_DEVICE_IMAGE3D_MAX_DEPTH, ::size_t) \
+ F(cl_device_info, CL_DEVICE_IMAGE_SUPPORT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MAX_PARAMETER_SIZE, ::size_t) \
+ F(cl_device_info, CL_DEVICE_MAX_SAMPLERS, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MEM_BASE_ADDR_ALIGN, cl_uint) \
+ F(cl_device_info, CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, cl_uint) \
+ F(cl_device_info, CL_DEVICE_SINGLE_FP_CONFIG, cl_device_fp_config) \
+ F(cl_device_info, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, cl_device_mem_cache_type) \
+ F(cl_device_info, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, cl_uint)\
+ F(cl_device_info, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, cl_ulong) \
+ F(cl_device_info, CL_DEVICE_GLOBAL_MEM_SIZE, cl_ulong) \
+ F(cl_device_info, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, cl_ulong) \
+ F(cl_device_info, CL_DEVICE_MAX_CONSTANT_ARGS, cl_uint) \
+ F(cl_device_info, CL_DEVICE_LOCAL_MEM_TYPE, cl_device_local_mem_type) \
+ F(cl_device_info, CL_DEVICE_LOCAL_MEM_SIZE, cl_ulong) \
+ F(cl_device_info, CL_DEVICE_ERROR_CORRECTION_SUPPORT, cl_bool) \
+ F(cl_device_info, CL_DEVICE_PROFILING_TIMER_RESOLUTION, ::size_t) \
+ F(cl_device_info, CL_DEVICE_ENDIAN_LITTLE, cl_bool) \
+ F(cl_device_info, CL_DEVICE_AVAILABLE, cl_bool) \
+ F(cl_device_info, CL_DEVICE_COMPILER_AVAILABLE, cl_bool) \
+ F(cl_device_info, CL_DEVICE_EXECUTION_CAPABILITIES, cl_device_exec_capabilities) \
+ F(cl_device_info, CL_DEVICE_QUEUE_PROPERTIES, cl_command_queue_properties) \
+ F(cl_device_info, CL_DEVICE_PLATFORM, cl_platform_id) \
+ F(cl_device_info, CL_DEVICE_NAME, STRING_CLASS) \
+ F(cl_device_info, CL_DEVICE_VENDOR, STRING_CLASS) \
+ F(cl_device_info, CL_DRIVER_VERSION, STRING_CLASS) \
+ F(cl_device_info, CL_DEVICE_PROFILE, STRING_CLASS) \
+ F(cl_device_info, CL_DEVICE_VERSION, STRING_CLASS) \
+ F(cl_device_info, CL_DEVICE_EXTENSIONS, STRING_CLASS) \
+ \
+ F(cl_context_info, CL_CONTEXT_REFERENCE_COUNT, cl_uint) \
+ F(cl_context_info, CL_CONTEXT_DEVICES, VECTOR_CLASS) \
+ F(cl_context_info, CL_CONTEXT_PROPERTIES, VECTOR_CLASS) \
+ \
+ F(cl_event_info, CL_EVENT_COMMAND_QUEUE, cl::CommandQueue) \
+ F(cl_event_info, CL_EVENT_COMMAND_TYPE, cl_command_type) \
+ F(cl_event_info, CL_EVENT_REFERENCE_COUNT, cl_uint) \
+ F(cl_event_info, CL_EVENT_COMMAND_EXECUTION_STATUS, cl_uint) \
+ \
+ F(cl_profiling_info, CL_PROFILING_COMMAND_QUEUED, cl_ulong) \
+ F(cl_profiling_info, CL_PROFILING_COMMAND_SUBMIT, cl_ulong) \
+ F(cl_profiling_info, CL_PROFILING_COMMAND_START, cl_ulong) \
+ F(cl_profiling_info, CL_PROFILING_COMMAND_END, cl_ulong) \
+ \
+ F(cl_mem_info, CL_MEM_TYPE, cl_mem_object_type) \
+ F(cl_mem_info, CL_MEM_FLAGS, cl_mem_flags) \
+ F(cl_mem_info, CL_MEM_SIZE, ::size_t) \
+ F(cl_mem_info, CL_MEM_HOST_PTR, void*) \
+ F(cl_mem_info, CL_MEM_MAP_COUNT, cl_uint) \
+ F(cl_mem_info, CL_MEM_REFERENCE_COUNT, cl_uint) \
+ F(cl_mem_info, CL_MEM_CONTEXT, cl::Context) \
+ \
+ F(cl_image_info, CL_IMAGE_FORMAT, cl_image_format) \
+ F(cl_image_info, CL_IMAGE_ELEMENT_SIZE, ::size_t) \
+ F(cl_image_info, CL_IMAGE_ROW_PITCH, ::size_t) \
+ F(cl_image_info, CL_IMAGE_SLICE_PITCH, ::size_t) \
+ F(cl_image_info, CL_IMAGE_WIDTH, ::size_t) \
+ F(cl_image_info, CL_IMAGE_HEIGHT, ::size_t) \
+ F(cl_image_info, CL_IMAGE_DEPTH, ::size_t) \
+ \
+ F(cl_sampler_info, CL_SAMPLER_REFERENCE_COUNT, cl_uint) \
+ F(cl_sampler_info, CL_SAMPLER_CONTEXT, cl::Context) \
+ F(cl_sampler_info, CL_SAMPLER_NORMALIZED_COORDS, cl_addressing_mode) \
+ F(cl_sampler_info, CL_SAMPLER_ADDRESSING_MODE, cl_filter_mode) \
+ F(cl_sampler_info, CL_SAMPLER_FILTER_MODE, cl_bool) \
+ \
+ F(cl_program_info, CL_PROGRAM_REFERENCE_COUNT, cl_uint) \
+ F(cl_program_info, CL_PROGRAM_CONTEXT, cl::Context) \
+ F(cl_program_info, CL_PROGRAM_NUM_DEVICES, cl_uint) \
+ F(cl_program_info, CL_PROGRAM_DEVICES, VECTOR_CLASS) \
+ F(cl_program_info, CL_PROGRAM_SOURCE, STRING_CLASS) \
+ F(cl_program_info, CL_PROGRAM_BINARY_SIZES, VECTOR_CLASS< ::size_t>) \
+ F(cl_program_info, CL_PROGRAM_BINARIES, VECTOR_CLASS) \
+ \
+ F(cl_program_build_info, CL_PROGRAM_BUILD_STATUS, cl_build_status) \
+ F(cl_program_build_info, CL_PROGRAM_BUILD_OPTIONS, STRING_CLASS) \
+ F(cl_program_build_info, CL_PROGRAM_BUILD_LOG, STRING_CLASS) \
+ \
+ F(cl_kernel_info, CL_KERNEL_FUNCTION_NAME, STRING_CLASS) \
+ F(cl_kernel_info, CL_KERNEL_NUM_ARGS, cl_uint) \
+ F(cl_kernel_info, CL_KERNEL_REFERENCE_COUNT, cl_uint) \
+ F(cl_kernel_info, CL_KERNEL_CONTEXT, cl::Context) \
+ F(cl_kernel_info, CL_KERNEL_PROGRAM, cl::Program) \
+ \
+ F(cl_kernel_work_group_info, CL_KERNEL_WORK_GROUP_SIZE, ::size_t) \
+ F(cl_kernel_work_group_info, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, cl::size_t<3>) \
+ F(cl_kernel_work_group_info, CL_KERNEL_LOCAL_MEM_SIZE, cl_ulong) \
+ \
+ F(cl_command_queue_info, CL_QUEUE_CONTEXT, cl::Context) \
+ F(cl_command_queue_info, CL_QUEUE_DEVICE, cl::Device) \
+ F(cl_command_queue_info, CL_QUEUE_REFERENCE_COUNT, cl_uint) \
+ F(cl_command_queue_info, CL_QUEUE_PROPERTIES, cl_command_queue_properties)
+
+#if defined(CL_VERSION_1_1)
+#define __PARAM_NAME_INFO_1_1(F) \
+ F(cl_context_info, CL_CONTEXT_NUM_DEVICES, cl_uint)\
+ F(cl_device_info, CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE, cl_uint) \
+ F(cl_device_info, CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF, cl_uint) \
+ F(cl_device_info, CL_DEVICE_DOUBLE_FP_CONFIG, cl_device_fp_config) \
+ F(cl_device_info, CL_DEVICE_HALF_FP_CONFIG, cl_device_fp_config) \
+ F(cl_device_info, CL_DEVICE_HOST_UNIFIED_MEMORY, cl_bool) \
+ \
+ F(cl_mem_info, CL_MEM_ASSOCIATED_MEMOBJECT, cl::Memory) \
+ F(cl_mem_info, CL_MEM_OFFSET, ::size_t) \
+ \
+ F(cl_kernel_work_group_info, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, ::size_t) \
+ F(cl_kernel_work_group_info, CL_KERNEL_PRIVATE_MEM_SIZE, cl_ulong) \
+ \
+ F(cl_event_info, CL_EVENT_CONTEXT, cl::Context)
+#endif // CL_VERSION_1_1
+
+#if defined(USE_CL_DEVICE_FISSION)
+#define __PARAM_NAME_DEVICE_FISSION(F) \
+ F(cl_device_info, CL_DEVICE_PARENT_DEVICE_EXT, cl_device_id) \
+ F(cl_device_info, CL_DEVICE_PARTITION_TYPES_EXT, VECTOR_CLASS) \
+ F(cl_device_info, CL_DEVICE_AFFINITY_DOMAINS_EXT, VECTOR_CLASS) \
+ F(cl_device_info, CL_DEVICE_REFERENCE_COUNT_EXT , cl_uint) \
+ F(cl_device_info, CL_DEVICE_PARTITION_STYLE_EXT, VECTOR_CLASS)
+#endif // USE_CL_DEVICE_FISSION
+
+template
+struct param_traits {};
+
+#define __DECLARE_PARAM_TRAITS(token, param_name, T) \
+struct token; \
+template<> \
+struct param_traits \
+{ \
+ enum { value = param_name }; \
+ typedef T param_type; \
+};
+
+__PARAM_NAME_INFO_1_0(__DECLARE_PARAM_TRAITS)
+#if defined(CL_VERSION_1_1)
+__PARAM_NAME_INFO_1_1(__DECLARE_PARAM_TRAITS)
+#endif // CL_VERSION_1_1
+
+#if defined(USE_CL_DEVICE_FISSION)
+__PARAM_NAME_DEVICE_FISSION(__DECLARE_PARAM_TRAITS);
+#endif // USE_CL_DEVICE_FISSION
+
+#undef __DECLARE_PARAM_TRAITS
+
+// Convenience functions
+
+template
+inline cl_int
+getInfo(Func f, cl_uint name, T* param)
+{
+ return GetInfoHelper::get(f, name, param);
+}
+
+template
+struct GetInfoFunctor0
+{
+ Func f_; const Arg0& arg0_;
+ cl_int operator ()(
+ cl_uint param, ::size_t size, void* value, ::size_t* size_ret)
+ { return f_(arg0_, param, size, value, size_ret); }
+};
+
+template
+struct GetInfoFunctor1
+{
+ Func f_; const Arg0& arg0_; const Arg1& arg1_;
+ cl_int operator ()(
+ cl_uint param, ::size_t size, void* value, ::size_t* size_ret)
+ { return f_(arg0_, arg1_, param, size, value, size_ret); }
+};
+
+template
+inline cl_int
+getInfo(Func f, const Arg0& arg0, cl_uint name, T* param)
+{
+ GetInfoFunctor0 f0 = { f, arg0 };
+ return GetInfoHelper, T>
+ ::get(f0, name, param);
+}
+
+template
+inline cl_int
+getInfo(Func f, const Arg0& arg0, const Arg1& arg1, cl_uint name, T* param)
+{
+ GetInfoFunctor1 f0 = { f, arg0, arg1 };
+ return GetInfoHelper, T>
+ ::get(f0, name, param);
+}
+
+template
+struct ReferenceHandler
+{ };
+
+template <>
+struct ReferenceHandler
+{
+ // cl_device_id does not have retain().
+ static cl_int retain(cl_device_id)
+ { return CL_INVALID_DEVICE; }
+ // cl_device_id does not have release().
+ static cl_int release(cl_device_id)
+ { return CL_INVALID_DEVICE; }
+};
+
+template <>
+struct ReferenceHandler
+{
+ // cl_platform_id does not have retain().
+ static cl_int retain(cl_platform_id)
+ { return CL_INVALID_PLATFORM; }
+ // cl_platform_id does not have release().
+ static cl_int release(cl_platform_id)
+ { return CL_INVALID_PLATFORM; }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_context context)
+ { return ::clRetainContext(context); }
+ static cl_int release(cl_context context)
+ { return ::clReleaseContext(context); }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_command_queue queue)
+ { return ::clRetainCommandQueue(queue); }
+ static cl_int release(cl_command_queue queue)
+ { return ::clReleaseCommandQueue(queue); }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_mem memory)
+ { return ::clRetainMemObject(memory); }
+ static cl_int release(cl_mem memory)
+ { return ::clReleaseMemObject(memory); }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_sampler sampler)
+ { return ::clRetainSampler(sampler); }
+ static cl_int release(cl_sampler sampler)
+ { return ::clReleaseSampler(sampler); }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_program program)
+ { return ::clRetainProgram(program); }
+ static cl_int release(cl_program program)
+ { return ::clReleaseProgram(program); }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_kernel kernel)
+ { return ::clRetainKernel(kernel); }
+ static cl_int release(cl_kernel kernel)
+ { return ::clReleaseKernel(kernel); }
+};
+
+template <>
+struct ReferenceHandler
+{
+ static cl_int retain(cl_event event)
+ { return ::clRetainEvent(event); }
+ static cl_int release(cl_event event)
+ { return ::clReleaseEvent(event); }
+};
+
+template
+class Wrapper
+{
+public:
+ typedef T cl_type;
+
+protected:
+ cl_type object_;
+
+public:
+ Wrapper() : object_(NULL) { }
+
+ ~Wrapper()
+ {
+ if (object_ != NULL) { release(); }
+ }
+
+ Wrapper(const Wrapper& rhs)
+ {
+ object_ = rhs.object_;
+ if (object_ != NULL) { retain(); }
+ }
+
+ Wrapper& operator = (const Wrapper& rhs)
+ {
+ if (object_ != NULL) { release(); }
+ object_ = rhs.object_;
+ if (object_ != NULL) { retain(); }
+ return *this;
+ }
+
+ cl_type operator ()() const { return object_; }
+
+ cl_type& operator ()() { return object_; }
+
+protected:
+
+ cl_int retain() const
+ {
+ return ReferenceHandler::retain(object_);
+ }
+
+ cl_int release() const
+ {
+ return ReferenceHandler::release(object_);
+ }
+};
+
+#if defined(__CL_ENABLE_EXCEPTIONS)
+static inline cl_int errHandler (
+ cl_int err,
+ const char * errStr = NULL) throw(Error)
+{
+ if (err != CL_SUCCESS) {
+ throw Error(err, errStr);
+ }
+ return err;
+}
+#else
+static inline cl_int errHandler (cl_int err, const char * errStr = NULL)
+{
+ return err;
+}
+#endif // __CL_ENABLE_EXCEPTIONS
+
+} // namespace detail
+//! \endcond
+
+/*! \stuct ImageFormat
+ * \brief ImageFormat interface fro cl_image_format.
+ */
+struct ImageFormat : public cl_image_format
+{
+ ImageFormat(){}
+
+ ImageFormat(cl_channel_order order, cl_channel_type type)
+ {
+ image_channel_order = order;
+ image_channel_data_type = type;
+ }
+
+ ImageFormat& operator = (const ImageFormat& rhs)
+ {
+ if (this != &rhs) {
+ this->image_channel_data_type = rhs.image_channel_data_type;
+ this->image_channel_order = rhs.image_channel_order;
+ }
+ return *this;
+ }
+};
+
+/*! \class Device
+ * \brief Device interface for cl_device_id.
+ */
+class Device : public detail::Wrapper
+{
+public:
+ Device(cl_device_id device) { object_ = device; }
+
+ Device() : detail::Wrapper() { }
+
+ Device(const Device& device) : detail::Wrapper(device) { }
+
+ Device& operator = (const Device& rhs)
+ {
+ if (this != &rhs) {
+ detail::Wrapper::operator=(rhs);
+ }
+ return *this;
+ }
+
+ template
+ cl_int getInfo(cl_device_info name, T* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetDeviceInfo, object_, name, param),
+ __GET_DEVICE_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_device_info, name>::param_type param;
+ cl_int result = getInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+
+#if defined(USE_CL_DEVICE_FISSION)
+ cl_int createSubDevices(
+ const cl_device_partition_property_ext * properties,
+ VECTOR_CLASS* devices)
+ {
+ typedef CL_API_ENTRY cl_int
+ ( CL_API_CALL * PFN_clCreateSubDevicesEXT)(
+ cl_device_id /*in_device*/,
+ const cl_device_partition_property_ext * /* properties */,
+ cl_uint /*num_entries*/,
+ cl_device_id * /*out_devices*/,
+ cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+ static PFN_clCreateSubDevicesEXT pfn_clCreateSubDevicesEXT = NULL;
+ __INIT_CL_EXT_FCN_PTR(clCreateSubDevicesEXT);
+
+ cl_uint n = 0;
+ cl_int err = pfn_clCreateSubDevicesEXT(object_, properties, 0, NULL, &n);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __CREATE_SUB_DEVICES);
+ }
+
+ cl_device_id* ids = (cl_device_id*) alloca(n * sizeof(cl_device_id));
+ err = pfn_clCreateSubDevicesEXT(object_, properties, n, ids, NULL);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __CREATE_SUB_DEVICES);
+ }
+
+ devices->assign(&ids[0], &ids[n]);
+ return CL_SUCCESS;
+ }
+#endif
+};
+
+/*! \class Platform
+ * \brief Platform interface.
+ */
+class Platform : public detail::Wrapper
+{
+public:
+ static const Platform null();
+
+ Platform(cl_platform_id platform) { object_ = platform; }
+
+ Platform() : detail::Wrapper() { }
+
+ Platform(const Platform& platform) : detail::Wrapper(platform) { }
+
+ Platform& operator = (const Platform& rhs)
+ {
+ if (this != &rhs) {
+ detail::Wrapper::operator=(rhs);
+ }
+ return *this;
+ }
+
+ cl_int getInfo(cl_platform_info name, STRING_CLASS* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetPlatformInfo, object_, name, param),
+ __GET_PLATFORM_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_platform_info, name>::param_type param;
+ cl_int result = getInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+
+ cl_int getDevices(
+ cl_device_type type,
+ VECTOR_CLASS* devices) const
+ {
+ cl_uint n = 0;
+ cl_int err = ::clGetDeviceIDs(object_, type, 0, NULL, &n);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_DEVICE_IDS_ERR);
+ }
+
+ cl_device_id* ids = (cl_device_id*) alloca(n * sizeof(cl_device_id));
+ err = ::clGetDeviceIDs(object_, type, n, ids, NULL);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_DEVICE_IDS_ERR);
+ }
+
+ devices->assign(&ids[0], &ids[n]);
+ return CL_SUCCESS;
+ }
+
+#if defined(USE_DX_INTEROP)
+ /*! \brief Get the list of available D3D10 devices.
+ *
+ * \param d3d_device_source.
+ *
+ * \param d3d_object.
+ *
+ * \param d3d_device_set.
+ *
+ * \param devices returns a vector of OpenCL D3D10 devices found. The cl::Device
+ * values returned in devices can be used to identify a specific OpenCL
+ * device. If \a devices argument is NULL, this argument is ignored.
+ *
+ * \return One of the following values:
+ * - CL_SUCCESS if the function is executed successfully.
+ *
+ * The application can query specific capabilities of the OpenCL device(s)
+ * returned by cl::getDevices. This can be used by the application to
+ * determine which device(s) to use.
+ *
+ * \note In the case that exceptions are enabled and a return value
+ * other than CL_SUCCESS is generated, then cl::Error exception is
+ * generated.
+ */
+ cl_int getDevices(
+ cl_d3d10_device_source_khr d3d_device_source,
+ void * d3d_object,
+ cl_d3d10_device_set_khr d3d_device_set,
+ VECTOR_CLASS* devices) const
+ {
+ typedef CL_API_ENTRY cl_int (CL_API_CALL *PFN_clGetDeviceIDsFromD3D10KHR)(
+ cl_platform_id platform,
+ cl_d3d10_device_source_khr d3d_device_source,
+ void * d3d_object,
+ cl_d3d10_device_set_khr d3d_device_set,
+ cl_uint num_entries,
+ cl_device_id * devices,
+ cl_uint* num_devices);
+
+ static PFN_clGetDeviceIDsFromD3D10KHR pfn_clGetDeviceIDsFromD3D10KHR = NULL;
+ __INIT_CL_EXT_FCN_PTR(clGetDeviceIDsFromD3D10KHR);
+
+ cl_uint n = 0;
+ cl_int err = pfn_clGetDeviceIDsFromD3D10KHR(
+ object_,
+ d3d_device_source,
+ d3d_object,
+ d3d_device_set,
+ 0,
+ NULL,
+ &n);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_DEVICE_IDS_ERR);
+ }
+
+ cl_device_id* ids = (cl_device_id*) alloca(n * sizeof(cl_device_id));
+ err = pfn_clGetDeviceIDsFromD3D10KHR(
+ object_,
+ d3d_device_source,
+ d3d_object,
+ d3d_device_set,
+ n,
+ ids,
+ NULL);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_DEVICE_IDS_ERR);
+ }
+
+ devices->assign(&ids[0], &ids[n]);
+ return CL_SUCCESS;
+ }
+#endif
+
+ static cl_int get(
+ VECTOR_CLASS* platforms)
+ {
+ cl_uint n = 0;
+ cl_int err = ::clGetPlatformIDs(0, NULL, &n);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_PLATFORM_IDS_ERR);
+ }
+
+ cl_platform_id* ids = (cl_platform_id*) alloca(
+ n * sizeof(cl_platform_id));
+ err = ::clGetPlatformIDs(n, ids, NULL);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_PLATFORM_IDS_ERR);
+ }
+
+ platforms->assign(&ids[0], &ids[n]);
+ return CL_SUCCESS;
+ }
+};
+
+static inline cl_int
+UnloadCompiler()
+{
+ return ::clUnloadCompiler();
+}
+
+class Context : public detail::Wrapper
+{
+public:
+ Context(
+ const VECTOR_CLASS& devices,
+ cl_context_properties* properties = NULL,
+ void (CL_CALLBACK * notifyFptr)(
+ const char *,
+ const void *,
+ ::size_t,
+ void *) = NULL,
+ void* data = NULL,
+ cl_int* err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateContext(
+ properties, (cl_uint) devices.size(),
+ (cl_device_id*) &devices.front(),
+ notifyFptr, data, &error);
+
+ detail::errHandler(error, __CREATE_CONTEXT_FROM_TYPE_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Context(
+ cl_device_type type,
+ cl_context_properties* properties = NULL,
+ void (CL_CALLBACK * notifyFptr)(
+ const char *,
+ const void *,
+ ::size_t,
+ void *) = NULL,
+ void* data = NULL,
+ cl_int* err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateContextFromType(
+ properties, type, notifyFptr, data, &error);
+
+ detail::errHandler(error, __CREATE_CONTEXT_FROM_TYPE_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Context() : detail::Wrapper() { }
+
+ Context(const Context& context) : detail::Wrapper(context) { }
+
+ Context& operator = (const Context& rhs)
+ {
+ if (this != &rhs) {
+ detail::Wrapper::operator=(rhs);
+ }
+ return *this;
+ }
+
+ template
+ cl_int getInfo(cl_context_info name, T* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetContextInfo, object_, name, param),
+ __GET_CONTEXT_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_context_info, name>::param_type param;
+ cl_int result = getInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+
+ cl_int getSupportedImageFormats(
+ cl_mem_flags flags,
+ cl_mem_object_type type,
+ VECTOR_CLASS* formats) const
+ {
+ cl_uint numEntries;
+ cl_int err = ::clGetSupportedImageFormats(
+ object_,
+ flags,
+ type,
+ 0,
+ NULL,
+ &numEntries);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_SUPPORTED_IMAGE_FORMATS_ERR);
+ }
+
+ ImageFormat* value = (ImageFormat*)
+ alloca(numEntries * sizeof(ImageFormat));
+ err = ::clGetSupportedImageFormats(
+ object_,
+ flags,
+ type,
+ numEntries,
+ (cl_image_format*) value,
+ NULL);
+ if (err != CL_SUCCESS) {
+ return detail::errHandler(err, __GET_SUPPORTED_IMAGE_FORMATS_ERR);
+ }
+
+ formats->assign(&value[0], &value[numEntries]);
+ return CL_SUCCESS;
+ }
+};
+
+__GET_INFO_HELPER_WITH_RETAIN(cl::Context)
+
+/*! \class Event
+ * \brief Event interface for cl_event.
+ */
+class Event : public detail::Wrapper
+{
+public:
+ Event() : detail::Wrapper() { }
+
+ Event(const Event& event) : detail::Wrapper(event) { }
+
+ Event& operator = (const Event& rhs)
+ {
+ if (this != &rhs) {
+ detail::Wrapper::operator=(rhs);
+ }
+ return *this;
+ }
+
+ template
+ cl_int getInfo(cl_event_info name, T* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetEventInfo, object_, name, param),
+ __GET_EVENT_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_event_info, name>::param_type param;
+ cl_int result = getInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+
+ template
+ cl_int getProfilingInfo(cl_profiling_info name, T* param) const
+ {
+ return detail::errHandler(detail::getInfo(
+ &::clGetEventProfilingInfo, object_, name, param),
+ __GET_EVENT_PROFILE_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getProfilingInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_profiling_info, name>::param_type param;
+ cl_int result = getProfilingInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+
+ cl_int wait() const
+ {
+ return detail::errHandler(
+ ::clWaitForEvents(1, &object_),
+ __WAIT_FOR_EVENTS_ERR);
+ }
+
+#if defined(CL_VERSION_1_1)
+ cl_int setCallback(
+ cl_int type,
+ void (CL_CALLBACK * pfn_notify)(cl_event, cl_int, void *),
+ void * user_data = NULL)
+ {
+ return detail::errHandler(
+ ::clSetEventCallback(
+ object_,
+ type,
+ pfn_notify,
+ user_data),
+ __SET_EVENT_CALLBACK_ERR);
+ }
+#endif
+
+ static cl_int
+ waitForEvents(const VECTOR_CLASS& events)
+ {
+ return detail::errHandler(
+ ::clWaitForEvents(
+ (cl_uint) events.size(), (cl_event*)&events.front()),
+ __WAIT_FOR_EVENTS_ERR);
+ }
+};
+
+__GET_INFO_HELPER_WITH_RETAIN(cl::Event)
+
+#if defined(CL_VERSION_1_1)
+/*! \class UserEvent
+ * \brief User event interface for cl_event.
+ */
+class UserEvent : public Event
+{
+public:
+ UserEvent(
+ const Context& context,
+ cl_int * err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateUserEvent(
+ context(),
+ &error);
+
+ detail::errHandler(error, __CREATE_USER_EVENT_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ UserEvent() : Event() { }
+
+ UserEvent(const UserEvent& event) : Event(event) { }
+
+ UserEvent& operator = (const UserEvent& rhs)
+ {
+ if (this != &rhs) {
+ Event::operator=(rhs);
+ }
+ return *this;
+ }
+
+ cl_int setStatus(cl_int status)
+ {
+ return detail::errHandler(
+ ::clSetUserEventStatus(object_,status),
+ __SET_USER_EVENT_STATUS_ERR);
+ }
+};
+#endif
+
+inline static cl_int
+WaitForEvents(const VECTOR_CLASS& events)
+{
+ return detail::errHandler(
+ ::clWaitForEvents(
+ (cl_uint) events.size(), (cl_event*)&events.front()),
+ __WAIT_FOR_EVENTS_ERR);
+}
+
+/*! \class Memory
+ * \brief Memory interface for cl_mem.
+ */
+class Memory : public detail::Wrapper
+{
+public:
+ Memory() : detail::Wrapper() { }
+
+ Memory(const Memory& memory) : detail::Wrapper(memory) { }
+
+ Memory& operator = (const Memory& rhs)
+ {
+ if (this != &rhs) {
+ detail::Wrapper::operator=(rhs);
+ }
+ return *this;
+ }
+
+ template
+ cl_int getInfo(cl_mem_info name, T* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetMemObjectInfo, object_, name, param),
+ __GET_MEM_OBJECT_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_mem_info, name>::param_type param;
+ cl_int result = getInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+
+#if defined(CL_VERSION_1_1)
+ cl_int setDestructorCallback(
+ void (CL_CALLBACK * pfn_notify)(cl_mem, void *),
+ void * user_data = NULL)
+ {
+ return detail::errHandler(
+ ::clSetMemObjectDestructorCallback(
+ object_,
+ pfn_notify,
+ user_data),
+ __SET_MEM_OBJECT_DESTRUCTOR_CALLBACK_ERR);
+ }
+#endif
+
+};
+
+__GET_INFO_HELPER_WITH_RETAIN(cl::Memory)
+
+/*! \class Buffer
+ * \brief Memory buffer interface.
+ */
+class Buffer : public Memory
+{
+public:
+ Buffer(
+ const Context& context,
+ cl_mem_flags flags,
+ ::size_t size,
+ void* host_ptr = NULL,
+ cl_int* err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateBuffer(context(), flags, size, host_ptr, &error);
+
+ detail::errHandler(error, __CREATE_BUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Buffer() : Memory() { }
+
+ Buffer(const Buffer& buffer) : Memory(buffer) { }
+
+ Buffer& operator = (const Buffer& rhs)
+ {
+ if (this != &rhs) {
+ Memory::operator=(rhs);
+ }
+ return *this;
+ }
+
+#if defined(CL_VERSION_1_1)
+ Buffer createSubBuffer(
+ cl_mem_flags flags,
+ cl_buffer_create_type buffer_create_type,
+ const void * buffer_create_info,
+ cl_int * err = NULL)
+ {
+ Buffer result;
+ cl_int error;
+ result.object_ = ::clCreateSubBuffer(
+ object_,
+ flags,
+ buffer_create_type,
+ buffer_create_info,
+ &error);
+
+ detail::errHandler(error, __CREATE_SUBBUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+
+ return result;
+ }
+#endif
+};
+
+#if defined (USE_DX_INTEROP)
+class BufferD3D10 : public Buffer
+{
+public:
+ typedef CL_API_ENTRY cl_mem (CL_API_CALL *PFN_clCreateFromD3D10BufferKHR)(
+ cl_context context, cl_mem_flags flags, ID3D10Buffer* buffer,
+ cl_int* errcode_ret);
+
+ BufferD3D10(
+ const Context& context,
+ cl_mem_flags flags,
+ ID3D10Buffer* bufobj,
+ cl_int * err = NULL)
+ {
+ static PFN_clCreateFromD3D10BufferKHR pfn_clCreateFromD3D10BufferKHR = NULL;
+ __INIT_CL_EXT_FCN_PTR(clCreateFromD3D10BufferKHR);
+
+ cl_int error;
+ object_ = pfn_clCreateFromD3D10BufferKHR(
+ context(),
+ flags,
+ bufobj,
+ &error);
+
+ detail::errHandler(error, __CREATE_GL_BUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ BufferD3D10() : Buffer() { }
+
+ BufferD3D10(const BufferD3D10& buffer) : Buffer(buffer) { }
+
+ BufferD3D10& operator = (const BufferD3D10& rhs)
+ {
+ if (this != &rhs) {
+ Buffer::operator=(rhs);
+ }
+ return *this;
+ }
+};
+#endif
+
+/*! \class BufferGL
+ * \brief Memory buffer interface for GL interop.
+ */
+class BufferGL : public Buffer
+{
+public:
+ BufferGL(
+ const Context& context,
+ cl_mem_flags flags,
+ GLuint bufobj,
+ cl_int * err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateFromGLBuffer(
+ context(),
+ flags,
+ bufobj,
+ &error);
+
+ detail::errHandler(error, __CREATE_GL_BUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ BufferGL() : Buffer() { }
+
+ BufferGL(const BufferGL& buffer) : Buffer(buffer) { }
+
+ BufferGL& operator = (const BufferGL& rhs)
+ {
+ if (this != &rhs) {
+ Buffer::operator=(rhs);
+ }
+ return *this;
+ }
+
+ cl_int getObjectInfo(
+ cl_gl_object_type *type,
+ GLuint * gl_object_name)
+ {
+ return detail::errHandler(
+ ::clGetGLObjectInfo(object_,type,gl_object_name),
+ __GET_GL_OBJECT_INFO_ERR);
+ }
+};
+
+/*! \class BufferRenderGL
+ * \brief Memory buffer interface for GL interop with renderbuffer.
+ */
+class BufferRenderGL : public Buffer
+{
+public:
+ BufferRenderGL(
+ const Context& context,
+ cl_mem_flags flags,
+ GLuint bufobj,
+ cl_int * err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateFromGLRenderbuffer(
+ context(),
+ flags,
+ bufobj,
+ &error);
+
+ detail::errHandler(error, __CREATE_GL_BUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ BufferRenderGL() : Buffer() { }
+
+ BufferRenderGL(const BufferGL& buffer) : Buffer(buffer) { }
+
+ BufferRenderGL& operator = (const BufferRenderGL& rhs)
+ {
+ if (this != &rhs) {
+ Buffer::operator=(rhs);
+ }
+ return *this;
+ }
+
+ cl_int getObjectInfo(
+ cl_gl_object_type *type,
+ GLuint * gl_object_name)
+ {
+ return detail::errHandler(
+ ::clGetGLObjectInfo(object_,type,gl_object_name),
+ __GET_GL_OBJECT_INFO_ERR);
+ }
+};
+
+/*! \class Image
+ * \brief Base class interface for all images.
+ */
+class Image : public Memory
+{
+protected:
+ Image() : Memory() { }
+
+ Image(const Image& image) : Memory(image) { }
+
+ Image& operator = (const Image& rhs)
+ {
+ if (this != &rhs) {
+ Memory::operator=(rhs);
+ }
+ return *this;
+ }
+public:
+ template
+ cl_int getImageInfo(cl_image_info name, T* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetImageInfo, object_, name, param),
+ __GET_IMAGE_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getImageInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_image_info, name>::param_type param;
+ cl_int result = getImageInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+};
+
+/*! \class Image2D
+ * \brief Image interface for 2D images.
+ */
+class Image2D : public Image
+{
+public:
+ Image2D(
+ const Context& context,
+ cl_mem_flags flags,
+ ImageFormat format,
+ ::size_t width,
+ ::size_t height,
+ ::size_t row_pitch = 0,
+ void* host_ptr = NULL,
+ cl_int* err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateImage2D(
+ context(), flags,&format, width, height, row_pitch, host_ptr, &error);
+
+ detail::errHandler(error, __CREATE_IMAGE2D_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Image2D() { }
+
+ Image2D(const Image2D& image2D) : Image(image2D) { }
+
+ Image2D& operator = (const Image2D& rhs)
+ {
+ if (this != &rhs) {
+ Image::operator=(rhs);
+ }
+ return *this;
+ }
+};
+
+/*! \class Image2DGL
+ * \brief 2D image interface for GL interop.
+ */
+class Image2DGL : public Image2D
+{
+public:
+ Image2DGL(
+ const Context& context,
+ cl_mem_flags flags,
+ GLenum target,
+ GLint miplevel,
+ GLuint texobj,
+ cl_int * err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateFromGLTexture2D(
+ context(),
+ flags,
+ target,
+ miplevel,
+ texobj,
+ &error);
+
+ detail::errHandler(error, __CREATE_GL_BUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Image2DGL() : Image2D() { }
+
+ Image2DGL(const Image2DGL& image) : Image2D(image) { }
+
+ Image2DGL& operator = (const Image2DGL& rhs)
+ {
+ if (this != &rhs) {
+ Image2D::operator=(rhs);
+ }
+ return *this;
+ }
+};
+
+/*! \class Image3D
+ * \brief Image interface for 3D images.
+ */
+class Image3D : public Image
+{
+public:
+ Image3D(
+ const Context& context,
+ cl_mem_flags flags,
+ ImageFormat format,
+ ::size_t width,
+ ::size_t height,
+ ::size_t depth,
+ ::size_t row_pitch = 0,
+ ::size_t slice_pitch = 0,
+ void* host_ptr = NULL,
+ cl_int* err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateImage3D(
+ context(), flags, &format, width, height, depth, row_pitch,
+ slice_pitch, host_ptr, &error);
+
+ detail::errHandler(error, __CREATE_IMAGE3D_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Image3D() { }
+
+ Image3D(const Image3D& image3D) : Image(image3D) { }
+
+ Image3D& operator = (const Image3D& rhs)
+ {
+ if (this != &rhs) {
+ Image::operator=(rhs);
+ }
+ return *this;
+ }
+};
+
+/*! \class Image2DGL
+ * \brief 2D image interface for GL interop.
+ */
+class Image3DGL : public Image3D
+{
+public:
+ Image3DGL(
+ const Context& context,
+ cl_mem_flags flags,
+ GLenum target,
+ GLint miplevel,
+ GLuint texobj,
+ cl_int * err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateFromGLTexture3D(
+ context(),
+ flags,
+ target,
+ miplevel,
+ texobj,
+ &error);
+
+ detail::errHandler(error, __CREATE_GL_BUFFER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Image3DGL() : Image3D() { }
+
+ Image3DGL(const Image3DGL& image) : Image3D(image) { }
+
+ Image3DGL& operator = (const Image3DGL& rhs)
+ {
+ if (this != &rhs) {
+ Image3D::operator=(rhs);
+ }
+ return *this;
+ }
+};
+
+/*! \class Sampler
+ * \brief Sampler interface for cl_sampler.
+ */
+class Sampler : public detail::Wrapper
+{
+public:
+ Sampler() { }
+
+ Sampler(
+ const Context& context,
+ cl_bool normalized_coords,
+ cl_addressing_mode addressing_mode,
+ cl_filter_mode filter_mode,
+ cl_int* err = NULL)
+ {
+ cl_int error;
+ object_ = ::clCreateSampler(
+ context(),
+ normalized_coords,
+ addressing_mode,
+ filter_mode,
+ &error);
+
+ detail::errHandler(error, __CREATE_SAMPLER_ERR);
+ if (err != NULL) {
+ *err = error;
+ }
+ }
+
+ Sampler(const Sampler& sampler) : detail::Wrapper(sampler) { }
+
+ Sampler& operator = (const Sampler& rhs)
+ {
+ if (this != &rhs) {
+ detail::Wrapper::operator=(rhs);
+ }
+ return *this;
+ }
+
+ template
+ cl_int getInfo(cl_sampler_info name, T* param) const
+ {
+ return detail::errHandler(
+ detail::getInfo(&::clGetSamplerInfo, object_, name, param),
+ __GET_SAMPLER_INFO_ERR);
+ }
+
+ template typename
+ detail::param_traits::param_type
+ getInfo(cl_int* err = NULL) const
+ {
+ typename detail::param_traits<
+ detail::cl_sampler_info, name>::param_type param;
+ cl_int result = getInfo(name, ¶m);
+ if (err != NULL) {
+ *err = result;
+ }
+ return param;
+ }
+};
+
+__GET_INFO_HELPER_WITH_RETAIN(cl::Sampler)
+
+class Program;
+class CommandQueue;
+class Kernel;
+
+/*! \class NDRange
+ * \brief NDRange interface
+ */
+class NDRange
+{
+private:
+ size_t<3> sizes_;
+ cl_uint dimensions_;
+
+public:
+ NDRange()
+ : dimensions_(0)
+ { }
+
+ NDRange(::size_t size0)
+ : dimensions_(1)
+ {
+ sizes_.push_back(size0);
+ }
+
+ NDRange(::size_t size0, ::size_t size1)
+ : dimensions_(2)
+ {
+ sizes_.push_back(size0);
+ sizes_.push_back(size1);
+ }
+
+ NDRange(::size_t size0, ::size_t size1, ::size_t size2)
+ : dimensions_(3)
+ {
+ sizes_.push_back(size0);
+ sizes_.push_back(size1);
+ sizes_.push_back(size2);
+ }
+
+ operator const ::size_t*() const { return (const ::size_t*) sizes_; }
+ ::size_t dimensions() const { return dimensions_; }
+};
+
+static const NDRange NullRange;
+
+/*!
+ * \struct LocalSpaceArg
+ * \brief Local address raper for use with Kernel::setArg
+ */
+struct LocalSpaceArg
+{
+ ::size_t size_;
+};
+
+namespace detail {
+
+template
+struct KernelArgumentHandler
+{
+ static ::size_t size(const T&) { return sizeof(T); }
+ static T* ptr(T& value) { return &value; }
+};
+
+template <>
+struct KernelArgumentHandler
+{
+ static ::size_t size(const LocalSpaceArg& value) { return value.size_; }
+ static void* ptr(LocalSpaceArg&) { return NULL; }
+};
+
+}
+//! \endcond
+
+inline LocalSpaceArg
+__local(::size_t size)
+{
+ LocalSpaceArg ret = { size };
+ return ret;
+}
+
+class KernelFunctor;
+
+/*! \class Kernel
+ * \brief Kernel interface that implements cl_kernel
+ */
+class Kernel : public detail::Wrapper