sgminer を高速化サセようぜ
1 ：ハッテン場五段：2016/04/10 02:01:05 (8年前) 512.34MONA/5人
Lyra2REv2 での採掘は現状 Radeon は圧倒的不利な状況です。K6 時代からの AMD 信者としては悔しくてたまらない。なんとかして高速化サセねば！
言うばっかではなんなので、グラボでマワしたら８％ほど高速化できたものを公開します。APU では若干速度が落ちます。sgminer の kernel ディレクトリにつっこんで下さい。
http://kernel.hattenba.xyz/kernel-20160409.zip
【参考文献】
scryptがGPUに破られる時 - びりあるの研究ノート
https://blog.visvirial.com/articles/519
2 ：izuna五段：2016/04/10 07:20:05 (8年前) 0.00000184MONA/1人
去年の9月頃に、いろいろと遊んでいたのだけど、あれから状況が変わっていなければ、-g 1から -g 2にしてスレッドを1から2にしても、性能が10%以上、高速化される。(R7 250XEによる測定結果)
そのときに考えた、高速化される、理由は
Lyra2REv2は、いろいろなハッシュ関数を７個ならべたもので、真ん中の１つだけが、メモリ帯域を大食いする関数。他のハッシュ関数は、メモリ帯域をあまり食わない。
２スレッドにすれば、１スレッド目がメモリ帯域を大食いする関数で演算器がアイドルになっている状況で、２スレッド目が、メモリ帯域を食わない関数を実行できれば、-g 2で性能が向上する理由になる。
>>1 の高速化版は、-g オプションを使うと、性能は、どれくらいになるのでしょうか？
3 ：名無し名誉名人教士：2016/04/10 08:10:47 (8年前) 0.1MONA/1人
>>1
Lyra2REv2の7個の関数の真ん中がプロセスあたり1536Byte使用しているため、キャッシュに収まりきらない…
R9 280Xで考えると、2,048コア、1048576Byte(1MB)キャッシュだから、
1048576Byte÷2048コア＝512Byte/コア
ってことで、使用するメモリを1/3にできれば(近づければ)高速化できる。
なお、Geforceでは…
Geforce GTX970は1664コア、1835008Byte(1792kB)キャッシュ1835008Byte÷1664コア＝1102.8Byte/コア
Geforce GTX750Tiは640コア、2097152Byte(2MB)キャッシュ
2097152Byte÷640コア＝3276.8Byte/コア
4 ：なむやん七段教士：2016/04/10 09:34:43 (8年前) 1.1MONA/2人
オレのR7 250が火を噴くぜ！
5 ：izuna五段：2016/04/10 10:01:22 (8年前) 0MONA/0人
>>3
使用するメモリを1/3にするのは、難しいかもしれないが、同時に使用するメモリを1/3にしていく方向は、あるかもしれない。
真ん中の関数が、一斉にメモリにアクセスするから、キャッシュから、あふれて、性能がかなり低下する。
1計算スレッドで同時に3っつ計算してみよう。このとき確保するメモリは1536Byteで、これを3っつで共用する。
1つめの計算では、最初に真ん中のハッシュをする。
2つめの計算では、真ん中で真ん中のハッシュをする。
3つめの計算では、最後に真ん中のハッシュをする。
1つめの計算で、最初に真ん中のハッシュを計算できるようにする工夫が、必要だけど、これで使用するメモリが1/3になるかな。
あ、全然、適当なアイディアだから、うまくいくとは、限らないよ。
6 ：izuna五段：2016/04/10 10:08:01 (8年前) 0MONA/0人
>>4
織田信長の３段撃ち、をするのだ！
7 ：ハッテン場五段：2016/04/12 10:29:22 (8年前) 0.1MONA/1人
みなさま、投げモナとご意見ありがとです。
かなり濃いモナをぶつけて下さったようで恐縮です。
色々とヤル気がでてきました。
現実的なのはメモリ再配置か、グローバルメモリ、プログラム中で言えば DMatrix 配列への読み書きをいかに減らすか、はたまたその両方かと思ってます。
1536 byte 全部読まなくていいんだよ！とか、全部書き込まなくてもいいんだよ！とかそういう状況を見つけ出せたら。
いろいろと試しちゃいるんですが、結果がぜんぜんでません。
仕事もモンハンも射精活動もかわいい子の物色もしないといけなくて、なかなか進まない。
とりあえず休みたい
8 ：hururamu六段：2016/04/12 20:45:11 (8年前) 0.2MONA/1人
>全部読まなくていいんだよ！とか、全部書き込まなくてもいいんだよ！
>仕事もモンハンも射精活動もかわいい子の物色もしないといけなくて、なかなか進まない。
全部やらなくてもいいんだよ！
9 ：ハッテン場五段：2016/04/13 04:01:10 (8年前) 0.1MONA/1人
はーい 当社比１２％性能アップ版が作れましたよー イクイクー
http://kernel.hattenba.xyz/kernel-20160413.zip
OpenCL kernel のほうじゃなくて sgminer のソース見てて気づいたんですけど、lyra2v2 のとこで使ってる 1536 byte のデータって、ほかのところでも使ってるわけじゃなくて、lyra2v2 のとこだけでしか使ってないみたいですね。
ということはデータの整合性がとれる範囲であれば、なぶっちゃっても問題なし、と。
そこんとこ省いたら性能が４％ほどアップ致しました。
10 ：ハッテン場五段：2016/04/13 04:05:29 (8年前) 0MONA/0人
今回のは、DMatrix への最後の方の書き込みを一定の条件下であればヤラないことにしました。OpenCL の苦手な比較演算が激遅グローバルメモリへのアクセスに打ち勝ったのです。
>>8
オスとして・・その衝動は止められない！！
11 ：ハッテン場五段：2016/04/13 14:58:53 (8年前) 0.1MONA/1人
あ、すみません、アップロードするファイルを間違えてたっぽいです。
差し替えてますんで、コンパイルが通らなかった方は今一度ダウンロードお願いします。このファイルは 17402 バイトです。間違えてたのは 17403 バイトです。
http://kernel.hattenba.xyz/kernel-20160413.zip
12 ：リキプロマン六段：2016/04/13 18:49:37 (8年前) 0MONA/0人
reduceDuplexRowf を3回ループで回す際に、最後の1回の時に別のメソッドを用いて条件分岐させて、DMatrixの書き込みを減らしているのですね。
ところで、s1のこの条件はどうやって見つけたのですか？自分もLyra2v2の最適化に興味があるので良ければ教えて下さい。
13 ：ハッテン場五段：2016/04/14 04:08:59 (8年前) 0.5MONA/1人
>>12
当社比 12% アップ版の Lyra2REv2.cl の 348 行目から 351 行目にご注目下さい。
変数 shift が取りうる値って最小が 0、最大が 36 です。
351 行目で読み込む DMatrix のデータ範囲は、j が取りうる範囲から考えても DMatrix[0] から DMatrix[38] になります。
だから初めは 0 <= s1 <= 38 でよかろうと思ったんです。
その名残が Lyra2v2.cl の 142 行目に残ってます。
でももうちょっとエレガントにやりたいなって思いまして。
変数 rowa、Lyra2v2.cl において言えば変数 rowInOut が取りうる値それぞれで場合分けして考えました。
rowa=rowInOut はそもそも ulong & 3 で計算されてますから 0 <= rowa=rowInOut =< 3 にしかなりえませぬ。
したがって Lyra2REv2.cl の 351 行目において DMatrix から読み込むデータの範囲は
14 ：ハッテン場五段：2016/04/14 04:09:20 (8年前) 0.5MONA/1人
rowa=rowInOut = 0 の時、DMatrix[0] から DMatrix[2] まで、
rowa=rowInOut = 1 の時、DMatrix[12] から DMatrix[14] まで、
rowa=rowInOut = 2 の時、DMatrix[24] から DMatrix[26] まで、
rowa=rowInOut = 3 の時、DMatrix[36] から DMatrix[38] まで
になります。一般化すると書き込むのは DMatrix[ rowa * 12 ] から DMatrix[ rowa * 12 + 2 ] であればよく、
s1 の範囲は rowa * 12 <= s1 <= rowa * 12 + 2 であればいい、と言えます。
んで関数内にはちょうど ps1 とかいう、計算に利用できそうな変数があったんで準用しております。
今気づいたんですけど、当社比 12% 版はもっと効率化できそうですね。
時間が許すかぎり頑張ってみたいと思います。
15 ：ハッテン場五段：2016/04/14 05:20:04 (8年前) 0MONA/0人
>>14 で書いた企ては見事に失敗しました。動くけど遅くなる！
16 ：リキプロマン六段：2016/04/14 20:09:03 (8年前) 0MONA/0人
とても詳しく解説して頂きありがとうございました。
そのshiftのとりうる値に合わせてDmatrixの読み込む大きさを変化させるわけですね。
逆に遅くなるのはなぜなんでしょうか。
自分もいろいろ考えていますが決定的な案が浮かびませんね…
17 ：名無し名誉名人教士：2016/04/14 21:19:48 (8年前) 0.1MONA/1人
>>16
GPUは分岐が苦手だから…
18 ：ハッテン場五段：2016/04/18 07:05:36 (8年前) 0MONA/0人
>>17
その一言に尽きますな
19 ：ハッテン場五段：2016/04/21 07:50:10 (8年前) 0MONA/0人
条件分岐をなくそうと bitselect 関数を使おうとするとなぜか OS が固まるこの頃、皆様イカがお過ごしでしょうか。なかなかうまくいきません。つか地震ばっかでそれどころじゃないです。すみません。
20 ：ハッテン場五段：2016/04/21 19:19:26 (8年前) 0MONA/0人
当社比17％アップ版ができましたよー イクイクー
http://kernel.hattenba.xyz/kernel-20160421.zip
でもこれ開発マシン(OpenCL2.0)とは別マシン(OpenCL1.2)で動かしてみたら計算速度が４割ほど落ちました。環境を選びます。ドライバが新しくないとダメなのかな？別マシンは Vista なんで古いドライバを未だに使ってます。
21 ：ハッテン場五段：2016/04/21 19:26:41 (8年前) 0MONA/0人
今回の変更は大きく分けて２つ。
一つはちょこっとお手入れして、掛け算の回数を減らしました。掛け算や割り算の回数を減らすのは計算量減らしの基本っすもんね。昔学校で習った。でもこれじゃ全然早くなってないんすよね。おまけ状態です。ププ
もう一つは激遅グローバルメモリに居座る DMatrix 配列を関数でカプセル化して、データの一部をプライベートメモリに移したことです。条件分岐の回数が格段に増えましたが、それでも以前よりも更に５％速度が向上しました。
本当は DMatrix 配列のデータは全部プライベートメモリに移したいのですが、皆様御存知の通り、Radeon ではコアあたりのメモリ量が全然足りません。なんで一部だけ。
22 ：siv三段：2016/04/21 22:31:21 (8年前) 0.1MONA/1人
>>20のファイルをnicehash minerのsgminer5.1.0-optimizedのkernelに入れてLyra2REv2回してみましたが特に変化は見られなかったですね(nicehash→6.810MH/s >>20 6.795MH/s)
自分の試し方が悪いのかもしれません
Radeon HD7970(16.150.2211.0)
23 ：ハッテン場五段：2016/04/21 22:55:32 (8年前) 0MONA/0人
>>22
ご報告ありがとうございます。もしかしたら環境にすごく左右されることばっか俺はやってるのかもですね。
Lyra2v2.cl の 71 行めにある DBUFLEN の大きさを変えると多少違うかもしれません。 0 から 48 の間で選べます。うちの開発マシンでは 4 が最適でした。
24 ：ハッテン場五段：2016/04/22 18:50:36 (8年前) 0MONA/0人
R9 290x + OpenCL1.2 なマシンで動かしてみたら、DBUFLEN はどうも 32 が最適みたいな感じでした。環境によってかなり数字が変わるなー
25 ：名前はまだ無い四段：2016/04/24 17:12:23 (8年前) 1.1MONA/2人
Lyra2v2部分のメモリ使用量が多い問題は、複数のスレッドで1つのハッシュを計算するようにすれば
スレッドあたりのメモリ使用量が減って回避できるように思えます。
4スレッドで1つのハッシュの場合は、64スレッドで16のハッシュを計算となり
DMatrixが使用するメモリは1536x16の24576バイトなので、CUあたり32kBのLDSに納められそうです。
steteやDMatrix等がulong4のベクトル型になっているので、それらをulongのスカラー型として4スレッドで処理するのはどうでしょうか？
26 ：ハッテン場五段：2016/04/24 19:31:18 (8年前) 0MONA/0人
>>25
うーむ、すみませんが、ulong4 を ulong の配列に変えることと、４スレッドで動かすことの関連性がよくわからないです。申し訳ありません。
APU はともかく、GPU はベクトル型の扱いは苦手ですから ulong4 を ulong の配列に変えてしまえば多少早くなると思います。４スレッドで動かせばメモリアクセスの空白期間が少なくなってムダが減るのもわかります。でもその２つの関連性がよくわからないです。
ちなみに俺はいつも -g 4 で動かしてまして、当社比の数字も -g 4 を指定した時の数字を使っています。
27 ：izuna五段：2016/04/24 19:51:15 (8年前) 0MONA/0人
>>25
ハッシュの計算って逐次で並列性の少ない計算だと思うのだが
並列性のないハッシュの計算を４スレッドで計算できるの？
28 ：名前はまだ無い四段：2016/04/25 23:41:47 (8年前) 1MONA/1人
>>26
ホスト側のスレッドではなくデバイス側のスレッドです。
GCNアーキテクチャではCU毎に64本の演算パイプがあり、WORKSIZEを64にして64スレッド単位で実行するのは良いのですが
Lyra2v2部分では1スレッドが1つのハッシュを計算するというストレートな実装ではスレッドあたりのメモリ使用量が多くなり色々と大変です。
4スレッドで1つのハッシュを計算し、それを4回繰り返すようにすればスレッドあたりのメモリ使用量が1/4となり
キャッシュにヒットしやすくなったりLDSに納められるようになったりします。
（DMatrixの配置の関係で同一ワークグループに属するスレッドのグローバルメモリへのアクセスが
複数回の命令になったりキャッシュヒット率が下がっていたりするようにも思いますけど。）
>>27
ハッシュの計算は並列性が低いことが多いと思いますが、Lyra2v2の場合はulong4のコピーやXOR、加算等が多いようですし、
round_lyraも前半と後半に分ければそれぞれ4並列で行っても問題ないと思います。
29 ：izuna五段：2016/04/26 00:28:57 (8年前) 0MONA/0人
>>28
> ハッシュの計算は並列性が低いことが多いと思いますが…
XORや加算が多いからといって、並列性が高いとは限らない。
4スレッドで2倍の性能がでても、最初から性能が半分になっている。
もう少し、具体的に実装を考えてみないと、どうとも言えない。
Lyra2REv2で使われるハッシュ関数の、ほとんどは、SHA-3候補のなんだけど、これらを4スレッドで実行するの？
30 ：ハッテン場五段：2016/04/26 02:48:06 (8年前) 0MONA/0人
>>28
よくわかんないんで実際に sgminer の OpenCL カーネルいじって見本みせて頂けますか？
31 ：ハッテン場五段：2016/04/28 08:03:05 (8年前) 0MONA/0人
前に配ったカーネルの Lyra2v2.cl の reduceDuplexRowf2 関数内で変数 max を使った条件分岐がいくつかあります。が、手で動きをトレースしたところ、それら条件分岐は必ず真または必ず偽にしかなりようがないことが分かりました。
で、条件文を取り去って動かしたのですが、速さはまったく変わりませんでした。というわけで訂正したものの配布は必要なさそうなので行いませぬ。あしからず。
32 ：ハッテン場五段：2016/05/01 05:59:51 (8年前) 1.114MONA/2人
当社比 25% アップ版ができました。イクイクー
http://kernel.hattenba.xyz/kernel-20160501.zip
・グローバルメモリへのアクセス回数を減らしました。
・ムダな条件分岐を取り払いました。
・あと何か変えたような気がします。
うちの R9 380 では 25% アップしましたが、R9 290x では 14% しか上がりませんでした。あしからず。
33 ：なむやん七段教士：2016/05/01 22:37:22 (8年前) 0MONA/0人
１０％以上アップとかすごすぎ
R7 250無印だとどんだけ上がるんだろう？
34 ：ハッテン場五段：2016/05/02 04:28:02 (8年前) 0MONA/0人
>>33
もし試す機会がきたら動作報告よろしくお願いします！
35 ：名無し名誉名人教士：2016/05/02 16:43:20 (8年前) 0.1MONA/1人
Radeon R9 290Xでみると、
1CUあたり、(スカラユニット×16)×4基に対して、レジスタファイル(Vector Refisters)が64kB×4基…
スカラユニット1つあたり4kBも確保できるから、ローカル変数で1536バイト確保できちゃう？
素人考えですまぬ…
36 ：ハッテン場五段：2016/05/03 04:51:33 (8年前) 0MONA/0人
>>35
ローカルメモリ！
sgminer の本体のほうを手入れしないとできませんが、幸いにも何ヶ月か前に自分でコンパイルしてたりしたのが残ってまして再利用すれば試せるかも？です。
37 ：ハッテン場五段：2016/05/03 04:53:29 (8年前) 0MONA/0人
問題は、./autogen.sh とか ./configure に色々オプションをつけて動かしてたのは覚えてるんですが、それらをすっかり完全に忘れてしまってることです。どこかにも書いてもないみたい。ぐぬぬ
38 ：名無し名誉名人教士：2016/05/03 06:10:49 (8年前) 0MONA/0人
>>37
シェアードメモリでも速くなるかな？
(OpenCLにシェアードメモリってあったっけ？)
39 ：名無し名誉名人教士：2016/05/03 06:15:21 (8年前) 0MONA/0人
ちなみに、ccminerのほうでreduceDuplexRowf2を試してみたけど、逆に遅くなった…
if分で相当遅くなるのかな？
(GTX970で改造前9.5MH/s、改造後8.5MH/s)
いま、if(rioneq2)無しで試してみる。
40 ：名無し名誉名人教士：2016/05/03 07:13:38 (8年前) 1MONA/1人
ccminerでreduceDuplexRowf2を試した結果…(GTX970で実験)
･ソースコードの通りに組んでみた場合
→改造前9.5MH/s、改造後8.5MH/s
･reduceDuplexRowf回り、unrollを見直し
(for文復活、最終のみreduceDuplexRowf2を使用)
→改造後10.0MH/s
･if(rioneq2)による分岐を排除
→改造後10.4MH/s
不用意なunrollは速度低下を招く…
それにしても、最適化ぱねぇ…
41 ：ハッテン場五段：2016/05/03 07:34:01 (8年前) 0MONA/0人
sgminer をビルドしなおして、グローバルメモリからローカルメモリに移行して動かしてみたんですが・・・ダメダメでした。どうもメモリが足りなすぎるようです。
-w 16 くらいなら動くんですが、すごく遅い。メモリ割り当て量や割り当て方を再検討してみます。
42 ：ハッテン場五段：2016/05/03 07:36:03 (8年前) 0MONA/0人
>>40
あの分岐は鬼門だったんですなｗ
43 ：ハッテン場五段：2016/05/03 09:43:52 (8年前) 0MONA/0人
色んなサイト見てローカルメモリについて分かったこと
・ローカルメモリは L2 キャッシュに置かれる
・びりあるさんが twitter にて 2015 年 10 月 13 日に「1.5KB くらいしか使わないので、普通に L2 キャッシュあたりに載っちゃう量」だと発言している
てことはやっぱりローカルメモリを使うこと自体は正解で、俺のやり方が間違ってるのかな。再検討！
44 ：ハッテン場五段：2016/05/03 10:33:31 (8年前) 0MONA/0人
clGetDeviceInfo に CL_DEVICE_LOCAL_MEM_SIZE を指定してローカルメモリの容量を調べたところ、R9 380 も APU も 32KB だと分かりました。Lyra2v2 の計算で使う 1536 バイトで割ると 21.33。
試しにローカルメモリを使った版で -w 21 を使うと遅いながらも一応ちゃんと動くのですが -w 22 にすると動作しなくなりました。ローカルメモリを使うとなるとワークアイテムは 21 よりも増やせないようです。
うーん、どうやって高速化させればいいんだろう？ スレッドたくさん作る？
45 ：名無し名誉名人教士：2016/05/03 20:33:03 (8年前) 0.1MONA/1人
>>32
ccminerでは、dbufによるメモリ分岐の最適化は逆効果でした…
GTX970で、改造前10.4MH/s、改造後8.5MH/s
Radeonでは効果はありましたか？
>>43
1536バイトを全部バラの変数にして(レジスタのみ使用しメモリを使用しないで)演算、関数も(引数はメモリを使用するから)inline化すれば…
かなり面倒くさいな…
46 ：ハッテン場五段：2016/05/04 04:04:55 (8年前) 0MONA/0人
>>45
うちの Radeon どもでは数％ですが効果ありましたよー どうも Radeon では条件分岐にかかるコストよりもグローバルメモリへのアクセスにかかるコストのほうが若干大きいようです。
ゲフォはメモリにもっと余裕がありそうなんで、DBUFLEN を増やしてみてもいいかもです。
1536 バイトを全部バラしますかｗ 手間が恐ろしいですｗ
でも実際いじってみて感じたんですが、同じ量のメモリを確保するのでも、でかい配列を一度に確保するよりは、小出しに確保していくほうがプログラムがすんなり動いてくれますね。まるで尿道結石のようです。
47 ：ハッテン場五段：2016/05/04 04:05:57 (8年前) 0MONA/0人
１ワークアイテムで使えるローカルメモリ量について。
ローカルメモリが x バイトあって ulong4 一つで 32 バイト使うから、１ワークグループで ulong4 が x/32 個使える。ということは sgminer のオプション -w で指定したワークサイズを n とすると、１ワークアイテムで使えるローカルメモリ上の ulong4 は x/(32*n) 個。
うちの R9 290x も R9 380 も APU どももどいつもこいつもローカルメモリの量は 32KB という表示が。ならば１ワークアイテムで使えるローカルメモリ上の ulong4 は
・ワークサイズ n が 16 の時は 64 個
・ワークサイズ n が 32 の時は 32 個
・ワークサイズ n が 64 の時は 16 個
・ワークサイズ n が 128 の時は 8 個
・ワークサイズ n が 256 の時は 4 個
とまあ結局はお情け程度の量の ulong4 しか使えないわけで。ちなみに ulong4 は 48 個必要です。これじゃキャッシュとかバッファくらいにしか使い物になんないかもですね。無いよりましか。
48 ：ハッテン場五段：2016/05/04 11:12:54 (8年前) 0.114MONA/1人
当社比 36 ％アップ版ができました イクイクー
http://kernel.hattenba.xyz/kernel-20160504.zip
今回は要注意！
OpenCL 2.0 上での R9 380 では 36 ％くらいアップしましたが、OpenCL 2.0 上での APU や OpenCL 1.2 上での R9 290x では 15 ％ほどの減になりました。
25 ％アップ版で使ったような、プライベートメモリを使ったグローバルメモリ読み書き回数減らしを更に適用させていきました。
49 ：しやく七段錬士：2016/05/04 12:24:41 (8年前) 0.1MONA/1人
トピの内容は難しくてさっぱりですが、速度向上等々ありましたのでご報告とお礼。
今のところ >>32 版が一番あってる、うちの7950君(1.1M → 1.4M)
>>9 のときには 1.25M位でした。
>>48 は速度低下(1.4M → 1.2M)してしまいましたので戻しました。
こういった最適化による改善の余地があるなんて、そしてそれを施しているトピ主さまはすごい！
ありがとうございます～
50 ：名無し名誉名人教士：2016/05/04 23:43:27 (8年前) 0.1MONA/1人
ccminerで1536バイトをローカル変数(レジスタ)で確保してやってみた。
手順：
･uint2でローカル変数を384個確保する
･関数をすべてinline展開
･すべてのfor文を手動unroll
･DMatrixをローカル変数(384個)に変更する
改造前10.4MH/s→改造後3.6MH/s
あれぇ…
51 ：ハッテン場五段：2016/05/05 04:51:01 (8年前) 0MONA/0人
>>49
テストご協力ありがとうございます！俺は信者ですから広めるのは当然なのです。
>>50
勇者現る！
速度減は残念ですが、何か別のいいことにつながるといいですね
uint2 をお使いのようですが、そちらのデバイスでは uint では 2 が最適なベクトル長なんでしょうか？
52 ：名無し名誉名人教士：2016/05/05 05:21:45 (8年前) 0.1MONA/1人
>>51
よくわからないが、uint2を4個でstruct→uint28として使われていたな…
…ってか、そもそもCUDAでuint2の加算に
"add.cc.u32 %0,%2,%4;"
"addc.u32 %1,%3,%5;"
とあったので、ひょっとして32ビットで作らなきゃダメ？
(そもそも、asm使っているからccminerは速いのか？)
53 ：名無し名誉名人教士：2016/05/05 16:37:53 (8年前) 0MONA/0人
DMatrixの読み込みを最適化できる？
並列スレッドでまとめて読み込めれば速くなるはず…
現状
DMatrix[0][1][2]… (thread_0用)
DMatrix[0][1][2]… (thread_1用)
DMatrix[0][1][2]… (thread_2用)
改善案
DMatrix[0]：(thread_0)(thread_1)(thread_2)…
DMatrix[1]：(thread_0)(thread_1)(thread_2)…
DMatrix[2]：(thread_0)(thread_1)(thread_2)…
(以下略)
54 ：名無し名誉名人教士：2016/05/05 17:07:33 (8年前) 0MONA/0人
>>53の方法でやってみたら、思いのほか速くなったでござる…
(ccminerでやってみた)
GTX970：改造前10.4MH/s→改造後21MH/s
…ぱねぇ！
でも、コードミスったのかエラーがチョイチョイ出る。ちょっと見直し中。
55 ：なむやん七段教士：2016/05/05 17:08:36 (8年前) 0.00000758MONA/1人
ふぁ！？倍！？打ち間違い？
56 ：名無し名誉名人教士：2016/05/05 23:08:36 (8年前) 0MONA/0人
>>54の最適化結果
数値がよかったのはエラーのせいだったようだ…
方針はいいと思ったんだけどな…
57 ：名無し名誉名人教士：2016/05/06 01:21:35 (8年前) 10MONA/1人
シェアードメモリで768バイト、グローバルメモリで768バイト
このパターンで作ってみた。
(ccminerでやってみた)
GTX970：改造前10.4MH/s
グローバルメモリ768バイトならL2キャッシュに収まる！
ここまで引き出せれば十分でしょ！
ハッテン場 様、参考までにソースをうｐしておきます。(CUDAですが)
https://onedrive.live.com/redir?resid=C7ABE390AB1575E7!782&authkey=!AI84vQzotFa0_bQ&ithint=file,zip
58 ：名無し七段教士：2016/05/06 02:12:58 (8年前) 1MONA/1人
当方環境 (R9 390, CL2.0.6.0, ドライバー16.3.2)においては以下のようになりました。4ケタ数字はzipの日付
ドライバーのせいでしょうか、結構差が出ました。
0413(7.4MH/s) > 0409(7.0MH/s)≒ Original(7.0MH/s)>> 0421(5.8MH/s) >> 0501(4.9MH/s)>> 0504(3.9MH/s)
59 ：ハッテン場五段：2016/05/06 03:24:14 (8年前) 0MONA/0人
>>57
素晴らしきプリケツ！ もうちょっとして手が空いたらじっくり鑑賞します！
>>58
うちはまだ 15.12 使ってます。つーか 16.3.2 に一度アップグレードしたんですが、sgminer がうまく動いてくれなくて戻しちゃいました。新しいドライバだと前のやつのほうがうまく動くんですね。ご報告感謝です！
60 ：ハッテン場五段：2016/05/06 08:46:20 (8年前) 0MONA/0人
>>57
じっくりねっとりと拝見しました。素晴らしいプリケツです。俺が書いた、関数を使ったバッファ？とは比べようもないエレガントさでした。
ただこちらの L2 キャッシュはどうも 32KB くらいしか確保ができないようで、１ワークアイテムあたり 768 バイトも確保するのは非常に難しそうです。せめて 384 バイトでも確保できればいいのですが。256 バイトがせいぜいかな。頭がいたい
61 ：名無し名誉名人教士：2016/05/06 17:23:34 (8年前) 0MONA/0人
>>60
RadeonはTonga、Hawaii、FijiにはLocalDataShareMemoryが64kBあるようですが、使えないのでしょうか？
62 ：名無し名誉名人教士：2016/05/06 23:15:02 (8年前) 10MONA/1人
>>57をさらに最適化してみた。もう、ゴールしてもいいよね…
シェアードメモリのバンクコンフリクトが面倒くせぇ…
ハッテン場 様、参考までにソースをうｐしておきます。(シェアードメモリの参考に)
https://onedrive.live.com/redir?resid=C7ABE390AB1575E7!783&authkey=!ACvnu6myOBoE8ME&ithint=file,zip
63 ：名無し名誉名人教士：2016/05/06 23:21:18 (8年前) 0MONA/0人
ちなみに>>53の最適化案はL2キャッシュがない場合は有効だと思うが、
シェアードメモリを併用したため、L2に入りきるサイズになったから、効果はないかもね…
64 ：名無し名誉名人教士：2016/05/06 23:25:21 (8年前) 0MONA/0人
冷静に考えると、sgminerを最適化するなら、
同じやり方でccminerも最適化できるから、差が縮まらないんじゃね？
ってか、ccminerはCUDAのインラインアセンブラまで手を出しているから、そう簡単に差は縮まらないと思われる…
65 ：名無し名誉名人教士：2016/05/07 03:19:16 (8年前) 0.1MONA/1人
>>62の最適化で1Block掘り当てた！！
66 ：名無し名誉名人教士：2016/05/07 03:22:49 (8年前) 0.1MONA/1人
ついさっき、もう1block掘れた！！
67 ：ハッテン場五段：2016/05/07 07:14:42 (8年前) 0MONA/0人
>>61
Tonga、Hawaii ともに試しましたが、32KB が限度のようです。むむむ。64KB 確保しようとすると sgminer 止まっちゃいます。
>>62
ありがとうございます！ よく分からん部分が結構ありますが、努力します。
>>63
OpenCL で言えばスレッド化はサブデバイス化にあたるのかなー サブデバイスはノータッチだったんでもっと勉強しないと
>>64
こないだ見せていただいたやつは、俺にとって初めて見る CUDA ソースだったんですが、アセンブラっぽいの見て「あ、これずるい！」と思いましたｗ
確かに差は縮まらないかもしれませんが、ケツを貸しあった者同士の距離は縮まったんじゃないかなと思いますぜ
68 ：名無し名誉名人教士：2016/05/07 16:49:11 (8年前) 0.1MONA/1人
>>67
sgminerのソースを見てみたら、--worksizeはデフォルトだと256らしい………いくらなんでも大きすぎないか？
ちなみにccminerではデフォルトで10程度。(最適化で16に変更。シェアードメモリ96kB÷4ブロック/SMM÷16スレッド/ブロック＝1536バイト。余裕を見てその半分を確保)
ワークグループ当たり32スレッド(--worksize 32と指定する)なら、32kB÷32＝1kBまで確保できる？
1536バイトの半分(768バイト)を確保するくらいなら何とかなりそうだが…
69 ：リキプロマン六段：2016/05/07 20:03:46 (8年前) 0MONA/0人
>>62
Ubuntu(gcc5.2 + CUDA 7.5)の環境で走らせてみましたが、どうもうまくいきません。
_ptr[0] = { data.x.x, data.x.y, data.y.x, data.y.y };
_ptr[stride] = { data.z.x, data.z.y, data.w.x, data.w.y };
の行がどうも怪しいようですが・・・
70 ：名無し名誉名人教士：2016/05/07 20:21:54 (8年前) 0MONA/0人
>>69
ccMiner 1.5.77-git Fork by SP for Maxwell
のソースを使ってやっています。(たぶん1.5.79でも大丈夫)
Windowsではちゃんと動いています。
Maxwellしか動かないので、別のforkでは書き換えが必要と思われます。
(そもそもKeplerではシェアードメモリが足りない)
それと、GTX970とGTX960では簡単なチューニングは行いましたが、その他のGPUでは適正値がわからないので、動かない可能性もあります。(特にGTX750Tiなど、compute capability 5.0は動作未確認)
71 ：リキプロマン六段：2016/05/07 21:29:05 (8年前) 0MONA/0人
>>70
動作環境はGTX 980です。ソースもccMiner 1.5.79を用いています。
>>57 のソースではコンパイルも動作も正常でした。
エラー文は、
lyra2/cuda_lyra2v2.cu(41): error: expected an expression
lyra2/cuda_lyra2v2.cu(42): error: expected an expression
でした。
今、gccとiccでのコンパイルの差がないか資料を探しています。
それと、lyra2v2で実行されるcubuhashとkeccakを高速化出来たのですが、これって需要ありますかね・・・
72 ：名無し名誉名人教士：2016/05/07 21:49:28 (8年前) 0MONA/0人
>>71
これで解決できればいいけど…(結局原因がわからない…)
http://d.hatena.ne.jp/aont/20110504/1304498046
ちなみに、http://askmona.org/3853にて、ソロマイニング用ではありますが、Windows向けバイナリをアップしています。
73 ：リキプロマン六段：2016/05/07 22:35:52 (8年前) 0MONA/0人
>>72
このエラーが発生している箇所がデバイス関数なので、cppには書くことが出来ないので無理ですね・・・
Windows+VisualStudio環境も用意しているので、こちらでコンパイルしてみますね。
74 ：名無し七段教士：2016/05/07 23:41:48 (8年前) 1MONA/1人
>>59
ちなみに、今更ですが
当方もバージョンあげて動かなくなりましたが、
%TEMP% に使用.clファイルコピペで解決しました
75 ：ハッテン場五段：2016/05/08 05:28:35 (8年前) 0MONA/0人
>>68
たしかに大きすぎるんですけど、64 とか 128 とか 256 じゃないとまともな速度で動いてくれんのです。
>>74
なんとそんな方法が！見るディレクトリが変わっちゃったのかな
76 ：ハッテン場五段：2016/05/08 05:32:39 (8年前) 1MONA/1人
ローカルメモリ活用のため、改造した sgminer を公開します。
http://kernel.hattenba.xyz/sgminer-lyra2rev2-only-20160508.zip
Windows 用バイナリもあります。多分 32bit
http://kernel.hattenba.xyz/sgminer_windows.zip
当社比 51% アップなんですが、やってることはこれまでと変わらないので、これまで速度が向上しなかった環境ではやはり向上しないと思われます。色々試しちゃいるんですが、今のところうまくいったのはこれだけなんです。すみません
77 ：名前はまだ無い四段：2016/05/09 02:21:17 (8年前) 1MONA/1人
>>29
4スレッドで1つのハッシュをというのはLyra2v2の部分の話です。
たとえばulong4の
state[j] ^= state1[j];
は
state[j].x ^= state1[j].x;
state[j].y ^= state1[j].y;
state[j].z ^= state1[j].z;
state[j].w ^= state1[j].w;
なので4並列で行っても問題ないと思います。
78 ：名前はまだ無い四段：2016/05/09 02:47:21 (8年前) 0.1MONA/1人
>>30
遅くなりましたが、今はテストに良さそうな環境が空いていませんし、テスト環境無しでさらっと書く自信は無いです。
（かなり手間がかかりそうで気が向かないというのが本音かもしれません）
流れとしてはlocal_idを取得してそれを元にどの入力のハッシュを処理するか
とulong4のどの次元を担当するかを決定する感じでしょうか。
ただ、AMDの資料を見てみるとwavefrontのCUでの実行のされ方などの関係もあり
1つのworkgroupでLDSを使い過ぎると効率が下がりそうですし、グローバルメモリへのアクセスが避けられそうになく
グローバルメモリでのデータ格納方法を見直す方が優先順位は高そうです。
64bitのデータではコアレスアクセスにならないという記述があったり、
分岐粒度のこともあったりとRadeonにとってはかなり不利なアルゴリズムの気がしてきました。
79 ：名無し名誉名人教士：2016/05/09 21:34:16 (8年前) 3.114114MONA/1人
>>77
あれ？これってひょっとして…
32スレッド÷4並列×1536バイト/スレッド＝12kBだから、案外、シェアードメモリ(OpenCLではローカルメモリ)が使えるんじゃね？
--worksizeが128の場合、128スレッド÷4並列×1536＝48kBとなり、1536バイトの半分程度ならローカルメモリで確保できそうだ。
…それだと、CUDAの場合は1536バイトの全領域がシェアードメモリで確保できるな。結局差が縮まらないわけか…
80 ：名無し名誉名人教士：2016/05/11 15:11:09 (8年前) 1MONA/1人
>>76
高速化乙です。
ローカルメモリのバンクコンフリクトが起こっているようなので、念のため。
ローカルメモリを確保した場合、アクセスは32バンクで行われます(メモリの窓口と考えると分かりやすい)
バンク当たり4バイト(8バイト？)アクセスするとして、long4の読み込みには32バイト使うため、1スレッドで8バンクにアクセスします。
つまり、32÷8=4スレッド同時にアクセスできます。
しかし、データは各スレッドでldbuflen個分、連続して使用しているため、たとえば、ldbuflen＝4の場合、local_id(0)のldbuf[0]～ldbuf[4]でバンク0～バンク31を占有し、local_id(1)以降のldbuf[0]と同一バンクになってしまいます。(つまり、ldbuf[0]を読むためにはバンク0～バンク7を各スレッドが同時にアクセスし、バンク8～バンク31が開いた状態になる。)
対策として、ldbuf[0]～ldbuf[4]と並べて確保するのではなく、「ldbuf[0]スレッド数分確保、その続きにldbuf[1]を…」とすると上手くいくはず。
また、ldbuflenが奇数になる場合、1スレッドずつバンクがズレるため、バンクへのアクセスが分散されます。(たとえば3の場合、local_id(0)がバンク0～7、local_id(1)がバンク24～31、local_id(2)がバンク16～23、と少しずつバンクがズレていく)
81 ：名無し名誉名人教士：2016/05/11 15:30:59 (8年前) 1MONA/1人
>>80 つづき
ちなみに「32バイトで読み込む」のではなく、「8バイトで4回読み込む」にすると、
バンクコンフリクトは1/4となって同時アクセスは少なくなりますが、読み込み回数が4倍になるため、
結局かわらないという微妙な状況になってしまいます。
(OpenCLの方で32バイトを分割読み込みしている場合を除く。CUDAでは「16バイト2回読み込み」のため、バンクコンフリクトは1/4、読み込み回数2倍、となり、一応メリットはありました)
82 ：ハッテン場五段：2016/05/12 03:57:32 (8年前) 0MONA/0人
>>78
テストならすぐできます。sgminer に添付されてる cl ファイルをテキストエディタでいじるだけですよ。
>>80
サンクスであります！
何かしらの単位ごとにメモリ配置をバラしていくんですね。やはりそこに手をつけずに済むことはありえない・・か！
ここ２，３日は SVM が使えないか奮闘しております。これが使えれば激速になるはず！でも SVM 使おうが使うまいが手元のローカルメモリの扱いも避けて通れそうにないので両方ともぼちぼちやっていきます。
GW 明けて仕事忙しくなってきて死にそうです。あと鯖たべたい
83 ：ハッテン場五段：2016/05/15 11:06:24 (8年前) 6.114MONA/3人
はーい 当社比 100 ％アップ版ができましたよー イクイクー
http://kernel.hattenba.xyz/kernel-20160515.zip
改造版の sgminer でしか動きませんが、ローカルメモリは使ってません。とりあえずグローバルメモリをどうにかするのが先じゃね？っと思いまして、名人たちからご教示いただいた方法を使ってみたら、２倍近くの速度がでるようになりました。ただし r9 290x だと 1.5 倍くらいにとどまってました。
実はそういう方法があるのは分かってはいたんですが、それをやるとどうして早くなるのか理解できず、納得できなくて着手していませんでした。なるほどメモリのバンク絡みの話だったんですな。
ローカルメモリの使い方は十分注意がいるようなので、今回はとりあえずローカルメモリはまったく使ってません。下手に使うと速度が落ちちゃう。
84 ：なむやん七段教士：2016/05/15 11:08:40 (8年前) 0MONA/0人
ふぁーww
凄まじい性能アップですな！
85 ：ハッテン場五段：2016/05/15 11:08:41 (8年前) 0MONA/0人
SVM は sgminer を 64bit 化しないと使えないのでコンパイルを通すべく奮闘してたのですが、うまくいかないのでとりあえず既存の CL ファイル周りをいじったらうまくいったので公開に至りました。
86 ：名無し名誉名人教士：2016/05/18 11:33:18 (8年前) 0.1MONA/1人
>>83
高速化乙です。
ccminerで行った高速化のソースを上げましたので、参考にしてください。
>>77のとおり、Lyra2を4分割にしたら結構いけたので…
1ブロック64スレッドとすると、1536バイト÷(4分割)×(64スレッド)＝24kバイト(＋スレッド間のデータ受け渡しに若干必要)なので、ローカルメモリに丸ごと収まる？
ソースコード(Lyra2REv2部分のみ)
https://onedrive.live.com/redir?resid=C7ABE390AB1575E7!806&authkey=!ANaEIqlJ6vqyABw&ithint=file,zip
87 ：ハッテン場五段：2016/05/19 08:22:14 (8年前) 0MONA/0人
>>86
ありがとうございます。情報サンクスです。
OpenCL でスレッド化ってどうするのか、方法がないのかいまいちよく分からないです。サブデバイス化というのはカーネルの局所局所をスレッド分岐するという技術ではなさそうです。今試しにサブデバイス化実験をやってるんですが clGetDeviceInfo でコア吐き出すという無間地獄中です。デイバス分割しても使えるメモリの量が減らないんなら分割し甲斐もあるのですが、それを調べる方法を使うと固まっててはどうしようもない。頭が痛いです。
あっちのスレは賑やかですなー こちらは寂しい限りです (´Д⊂ｸﾞｽﾝ 無理も無いか
88 ：名無し名誉名人教士：2016/05/19 08:36:28 (8年前) 0MONA/0人
>>87
私は単純に、スレッド数4倍にして動かしています。サブデバイス化(？)ではないはずです。
89 ：ハッテン場五段：2016/05/20 03:03:18 (8年前) 0MONA/0人
>>88
コードを３０回くらい読みなおしてようやく話が見えてきました！
raund_lyra した後に state を８バイト分ずらして state2 に xor する処理について、別スレッドからデータをもってきたりしてるんですね。
元のコードとどうして等価になるのかまだよく分からないのでもうちょっと読みなおしてきます。
90 ：名無し名誉名人教士：2016/05/20 04:36:12 (8年前) 0MONA/0人
>>89
raund_lyraも別スレッドから持ってきていますよ～
91 ：ハッテン場五段：2016/05/21 05:35:28 (8年前) 0MONA/0人
>>90
！ 見落としまくりでした。
でやっぱりどうして以前のコードとスレッド化させたコードが等価になるのかよくわからないです。
元の中身がちがう、別スレッドの state の中身を手元にもってきてどうして計算できてしまうのだろうかとずっと考えています。
もしや 全てのスレッドに渡すもともとのデータはすべて同じなのでしょうか。いやいやもしそうだったらスレッド間でのデータやりとりするどころかまったく同じ計算をたくさんのスレッドでやらせる意味自体がないハズ。なんですが
92 ：ハッテン場五段：2016/05/21 05:56:42 (8年前) 0MONA/0人
いや！ もしやワークグループでは、元々与えられるデータは同一なのかな！？ もしそうであれば話は通るし、何も不自然でないし、これまで見てきた OpenCL や CUDA との解説とも矛盾しない！
93 ：名無し名誉名人教士：2016/05/21 10:10:20 (8年前) 0.2MONA/2人
>>91
__shfl()の解説です。
Data1 = __shfl(Data1, threadIdx.x + 1, 4);
これを例に考える。threadIdx.xはスレッドID。まあ、何番目のスレッドかを表すものと考えてください。
①自身のIDを計算します。(threadIdx.x)%(第3引数)
②各スレッドでData1を用意します(第1引数)。
③どのIDのデータを受け取るかを計算します。(第2引数と第3引数の剰余)
④最後に、③で指定したIDのデータを返します。(返値)
この動作をまとめると以下の通り。
自身のスレッドIDが0なら、スレッド1からData1を受け取る。
自身のスレッドIDが1なら、スレッド2からData1を受け取る。
自身のスレッドIDが2なら、スレッド3からData1を受け取る。
自身のスレッドIDが3なら、スレッド0からData1を受け取る。
94 ：名無し名誉名人教士：2016/05/21 10:14:59 (8年前) 0.2MONA/2人
>>93 つづき
自身のスレッドIDが4なら、スレッド5からData1を受け取る。
自身のスレッドIDが5なら、スレッド6からData1を受け取る。
自身のスレッドIDが6なら、スレッド7からData1を受け取る。
自身のスレッドIDが7なら、スレッド4からData1を受け取る。
__shfl(state[0], 0, 4)
のような使い方をしている部分がありますが、
スレッド0～3はいずれも、スレッド0からstate[0]を受け取る。
スレッド4～7はいずれも、スレッド4からstate[0]を受け取る。
スレッド8～11はいずれも、スレッド8からstate[0]を受け取る。
スレッド12～15はいずれも、スレッド12からstate[0]を受け取る。
ということになります。
95 ：ハッテン場五段：2016/05/21 20:04:35 (8年前) 0MONA/0人
>>93
ありがとうございます！
段々と頭がすっきりしてきました。朝日が上る日も近い！と感じてきます。
96 ：名前はまだ無い四段：2016/05/21 23:46:01 (8年前) 0MONA/0人
>>82
そのテストに良さそうな環境が空いていないのです。
>>88
4並列で行っている部分はブロック数を4倍にしているのですね。
Gfunc_v35ではuint64_tのデータをuint2として処理しているようですが、
加算を上位32ビットと下位32ビットに分けて行うことになり問題が発生しないのでしょうか？
97 ：名無し名誉名人教士：2016/05/22 01:00:21 (8年前) 0.1MONA/1人
>>96
実はuint2は64bit加算を行っているんですよ～
add.cc.u32 %0,%2,%4;
addc.u32 %1,%3,%5;
このようにキャリー付の加算を行って、32bit演算の繰り上がりを処理して、64bit演算と同じ結果が得られます。
98 ：名前はまだ無い四段：2016/05/22 17:02:40 (8年前) 0MONA/0人
>>97
桁あふれが反映されるのですか。
知らないとはまりそうですが、ulongをuint2として処理する場合は便利ですね。
99 ：名無し名誉名人教士：2016/05/22 18:10:10 (8年前) 0.1MONA/1人
>>98
ちなみにこの計算の意味は、
%0=%2+%4 →CC(桁溢れ)
%1=%3+%5+CC
となり、下位を%0、%2、%4で、上位を%1、%3、%5を計算すれば、64bit計算になります。
100 ：名無し名誉名人教士：2016/05/23 16:38:08 (8年前) 0MONA/0人
Monacoinウォレットがバージョンアップしたようです。
プールのウォレット更新がちょっと心配。
ソロマイニングはエラー報告が来ていたので、採掘方法が何か変わっている？
101 ：ハッテン場五段：2016/05/24 06:28:12 (8年前) 0MONA/0人
>>100
え、ソロまたなんか採掘やりにくくなったんですか？
うちのプールの monacoind を更新しました。今んとこ問題ないようですが、ブロック見つけてモナを配分できるのを確認するまでは気が抜けませんです
102 ：名無し名誉名人教士：2016/05/24 08:15:16 (8年前) 0MONA/0人
>>101
プールでは大丈夫って報告もあったんですが…どうなんでしょう？
103 ：ハッテン場五段：2016/05/25 06:25:02 (8年前) 0MONA/0人
>>102
うちのプールはちゃんとブロック発見できて、ちゃんと分配できたようです。分配つか俺一人しかいませんけどｗ
104 ：siv三段：2016/05/26 17:56:49 (8年前) 0MONA/0人
そういえば当初のハッシュレートと比べてどの程度向上したんですか?
具体的な値があるといいんですが...
105 ：ハッテン場五段：2016/05/29 04:36:36 (8年前) 0.1MONA/1人
>>104
現状うちのマシンの R9 380 で言えば
3.755MH/s → 7.650MH/s
くらいにはなっとります。
106 ：siv三段：2016/05/29 05:25:36 (8年前) 0MONA/0人
>>105
なるほど
初期と比べて劇的に改善されてますね!
107 ：ハッテン場五段：2016/05/29 06:49:47 (8年前) 0MONA/0人
>>106
独力ではここまでは絶対にムリでした。
皆さんから頂いたアイデアとモナとご声援の賜物です。
108 ：ハッテン場五段：2016/05/29 06:51:56 (8年前) 0MONA/0人
最近の技術的課題は、いかにスレッド間でデータをやりとりするか？ってところでつまずいてます。
CUDA には便利な命令があるようですが、OpenCL でとなると、ローカルメモリやグローバルメモリを使ってやり取りするほかないかな？と思ってます。でもこれがまたなかなかうまくいかないです (´・ω・`)
109 ：名無し名誉名人教士：2016/05/29 09:37:06 (8年前) 0MONA/0人
>>108
ローカルメモリでいいんじゃないかな？>>93のコードの疑似コードを(CUDAで)作るならこんな感じ…かな？
__threadfence_block(): //シェアードメモリのフェンス(データ上書き防止、シェアードメモリ未使用なら省略可)
uint2 buf=sharedmem[thread&31]; //シェアードメモリの退避(シェアードメモリ未使用なら省略可)
sharedmem[thread&31]=src; //渡すデータをシェアードメモリに書込
__threadfence_block(): //シェアードメモリのフェンス(データ上書き防止、「省略不可」)
uint2 dst=sharedmem[((thread&31)&~(WarpWidth-1))+(SelectID&(WarpWidth-1))]; //受取データをシェアードメモリから選択し読込
sharedmem[thread&31]=buf; //シェアードメモリの復旧(シェアードメモリ未使用なら省略可)
__threadfence_block(): //シェアードメモリのフェンス(データ上書き防止、シェアードメモリ未使用なら省略可)
return dst;
110 ：名無し名誉名人教士：2016/06/06 08:55:52 (8年前) 0MONA/0人
>>108
ハッテン場さん…高速化、難航してます？
なかなかこっちでの発言が無いようなので…
111 ：名無し四段：2016/06/13 04:14:31 (8年前) 0MONA/0人
>>74
ありがとう！
user（ユーザー名）￥appdata￥local￥temp￥の直下に
Kernelの中身置いたら、新しいドライバでも動くようになった。
おかげでハッシュレートも6.4MH/s（cata14.4）→9.4MH/sになったよ。
本気で動いてるから発熱も上がって75℃→86℃付近になって部屋が暑い
ちな、R9-290 4GB＠RADEON Software 16.6.1hotfix
112 ：hoshi五段：2016/06/19 02:40:16 (8年前) 0.00114114MONA/1人
久々にスレ一覧を見てたら素晴らしいスレで凄い高速化されてる・・・！！
自分も無駄な分岐を除くだけしてHD7950で5Mhs(15%向上?)くらいで回してましたが、もうとっくにそういった次元ではなさそうですね
rx480/rx470もそろそろ出るので期待してます！！
(Cはあんまり出来ないクソ雑魚ですが、もし役に立てそうであれば自分も協力します！)
113 ：きさらぎ八段錬士：2016/06/20 21:44:29 (8年前) 0MONA/0人
関係ない話ですがEthereumとDecredを同時に採掘できるマイナーは何で同時に採掘できるんでしょう？
monaとkumaを同時に採掘、とか出来ないものでしょうか
114 ：なむやん七段教士：2016/06/20 22:23:54 (8年前) 0MONA/0人
kuma!
クマはオススメですよ(意味深
115 ：名無し初段：2016/06/21 13:52:25 (8年前) 0MONA/0人
久々に参加してみたら
6950で1.7Ｍ うーむ
116 ：siv三段：2016/06/28 20:16:14 (8年前) 0MONA/0人
そろそろRadeon RX4xx発売されますね
ハッシュレートどれぐらい出るんでしょうか
117 ：ハッテン場五段：2016/06/29 06:44:30 (8年前) 59.63114114MONA/2人
ご無沙汰です。ハッテンです。難航というか、開発に取り掛かれない状況に陥っていましてみなさま申し訳ないです。５月末から時系列で申しますと・・
Vista を 7 にアップグレード！ドライバも更新して R9 290X が 7.5MH/s から 10.6MH/s に！ヒャッハー！
↓
地震保険おりた！保険金で monacoin 全力買い！
↓
不動産屋「そのアパート熊本地震で倒壊の危険があるから引っ越してね
↓
アパート片付け＆引っ越す金もないけどモナは売りたくないので１００円単位で金策
↓
ｱ゛ｱ゛ｱ゛ｱ゛ｱ゛ｱ゛ｱ゛ｱ゛！ ←今ココ
ただでさえ仕事忙しい時期なのに死にそうです すみません うちの汚部屋が手強いです
118 ：ハッテン場五段：2016/06/29 06:45:09 (8年前) 0MONA/0人
正直 RX 480 ほしいけどこんな状況じゃ買えないですね アハハ
119 ：名無し四段：2016/06/29 23:07:19 (8年前) 0MONA/0人
>>117
熊本っすかｗ近所かもｗ
こっちはあまり被害無かったけど
120 ：名無し四段：2016/07/03 00:57:23 (8年前) 0MONA/0人
RADEON RX460が75Wらしいので
常時起動して掘れるようになりたいなぁ。
ハッシュレート次第だけど
121 ：なむやん七段教士：2016/07/03 03:22:06 (8年前) 0MONA/0人
あの爆熱御殿常連のRADEONが75Wだと・・？
122 ：名無し四段：2016/07/03 09:09:48 (8年前) 0MONA/0人
GFの14nm FinFETだからね。
半分のサイズになれば、電力効率なら
2倍程度にはなるらしい。
実測で1.9倍だったみたいだが
123 ：siv三段：2016/07/05 19:31:55 (8年前) 0MONA/0人
そういやいつかしらにNicehashでRadeonのLyra2rev2ハッシュレート大幅に改善されたんだっけ
6月2日だったかしら
自分のHD7970のハッシュレートは現在12.5MH/s
124 ：リキプロマン六段：2016/07/09 12:59:19 (8年前) 0.114114MONA/1人
wolfさんがいろんなアルゴリズムに対して最適化したOpenCLのカーネルを配布したようです。
http://cryptomining-blog.com/8071-wolf-has-publicly-released-opencl-optimized-kernels/
Lyra2rev2に含まれているCubeHashやSkeinもあるので、これを取り入れればより一層高速化出来るのではと思います。
125 ：名無しさん：2016/07/17 07:49:12 (8年前) 0MONA/0人
Nicehashのバイナリ使ったらこんな感じですよ。
R9 390 一枚
i.imgur.com/zzNg11W
126 ：名無し四段：2016/07/18 02:40:22 (8年前) 0MONA/0人
>>123
なんかエラーでちゃう・・・。
127 ：名無し四段：2016/07/30 21:30:54 (8年前) 0MONA/0人
16.7.3にしたら、またハッシュレート上がった。
128 ：名無しさん：2016/08/01 03:25:31 (8年前) 0MONA/0人
マイニングでさ AWSのGPUは採算どうなんだろう？
129 ：名無し四段：2016/08/03 11:18:36 (8年前) 0.114114MONA/1人
Anniversary Updateしたら
コマンドプロンプトのフォントサイズ小さくされたし
Intel HD graphicsとRADEONの順番が逆になったのか
batのgpu-platform 1を0にしたら直ったよ。
130 ：七段教士：2016/08/25 21:40:28 (8年前) 0MONA/0人
https://github.com/nicehash/sgminer/releasesのv5.5.0を使っているのですが、
50Kh/sほどしか出ないです...
R9 390でbatはhttp://pastebin.com/iLduWSXBです
131 ：名前はまだ無い四段：2016/08/28 15:40:18 (8年前) 3.9MONA/1人
>>130
デフォルトではintensityが8とかなり低くなっているので、「-I」オプションで環境にあった値を設定する必要があります。
実行中でもメイン画面から「g」「i」で変更可能です。
132 ：七段教士：2016/08/28 23:12:48 (8年前) 0MONA/0人
>>131
ありがとうございます!
133 ：kenichiA380初段：2016/10/22 04:31:53 (8年前) 0MONA/0人
Sapphire R9 390X Nitro使ってます
Iは19以上にすると不安定
OCはコア1200 メモリー1700 電圧+75mV 電力制限は+50%
でハッシュレートは10.64MH/sなんですがこれは正常なんですかね
134 ：kenichiA380初段：2016/11/01 02:35:56 (8年前) 0MONA/0人
Radeon Crimson 16.10.3では掘れないみたいですね...
135 ：kenichiA380初段：2016/11/01 03:33:55 (8年前) 0MONA/0人
>>134
掘れました
136 ：ハッテン場五段：2016/11/25 04:12:04 (7年前) 0MONA/0人
久しぶりに閲覧
グラボ片っ端から売っちゃったり切手買っては転売したりしてて開発環境まるごとふっとんじゃいました。今あるのは R9 290x 一枚だけ。電気代もこれ以上増やせないくらいに生活が厳しいです。ごめんなさいです。
オークモナーの商品が売れてくれれば・・まあメシ代くらいにはなるか
137 ：名無し名誉名人教士：2016/11/25 08:42:42 (7年前) 1.14MONA/1人
>>136
こっちの最適化にも手を出してみたいけど、開発環境がイマイチ作れない…
コンパイルするにも一部ソースが欠けているみたい…
私にはsgminerは難易度が高いのだろうか…
(そもそも改造しかできない私はプログラマーとは呼べない半端者なわけだが…)
138 ：ハッテン場五段：2016/11/25 08:52:43 (7年前) 0MONA/0人
>>137
メモリ管理システムを根本から変えるならともかく、メモリへのデータの出し入れや計算順序を最適化するだけならコンパイラは不要ですぜ。
kernel ディレクトリにある cl ファイルをテキストエディタでいじって、sgminer を起動するだけです。起動時に cl ファイルをコンパイルしてくれます。
139 ：名無し名誉名人教士：2016/11/25 12:42:39 (7年前) 0.00114114MONA/1人
>>138
いや、三段打ちメソッド(命名、izuna氏？>>5、>>6)とか、ソロマイニングへの対応とか…いろいろやってみたいじゃない？
…私の技量でできるかはともかくとして…
140 ：Defekt名人四段錬士：2016/11/29 23:49:47 (7年前) 0MONA/0人
俺のGTX960Mは一晩1モナ
141 ：名無しさん：2017/05/02 04:43:58 (7年前) 0MONA/0人
ハッテン場様、
初めまして。突然のお願いで大変恐縮なのですが、ハッテン場様が
lyra2rev2カーネルに追加されたコードのライセンス販売、もしくは権利の
販売にご興味がおありでしたら以下のメールアドレスにご連絡を
いただけないでしょうか。
zawawawawawawawawa@gmail.com
以上、よろしくお願いいたします。スレ汚し大変失礼しました。
zawawa @ bitcointalk.org
142 ：光の名無し二段：2017/11/18 19:26:42 (7年前) 0.0114MONA/1人
mkxminerって新しいminerがリリースされてた
R290Xで32MH/sくらいでてます…＾ｐ＾
143 ：名無し名誉名人教士：2017/11/18 19:59:01 (7年前) 0MONA/0人
>>142
めっちゃ気になるが…
現状はnicehashしか使えないみたい…
[19:58:11] Started mkxminer v1.0.1
[19:58:11] Sorry, this release of software only works with nicehash.com pools. Try mining on stratum+tcp://lyra2rev2.usa.nicehash.com with your BTC address as login.
こんなメッセージが出た…
144 ：名無し名誉名人教士：2017/11/18 20:02:08 (7年前) 0MONA/0人
>>143
あ、でも、DNSサーバを弄れば一応できるかな？
lyra2rev2.usa.nicehash.comをvippool.netにaliasすればいい…のかな？
できるか知らんけど…
145 ：きら五段錬士：2017/11/18 20:10:21 (7年前) 0MONA/0人
hostsはどうだろうか
146 ：光の名無し二段：2017/11/18 20:22:40 (7年前) 0MONA/0人
hostsでやってますよー
147 ：名無し名誉名人教士：2017/11/18 20:29:44 (7年前) 0.049MONA/2人
>>145
私もそれに気づいてやってみた。
やり方は以下の通り。
①ping vippool.net を実行する。出てきたIPアドレスを記録しておく。
②C:\Windows\System32\drivers\etc (windows10の場合。他は知らん)内のhostsを編集する。
xxx.xxx.xxx.xxx lyra2rev2.usa.nicehash.com
を追記すればいい。
一旦デスクトップとかに保存して、その後、hostsを上書きする。(じゃないと権限が無くて保存できない)
③接続先をlyra2rev2.usa.nicehash.comとして起動する。ユーザー名、パスワードはvippoolのものを使用する。
これで一応マイニングできる。
148 ：名無し名誉名人教士：2017/11/18 20:30:46 (7年前) 0MONA/0人
>>147 補足
xxx.xxx.xxx.xxxは①で調べたIPアドレス
149 ：名無し名誉名人教士：2017/11/18 20:32:28 (7年前) 0MONA/0人
ちなみに、>>147はRX560でやってみた。
150 ：名無し名誉名人教士：2017/11/18 21:05:57 (7年前) 0MONA/0人
Radeon買ってこようかしらね…
151 ：名無し五段錬士：2017/11/18 21:51:45 (7年前) 0MONA/0人
>>149
うちもRX560なんですけど、ドライバ+MOD無しのノーマル仕様で13Mh/s 現状で日給0.2monaと健闘してますね
コイツはすげえぜ
152 ：名無し名誉名人教士：2017/11/18 22:20:59 (7年前) 0MONA/0人
ドライバを更新したら、RX560で12MH/s出た。
期待できる…のか？
153 ：名無し初段：2017/11/18 23:43:16 (7年前) 0MONA/0人
>>147の説明に勝手に補足。
メモ帳を右クリックして管理者権限で起動した後、Hostsファイルを開かないと、Hostsは上書きできません。
154 ：namo一級：2017/11/18 23:56:19 (7年前) 0MONA/0人
280XだったけどRX560買ってくるわ
155 ：名無し四段：2017/11/19 00:41:11 (7年前) 0MONA/0人
苦汁をなめ続けていたRadeon勢に光明が…
156 ：名無し一級：2017/11/19 01:12:37 (7年前) 0MONA/0人
>>147の③がよく判らないので教えてらえませんか？
初心者ですみません。
157 ：名無し一級：2017/11/19 01:15:31 (7年前) 0MONA/0人
>>156自己解決しました。
158 ：名無し四段：2017/11/19 01:16:44 (7年前) 0.39MONA/1人
>>156
hosts書き換えまでうまくいったんなら
stratum+tcp://lyra2rev2.usa.nicehash.com:【お使いのプールのポート番号】
ワーカーとパスワードは前マイナーの設定そのままでいけるはず
うちだとこういう設定
mkxminer.exe -o stratum+tcp://lyra2rev2.usa.nicehash.com:8888 -u 【ワーカー名】 -p 【ワーカーパスワード】 -I 20
159 ：名無し一級：2017/11/19 02:14:20 (7年前) 0MONA/0人
>>158
ありがとうございます。
動いたと思ったら実は全く動いてなかったんです。。。
160 ：光の名無し二段：2017/11/19 02:50:56 (7年前) 0MONA/0人
まあ290Xですし採算度外視なあそびではありますがこのカーブは気持ちいい
161 ：名無し名誉名人教士：2017/11/19 07:07:58 (7年前) 0MONA/0人
mkxminerのリンク（？）をとりあえず張ってみる。
https://bitcointalk.org/index.php?topic=2360168
162 ：名無し名誉名人教士：2017/11/19 08:21:52 (7年前) 0MONA/0人
mkxminerはローカルメモリを使って最適化していると思う。
その場合、コア数、クロック数の比で性能がある程度比較できる…のかな？
RX560:1024コア、1275MHzで12MH/s
これをベースとして…
RX580:2304コア、1340MHz⇒12MH/s×(2304/1024)×(1340/1275)＝28.38MH/s
RX480:2304コア、1266MHz⇒12MH/s×(2304/1024)×(1266/1275)＝26.81MH/s
RX570:2048コア、1244MHz⇒12MH/s×(2048/1024)×(1244/1275)＝23.42MH/s
RX470:2048コア、1206MHz⇒12MH/s×(2048/1024)×(1206/1275)＝22.70MH/s
RX460:896コア、1200MHz⇒12MH/s×(896/1024)×(1200/1275)＝9.88MH/s
RX550:512コア、1183MHz⇒12MH/s×(512/1024)×(1183/1275)＝5.57MH/s
(世代の違いから最適化具合が違うので、信憑性が…)
VEGA64:4096コア、1546MHz⇒12MH/s×(4096/1024)×(1546/1275)＝58.20MH/s
VEGA56:3584コア、1471MHz⇒12MH/s×(3584/1024)×(1471/1275)＝48.47MH/s
FuryX:4096コア、1050MHz⇒12MH/s×(4096/1024)×(1050/1275)＝39.53MH/s
FuryNano:4096コア、1000MHz⇒12MH/s×(4096/1024)×(1000/1275)＝37.65MH/s
163 ：名無し名誉名人教士：2017/11/19 08:22:10 (7年前) 0MONA/0人
300円/MONA、Diff40000、電気代30円/kWhと考えて、
RX580:28.38MH/s、185W⇒0.3548MONA(106.4円)、電気代133.2円
RX480:26.81MH/s、150W⇒0.3351MONA(100.5円)、電気代108.0円
RX570:23.42MH/s、150W⇒0.2928MONA(87.8円)、電気代108.0円
RX470:22.70MH/s、120W⇒0.2838MONA(85.1円)、電気代86.4円
RX560:12.00MH/s、75W⇒0.1500MONA(45.0円)、電気代54.0円
RX460:9.88MH/s、75W⇒0.1235MONA(37.1円)、電気代54.0円
RX550:5.57MH/s、50W⇒0.0696MONA(20.9円)、電気代36.0円
(繰り返すが、世代の違いから最適化具合が違うので、信憑性が…)
VEGA64:58.20MH/s、295W⇒0.7275MONA(218.3円)、電気代212.4円
VEGA56:48.47MH/s、210W⇒0.6059MONA(181.8円)、電気代151.2円
FuryX:39.53MH/s、275W⇒0.4941MONA(148.2円)、電気代198円
FuryNano:37.65MH/s、175W⇒0.4706MONA(141.2円)、電気代126円
164 ：名無し名誉名人教士：2017/11/19 08:24:58 (7年前) 0MONA/0人
>>162、>>163
一応計算したけど、なんか違うような…
どれにMonacoinに向いているか検証したいので、各GPUでの動作確認求む。
165 ：しやく七段錬士：2017/11/19 09:19:25 (7年前) 0MONA/0人
今更使ってる人いなそうなので参考にならないかもだけど書いとく。
7890 Sgminer 1.3M mkxminer 1.2M で変化なし
166 ：アポロ五段：2017/11/19 11:40:59 (7年前) 0MONA/0人
RX470 4GB コア1242MHz メモリ1750*4=7000MHz 約29MH/s
Radeonでこんなに掘れるとは
167 ：名無し名誉名人教士：2017/11/19 19:47:12 (7年前) 0MONA/0人
中古でRX480を買ってやってみた
VEGA64も買ってみたんだが、電源不足で起動せず…
さて、どうしたものか…(出張先のため、すぐには試せない…)
168 ：名無し名誉名人教士：2017/11/19 19:50:26 (7年前) 0MONA/0人
>>167 続き
RX480は中古27800円で33MH/s…
大量買いに走るか？
169 ：名無し初段：2017/11/19 20:10:34 (7年前) 0MONA/0人
あれー？うちのRX480だと起動もしない・・
"initscr():Unable to creaye SP"とかメッセージが出て落ちる・・
ドライバが古い(Crimson 17.3.3)からかなぁ？
皆さんやはりドライバは最新(17.11.2)？
170 ：アポロ五段：2017/11/19 21:23:51 (7年前) 0MONA/0人
17.10.1でした
171 ：たれぞう一級：2017/11/19 21:37:09 (7年前) 0MONA/0人
radeon勢の方はbiosいじってますか？
172 ：名無し初段：2017/11/19 21:44:57 (7年前) 0MONA/0人
細かい検証はしてないけど、
ドライバはベータマイニングドライバ
Vega64はPL-25%で47MH/sだった。（クロックは確認せず）
RX470はPL+10%とかにしないと22MH/s行かない感じに見えました。
mkxminerで-20%だと17MH/sでした。
Vega64とRX470とGTX1080+Ryzen1700のマイニングを同じマシンで動かしてるから、
PL上げるのがきついので、しばらくCryptoNightになりそう。
173 ：名無し初段：2017/11/19 22:46:38 (7年前) 0MONA/0人
ドライバを最新(17.11.2)にしたら、動くようになりました。
ただ、うちのRX480だとintensityは22、だいたい28MH/sという感じですね。
174 ：名無し一級：2017/11/20 00:35:56 (7年前) 0MONA/0人
うーん掘れてるみたいだけどVIPPOOLのほうを見ると反映されてない...
誰かタスケテ
175 ：だれかさん七段教士：2017/11/20 01:02:36 (7年前) 0MONA/0人
>>174
hostsファイル書きかえた？
176 ：名無し初段：2017/11/20 01:17:34 (7年前) 0.0039MONA/1人
>>174
とりあえずpingの結果とhostsの内容を実行前に表示するようにしては？
例えばこんな感じ。
@echo off
echo ----偽装ＯＫ？----
ping -a vippool.net | find "送信しています"
echo ----
type C:\Windows\System32\drivers\etc\hosts | find "nicehash"
echo ----偽装ＯＫ？----
pause
177 ：名無し名誉名人教士：2017/11/20 04:59:06 (7年前) 0MONA/0人
--gpu-powertune -40
を入れたら、VEGA64で電源不足にならなくなった。(まあ、安物電源だしね…)
-40%で48MH/s！？
178 ：めめんま('ω')初段：2017/11/20 05:40:15 (7年前) 0MONA/0人
検証おつかれさまです('ω')
133.0wはCurrent表示ですかね？
GTX1080のPL73で大体48～50Mh/sなのでワッパも匹敵するような
なんかVega64欲しくなってきた...あとは価格がもっとこなれてくれば..
179 ：名無し名誉名人教士：2017/11/20 09:27:43 (7年前) 0MONA/0人
>>178
133WはCurrent表示です。
試してはいませんが、ソフトウェアスイッチでパワーセーブ(165W)、バランス(220W)、ターボ(253W)とあり、初期状態はバランス(220W)みたいだね…それの-40%ってことで、133Wになっているっぽい…
いろいろ実験してみたいけど、現状の電源では無理…
180 ：めめんま('ω')初段：2017/11/20 11:03:28 (7年前) 0MONA/0人
>>179
有益な情報ありがとうございます。
なんかまだまだ伸びしろありそうですね～
近々グラボ追加しようと考えてるんですが、Vegaも1枚入れてみようかと本気で悩む...(゜_゜ )
181 ：名無し初段：2017/11/20 14:47:10 (7年前) 0MONA/0人
cast_xmr-vegaを動かす関係でドライバ最新じゃないので、参考程度で
ワットモニターで以下の条件で計測
①アイドルとマイニング時の差分でGPUの消費電力を推測
②HBCCはON（OFFとの差は測定してないので不明）
mkxminer+Vega64 MemClock：1025MHzのパターン
PL：+7% ハッシュレート：50.8Mh/s VGA消費電力：285W
MemClock：800MHzのパターン
PL：+7% ハッシュレート：55.1Mh/s VGA消費電力：285W
PL：+0% ハッシュレート：52.0Mh/s VGA消費電力：255W
PL：-10% ハッシュレート：47.9Mh/s VGA消費電力：230W
PL：-20% ハッシュレート：45.9Mh/s VGA消費電力：200W
PL：-50% ハッシュレート：37.4Mh/s VGA消費電力：130W
182 ：名無し一級：2017/11/20 15:33:30 (7年前) 0MONA/0人
174です。
何とかVIPPOOLのほうで掘れるの確認できました。
原因はhostsに書き込むとき頭に#がついてたからでした...お騒がせしました.....
183 ：fee二級：2017/11/20 15:50:01 (7年前) 0MONA/0人
現状VegaならCryptoNightの方が負荷も軽く収益も上ですね
自分の環境(Ryzen 5 Vega56)では
lyra2REv2 48MH/s 290W $2.88/d
CryptoNight 1960H/s 250W $4.31/d
・CPUマイニング中、システム全体で計測
・GPUClock 1407MHz MemClock 1100MHz
今後の最適化次第でしょうか
184 ：ガリレオ二段：2017/11/20 17:50:53 (7年前) 0MONA/0人
うまくいかないです。
Hostsも変更していますし、
ドライバーも最新ですし、
cmdファイルも変更したんですが、
「何かキーを押してください。」
で動かないです。パズルのようです。
185 ：名無し名誉名人教士：2017/11/20 20:51:18 (7年前) 0.00114114MONA/1人
電源を新調したので、いろいろ試してみた。VEGA64は1300MHzあたりが最良でしょうか？
VEGA64 @1537MHz(Mem 800MHz) 60.84MH/s 284W (0.214MH/s/W)
VEGA64 @1500MHz(Mem 800MHz) 59.35MH/s 250W (0.237MH/s/W)
VEGA64 @1400MHz(Mem 800MHz) 55.99MH/s 197W (0.284MH/s/W)
VEGA64 @1300MHz(Mem 800MHz) 52.05MH/s 162W (0.321MH/s/W)
VEGA64 @1200MHz(Mem 800MHz) 48.22MH/s 154W (0.313MH/s/W)
186 ：名無し初段：2017/11/20 22:31:37 (7年前) 0MONA/0人
>>185
検証お疲れ様です。
暖房代わりなら1400、効率重視なら1300という感じですか。
187 ：㌝㌝㌥初段：2017/11/21 01:49:51 (7年前) 0MONA/0人
RX480でテスト
ASMオプションをオンにしてみるとたしかにハッシュレート上がりますね...
188 ：名無し一級：2017/11/21 19:08:30 (7年前) 0MONA/0人
480でASM付けたら28から32になりました。
ASMってナンダ
189 ：名無し初段：2017/11/21 21:01:03 (7年前) 0MONA/0人
>>188
"assembler kernels"ってなってるから、ゴリゴリに最適化してるんでは？
10%程性能upっぽい記述があるけど、カードによってはクラッシュするかも？って・・・
190 ：㌝㌝㌥初段：2017/11/21 22:15:06 (7年前) 0MONA/0人
Ellesmere系=RX470,RX480,RX570,RX580 がクラッシュの可能性アリと…
とりあえず自分のRX480は一日回したけど、何もなし
191 ：㌝㌝㌥初段：2017/11/21 22:17:00 (7年前) 0MONA/0人
Polaris10と20って言ったほうがわかりやすいか？
192 ：あ熱帯います七段教士：2017/11/21 22:48:47 (7年前) 0MONA/0人
RISKプロセッサで、命令先読みペナルティを減らすために分岐回数が減るように並べ、演算順序入れ替えも考慮した並びで、アクセスが遅くなる命令も減らし相対ジャンプだけにしてレジスタだけで演算し、レジスタ転送も減らするように直接アセンブラ(機械語)でカーネルを作ったのかな?
もはや人間には無理と言われていることを地道にやったのかもｗ
193 ：名無し名誉名人教士：2017/11/22 07:48:09 (7年前) 0MONA/0人
mkxminerのダウンロードリンクが消されてる…
194 ：ガリレオ二段：2017/11/22 13:59:07 (7年前) 0MONA/0人
エラー４って、出てます。
カーネルについてなにか言っているようなのですが、
わかる方がいらっしゃったら、お願いします。
195 ：名無し一級：2017/11/23 10:03:59 (7年前) 0MONA/0人
採掘量が減っているような…
今まで１日で0.8モナだったのに0.2モナくらいになってる…
196 ：ハムちゃん二段錬士：2017/11/23 11:27:56 (7年前) 0MONA/0人
diffもアホみたいに上がってるし本格的に掘られてるんだろうね
197 ：名無し初段：2017/11/23 13:04:11 (7年前) 0MONA/0人
普段10時間ほどで0.1くらい掘れてるのに、今日は0.05しか掘れず何事かと思ったらdiff2倍以上になってたのか...
198 ：名無し名誉名人教士：2017/11/28 20:57:53 (6年前) 0MONA/0人
ccminerの最適化をsgminerに輸入してみた(未完成)
https://1drv.ms/u/s!Aud1FauQ46vHhywfOL0Uqaj_Ks4D
どこを間違えたのか、見当がつかない…
(lyra2rev2.cl、lyra2v2.clを主に修正しました)
OpenCLの解る人、助けて！
199 ：名無し三級：2017/11/30 21:53:45 (6年前) 0.039MONA/1人
>>198
OpenCL には詳しくないですが、Linux で試すと Lyra2REv2.cl のビルドエラーでした。
sph_u64 keccak_gpu_state[25];
の型が間違っていそう（ulong？？）なのと、
search1〜6 の引数を変えていたので、
altorithm.c の queue_lyra2REv2_kernel に修正が必要でした。もし添付の sgminer.exe に既にその変更が含まれていたら的外れで申し訳ないです。
search4 の lmem がどこから渡すべきか良くわからなかったので、引数から削って、search4 のローカル変数に書き換えたところ、一応ビルドは通りましたが、何か間違っているようで accept が返ってきません…。
200 ：名無し名誉名人教士：2017/12/01 21:33:33 (6年前) 0MONA/0人
ccminerの最適化をsgminerに輸入してみた(未完成2)
https://1drv.ms/u/s!Aud1FauQ46vHhy0fOL0Uqaj_Ks4D
大半が失敗なんだけど、たまに掘れるのは何だろう…
計算自体はこれでいいと思うけど、メモリ周りかな？
OpenCLの解る人、助けて！
201 ：名無し名誉名人教士：2017/12/01 21:43:37 (6年前) 0MONA/0人
>>200 補足
-w 64を付けないと動かない…
LocalMemory32768バイト÷64スレッド＝512バイト/スレッド
1スレッドあたり384バイト使うし、仕方ないよね…
202 ：光の名無し二段：2017/12/04 19:49:50 (6年前) 0MONA/0人
多少わかるけどマイナーのプログラム自体を自力で理解するところからはじまっちゃう…
Radeon派としてはなんとかしたいところではあるが
203 ：名無し二級：2017/12/05 10:54:15 (6年前) 0MONA/0人
mkxminer2.0でてる
204 ：名無し名誉名人教士：2017/12/05 11:11:35 (6年前) 0MONA/0人
>>203
出てるけど、devfeeが仕掛けられてるようだ…
* Note: the actual mining speed is a bit higher as some of the hashrate is used as devfee. The above values are what you (the miner) gets.
Other miners with devfee show you the total hashrate and then you have to reduce it by devfee: for example,
if you have 30MHs with 5% devfee, then effectively it is 28.5MHs for you (30 * (1.0 - 0.05)).
This miner shows you exactly the MHs that you earn, so devfee value is irrelevant.
とりあえず、mkxminerが公開停止されると面倒なので、sgminerの改造は続行かな？
205 ：名無し。。。五段教士：2017/12/05 11:17:20 (6年前) 0MONA/0人
devfeeとは？？
ムチでスマソ・・
206 ：名無し二級：2017/12/05 11:34:41 (6年前) 0MONA/0人
>>205
使用料
207 ：namo一級：2017/12/05 15:29:18 (6年前) 0MONA/0人
マイニングスレ期待age
208 ：名無し三級：2017/12/05 19:17:10 (6年前) 0MONA/0人
プロトコルが tls専用になっていてVipPoolで掘れない。これは悲しい
209 ：名無し四段：2017/12/05 20:51:29 (6年前) 0MONA/0人
退役寸前の280Xにもうひと働きしてもらおうと思ったけど
温度的にちょっときついかな・・・
210 ：名無しのモナー三級：2017/12/06 01:52:37 (6年前) 0MONA/0人
mkxminer v2.0.0
RX570 x 3
こんな感じで掘れてます。(pool=suprnova)
https://imgur.com/a/NqMWb
211 ：名無しさん：2017/12/07 10:37:03 (6年前) 0MONA/0人
>>210
同じ条件で半分ぐらいしか速度でません・・・なぜ・・
212 ：名無しさん：2017/12/07 10:58:41 (6年前) 0MONA/0人
suprnovaでもmonaでやったら同じぐらいでました失礼しました
213 ：鰯二段：2017/12/07 11:00:28 (6年前) 0MONA/0人
ローエンドでほそぼそと掘ってる人間だけどMKXは0.1MH以下は切り捨てて表示されるみたいだね
速度が何ｋH出てるのかわからないのはうーん……
214 ：名無し初段：2017/12/07 20:30:29 (6年前) 0MONA/0人
mkxminerの旧バージョンどこかにないですかね？
215 ：うー四段：2017/12/08 01:40:49 (6年前) 0MONA/0人
mkxminer v1.0.1どこ…
216 ：初心者名無し二級：2017/12/08 21:41:58 (6年前) 0MONA/0人
レベルの低い話ですいません
RX470の4GBのモデルでavg9.4Mh/sは適当な値なのでしょうか？
電気代とトントンだし1mona掘るのに2週間程度もかかるしあえてじゃない限りやる意味は無い？
217 ：名無し一級：2017/12/08 23:14:19 (6年前) 0MONA/0人
＞＞216
何で掘ってるかによるけど
480の8GBOCでmkxminer1.01なら30は出るよ
sgminer5.0.1なら9ぐらいだね...
218 ：初心者名無し二級：2017/12/08 23:53:32 (6年前) 0MONA/0人
言い訳でも何でも無く低レベルでした
同じくRX470の4GBのモデル
スレ冒頭の話題のsgminerだと9.4MH/s
mkxminerはexamle.cmdだと20MH/s
実際に動かすといくらか分かりませんが勉強しないと設定方法がわからないので時間かかりそうです
219 ：うー四段：2017/12/09 00:00:38 (6年前) 0MONA/0人
mkxminer v1.0.1
誰か下さい…
220 ：初心者名無し二級：2017/12/09 00:03:29 (6年前) 0MONA/0人
>>217
有難うございますRX480だと結構出るんですね
あ、２６MH/sぐらい出るかもしれません
221 ：光の名無し二段：2017/12/09 03:11:29 (6年前) 0MONA/0人
290Xを1140MHzまでOCして38MH/sのDiff32ギリギリキープとかいう電気代を無視したマイニングやってます
222 ：名無し二段錬士：2017/12/09 03:15:26 (6年前) 0MONA/0人
HD7970@1050MHzで5.5MH/sくらいだけど
もう少しくらい出るのかね？
223 ：lae九段尊者：2017/12/09 12:24:20 (6年前) 0MONA/0人
ま〜たLinux版ないね。
224 ：lae九段尊者：2017/12/09 12:25:17 (6年前) 0MONA/0人
更にGPL違反だらけ
225 ：名無しさん：2017/12/09 15:19:13 (6年前) 0MONA/0人
mkxminer 1.0.1やっと見つけてここに書いてある方法試してみたけど凄いね
sgminerの約5倍出てる
226 ：名無し四段教士：2017/12/09 15:19:29 (6年前) 0MONA/0人
SgminerでもMkxminerでも数百H/sにしかならない。CPU並の速度。
ゲフォと混在だとそんなものでしょうか？
227 ：名無し名誉名人教士：2017/12/09 17:13:39 (6年前) 0MONA/0人
ccminerの最適化をsgminerに輸入してみた(とりあえず動作版)
https://1drv.ms/u/s!Aud1FauQ46vHhy4fOL0Uqaj_Ks4D
mkxminerには及ばないけど、1.0.1がダウンロードできないからね…
228 ：名無し四段：2017/12/09 18:31:16 (6年前) 0MONA/0人
>>227
乙です
280Xで20分ほど回してみました
プールにもちゃんと届いてます
カーソルが変な位置に止まってるのはおま環ですかね？
229 ：光の名無し二段：2017/12/09 18:47:18 (6年前) 0MONA/0人
試そうと思ったらDLL足りない
230 ：名無し四段：2017/12/09 19:04:45 (6年前) 0MONA/0人
>>229
(未完成2)の時だけど
足りないDLLはググって出てきた↓ここからDLして動いてます(が保証はしかねる)
http://repairdll.info/en/
231 ：名無し名誉名人教士：2017/12/09 19:07:29 (6年前) 0MONA/0人
>>229 >>230
何が足りない？必要なら同梱してみる。
(私の環境では特にDLL必要ないんだが…)
232 ：名無し名誉名人教士：2017/12/09 19:13:23 (6年前) 0MONA/0人
>>227 の改造は、
①3段打ちメソッドを採用
②blake256とkeccak256を融合、Lyra2の3段目とskein256を融合
③Lyra2の2段目を4スレッド分割(for文で4周しているだけ)にして、LocalMemoryを使用。(そのため、-w 64が必須)
を行いました。これでもまだmkxminerには遠く及ばないんだが…
233 ：光の名無し二段：2017/12/09 19:15:33 (6年前) 0MONA/0人
ぬう
234 ：光の名無し二段：2017/12/09 19:17:23 (6年前) 0MONA/0人
>>231
OpenSSL系のが全部足りないらしいです
235 ：名無し名誉名人教士：2017/12/09 19:17:26 (6年前) 0MONA/0人
なお、LocalMemoryが32768バイト確保できる前提で作られています。
LocalMemoryの容量がそれに満たない場合、動作しませんので悪しからず…
236 ：光の名無し二段：2017/12/09 19:20:39 (6年前) 0MONA/0人
sgminer -n
プラットフォームは変えてます
237 ：名無し名誉名人教士：2017/12/09 19:33:43 (6年前) 0MONA/0人
>>234
具体的に必要なDLLを教えてください。
当方、コンパイル環境があるせいか、何が不足しているのかわからない状況…
238 ：光の名無し二段：2017/12/09 19:38:31 (6年前) 0MONA/0人
libssl-1_1-x64.dll と libcrypto-1_1-x64.dll ですね
239 ：名無し名誉名人教士：2017/12/09 20:11:53 (6年前) 0.00114114MONA/1人
DLL同梱＆-w 64制限解除版(sgminer-5.2.0-r2)
https://1drv.ms/u/s!Aud1FauQ46vHhy4fOL0Uqaj_Ks4D
Lyra2の2段目だけを64スレッド制限にして、その他を256スレッドで動くようにした。
どうだろうか？
240 ：光の名無し二段：2017/12/09 20:27:38 (6年前) 0MONA/0人
https://pastebin.com/AfT9CUae
うーん
290Xだとビルドエラーがとれないですねえ
clBuildProgramの途中で3行同じのが出てる気はしますが
241 ：名無し四段：2017/12/09 20:49:35 (6年前) 0MONA/0人
>>239
-wの制限解除は内部的？
128や256指定したらエラー吐きました
-w 64でハッシュレート微増です
242 ：名無し四段：2017/12/09 21:28:05 (6年前) 0MONA/0人
batファイルに -g 2 追加でさらに微増
色々試行錯誤してたscrypt時代を思い出してきた
243 ：光の名無し二段：2017/12/09 21:28:21 (6年前) 0MONA/0人
なぜかこのビルドのsgminer.exeだと、kernelのパス指定が効いていないようで中身を全部tempフォルダにぶち込んだら動きました
+10MH/sくらいでてますねー
244 ：名無しさん：2017/12/09 21:46:25 (6年前) 0MONA/0人
RX470 8G
Sgminer 9MH/s
r1 15MH/s
r2 -g 2 17MH/s
245 ：名無し＃三段：2017/12/10 23:31:15 (6年前) 0MONA/0人
mkxminer 1.0.1が見つからず >>239もエラーで利用できず・・
既存のsgminer 5.1.0でRX560x2で800KH/sという低さなのでなんとかしたい
mkxminer 2.0でhosts上書きしてもmonaに利用できないのは規制されているからだよね？
246 ：光の名無し二段：2017/12/10 23:37:11 (6年前) 0MONA/0人
>>245
239で起きてるエラーが私と同じだったのなら、%temp%に*.clを全部いれちゃえば動くようになる
247 ：名無し＃三段：2017/12/11 00:20:35 (6年前) 0MONA/0人
>>246
ありがとう
アプリケーションエラー0xc000007bになりました
DirectXもVisual C++も入ってるんだけどなぁ
でもあとこれ解決すれば動きそうだ
248 ：光の名無し二段：2017/12/11 02:04:40 (6年前) 0.1MONA/1人
https://gpuopen.com/gaming-product/radeon-gpu-analyzer-rga/
mkxminerは多分これつかってiccつかったc++のような状態とみたんだ
オフラインコンパイルしようとしてるけど前提知識がたりんわ…
249 ：光の名無し二段：2017/12/11 02:13:59 (6年前) 0MONA/0人
binできた( ´∀｀)プロファイラかけたり遊べるぜ
250 ：光の名無し二段：2017/12/11 02:26:19 (6年前) 0MONA/0人
https://pastebin.com/UWNPnd6D
lyra2rev2.clをコンパイルした結果のアセンブラソース(Hawaii版
10年くらい前ならアセンブラ読むの平気だったけどもうツライ
251 ：光の名無し二段：2017/12/11 03:04:23 (6年前) 0MONA/0人
CodeXLのコンパイラに最適化させようとおもったらkeccak1600.clのインライン化の部分で文句を言われる
252 ：光の名無し二段：2017/12/11 03:26:46 (6年前) 0MONA/0人
L1キャッシュもうちょい使えるのと
SIMDちゃんと使えてない？
LDS使うのはやめてファイルレジスタに全部押し込んだんだっけ…
253 ：光の名無し二段：2017/12/11 11:40:55 (6年前) 0MONA/0人
最適化オプション全開で+2~3MH/sくらい
254 ：初心者名無し二級：2017/12/13 22:28:05 (6年前) 0MONA/0人
rx470でもocと普通のモデルだと全然結果違うんスね…
23mh/sと30mh/sで一割ぐらい違うならともかく3割違うって…
255 ：名無し名誉名人教士：2017/12/14 06:38:46 (6年前) 0.00114114MONA/1人
実験的にblake256&keccak256、lyra2の1段目、3段目で配列を排除してみた。(sgminer-5.2.0-r3)
https://1drv.ms/u/s!Aud1FauQ46vHhzAfOL0Uqaj_Ks4D
ただ、これらはそれほど重くないので、対して効果が見込めない…
(ってか、効果はあるのか？)
256 ：名無し四段：2017/12/14 20:51:50 (6年前) 0MONA/0人
sgminer-5.2.0-r3
280X 「-g 2」オプションあり
微増だけど数字が増えるのはいいもんだね
名無し名誉名人ありがとう!!
257 ：名無し名誉名人教士：2017/12/15 09:05:11 (6年前) 0MONA/0人
mkxminerの紹介を見ると、
Increase intensity. 23 is recommended for 4GB+ videocards.
って書いてある。
Lyra2REv2は全てグローバルメモリ(実メモリ)上に確保すると、1536バイト＋32バイト使うので、-i 23の場合、
(1536+32)×2^23÷2^20＝12544MBとなり、通常では確保できない。
全てローカルメモリに納めると、1536バイトが128バイト(スレッド間転送用)で良くなり、必要なくなり、
(128+32)×2^23÷2^20＝1280MBとなる。これなら、-i 24.5までイケるんだよな…
4分割を採用して、グローバルメモリを使う場合、1536バイト÷4＝384バイトで、
(384+32)×2^23÷2^20＝3328MBとなる。このあたりが符合するんだよね…
ひょっとして、ローカルメモリを(あまり)使っていないのかな…
258 ：光の名無し二段：2017/12/15 18:35:55 (6年前) 0MONA/0人
mkxで放置してて測るの忘れてた
--gpu-platform 1 -I 20 -w 64 --incognito -g 2
290X@1140/1350
結構速くなった気が？！
259 ：光の名無し二段：2017/12/15 18:39:35 (6年前) 0MONA/0人
>>257
290X(4GB)だとI24まであげれました
メモリギリギリィ
260 ：名無し三級：2017/12/15 19:34:39 (6年前) 0MONA/0人
api-ms-win-core-libraryloader-l1-2-0.dllがないと言われて
起動できないんですがどうすればいいのでしょう？
261 ：鰯二段：2017/12/23 12:35:01 (6年前) 0MONA/0人
件のmkxminerが3.0betaでSSL接続以外にも対応しましたね
私の環境では動作確認が取れないのですが、作者いわく「コアに大幅な変更が入った」とのことです
下記リンク
https://bitcointalk.org/index.php?topic=2360168.msg26810404#msg26810404
262 ：名無し四段教士：2017/12/23 13:21:52 (6年前) 0MONA/0人
mkxminerはfeeがあるのがなぁ。。。
263 ：AN二段：2017/12/23 13:32:20 (6年前) 0MONA/0人
う、、i5+GTX1080使いだけど、かつてK6-2でセレロン300A機とFPSで戦った身としては１が心に刺さった・・・
次はryzenにしようと思ってるけど、Lyra2REv2にRadeonか・・・
ちょくちょく覗かせてもらいます～ｗ
264 ：名無し名誉名人教士：2017/12/27 17:40:19 (6年前) 0MONA/0人
githubでsgminerを探していたら、こんなの見つけた
https://github.com/MaxDZ8/sgminer/tree/lyra2REv2-4w-monolithic
kernel名lyra2rev2.mdzで動くみたいだね。ローカルメモリはスレッド間のデータ転送にのみ使用するようだ。
帰ったら実験してみようかな…
265 ：わろちｗｗｗ三段：2017/12/27 17:45:32 (6年前) 0MONA/0人
わざわざLyra2REv2で掘らなくても、zpoolとかのマルチプールでAMD有利の仮想通貨を採掘して支払いをMONAにしてもらえば効率がいいと思うよ
俺はgtx1070だけど、monaのプールマイニングだと0.25mona/dayぐらいだったけどzpoolにしたら0.4mona/dayぐらいに上がったよ
AMDだともっと上がるかも
266 ：名無し名誉名人教士：2017/12/27 19:37:36 (6年前) 0MONA/0人
>>265
monacoinを維持するシステムとしてのマイニングなわけで、
どっちが有利とか、そういうスタンスで開発しているわけではないんですよ。
儲けが出ないからって、掘る人がいなくなったら、コインが維持できない…
(まあ、この問答は、信者と儲は相いれないわけだが…)
267 ：名無し名誉名人教士：2017/12/27 19:42:10 (6年前) 0MONA/0人
>>264 のマイナーを試してみたけど、>>255 の半分程度の速度だった…
やっぱり、ローカルメモリは必須かな？
268 ：名無し名誉名人教士：2017/12/30 22:08:26 (6年前) 0.00114114MONA/1人
現在迷走中…とりあえず若干速くなったと思う(mkxには遠く及ばないが…)
sgminer-5.2.0-r4
https://1drv.ms/u/s!Aud1FauQ46vHhzEfOL0Uqaj_Ks4D
今回はソースコードもうｐしてみます。(たいした改造はしていませんが…)
https://1drv.ms/u/s!Aud1FauQ46vHhzIfOL0Uqaj_Ks4D
269 ：名無し名誉名人教士：2017/12/30 22:11:34 (6年前) 0MONA/0人
ちなみに…
ローカルメモリをuintでアクセスする⇒逆に遅くなった
全部を1つのkernelにしてみた⇒めっちゃ遅くなった
Lyra2の2段目を起点として、前半、後半をそれぞれ纏めてみた⇒若干遅くなった
あと、何をすればいいのか…
270 ：光の名無し二段：2017/12/30 23:17:01 (6年前) 0MONA/0人
>>269
analyzerで解析するとレジスタがあまってるsearchとかSIMDが全然効いてないsearchとかあったり
sgminer(とLYRA2Rev2.cl)をVisualStudio+CodeXLでデバッグすると悲しいくらいGPUが回ってない
でも最適化する頭がなかった／(^o^)＼
271 ：光の名無し二段：2017/12/30 23:23:31 (6年前) 0MONA/0人
あ あと古いsgminerだとWindowsSDK10でビルドできなくてデバッグすらできなかったんで適当なforkさがしてきた
https://github.com/MaxDZ8/sgminer
272 ：光の名無し二段：2017/12/30 23:24:50 (6年前) 0MONA/0人
AMDAPPSDKも2.xもう配布してないんだよもう…めんどくさいやつめ
273 ：光の名無し二段：2017/12/30 23:26:30 (6年前) 0MONA/0人
OCL_SDK_Lightで通るけどなんかイヤ
274 ：名無し名誉名人教士：2017/12/31 07:44:53 (6年前) 0MONA/0人
search3を見てみると…(Lyra2の2段目)
SALU命令：281 (全コアで共通のデータを扱う)
VALU命令：2614 (各コア別々のデータを扱う)
LDS命令：604 (ローカルメモリのアクセスを扱う)
FLAT命令：8 (グローバルメモリのアクセスを扱う)
ちなみに、s_waitcnt命令が(SALU命令281個中)275個もある。これを減らせればいいのかな？
でもbarrier(CLK_LOCAL_MEM_FENCE)は外せないよね…
275 ：名前を無くした名無し二級：2017/12/31 11:11:20 (6年前) 0MONA/0人
>>272
ttps://developer.amd.com/amd-accelerated-parallel-processing-app-sdk/
2.9ではだめ？
276 ：名前を無くした名無し二級：2017/12/31 11:23:31 (6年前) 0MONA/0人
>>271
ナイスハッシュ系だけど、比較的新しいフォークは使い物にならない感じ？
ttps://github.com/nicehash/sgminer/releases/tag/5.6.1
277 ：名無し名誉名人教士：2017/12/31 12:08:05 (6年前) 0MONA/0人
>>276
なぜかコンパイル通らないんだよね…
278 ：名前を無くした名無し二級：2017/12/31 14:17:45 (6年前) 0MONA/0人
IDA Pro Freeでmkxの中身覗いてみようと思ったけど
俺にはスキルが足りなかったｗ
279 ：光の名無し二段：2017/12/31 15:48:43 (6年前) 0MONA/0人
>>276 >>277
Windows向けのビルド、ほとんどの奴がWindows7までにしか対応してないんだ（SDK7
それがMSVC2010かな
たまに2013とか書いてあるけどSDK7のまま
Windows10でSDK10いれてるからうちの環境じゃビルドできんっていう
280 ：光の名無し二段：2017/12/31 15:56:57 (6年前) 0MONA/0人
>>275
証明書期限切れでアクセスできないんだよね…AMDはよ直して
281 ：名無し名誉名人教士：2017/12/31 22:52:19 (6年前) 0MONA/0人
>>280
証明書期限切れでもアクセスする方法なかったっけ？
Edgeなら「詳細」⇒「Web ページへ移動 (非推奨)」で行けたはず。
282 ：名無し一級：2018/01/02 23:05:42 (6年前) 0MONA/0人
RX480でmkx1.0.1だと34MH/s出るけどsgどうなりましたかね、前sg使ってた時は9MH/sしか出なかったんですけど
283 ：光の名無し二段：2018/01/02 23:36:24 (6年前) 0MONA/0人
284 ：名前を無くした名無し二級：2018/01/03 00:24:28 (6年前) 0MONA/0人
>>282
R9 290@1100MHzでこんな感じだね。
285 ：光の名無し二段：2018/01/05 00:13:16 (6年前) 0MONA/0人
https://pastebin.com/L9AFHpKu
ガイド付きプロファイルしたら文句ばかり
Radeon特有のOpenCLに渡す前と結果を受け取る部分の最適化でもしようとおもったらそれどころじゃなかった
286 ：光の名無し二段：2018/01/05 01:24:50 (6年前) 0MONA/0人
>>268
CLだけもってきてみようと思ったら
Error -11: Building Program (clBuildProgram)
そして案の定WindowsSDK10でそのsgminerはビルドできないので／(^o^)＼ｷﾞｷﾞｷﾞ
287 ：光の名無し二段：2018/01/05 01:54:49 (6年前) 0.2MONA/2人
元々のsearch3が遅すぎてホスト側を最適化するきがおきないｗｗ
高速化したkernelのこれがみたいんだよねー
288 ：光の名無し二段：2018/01/05 02:02:31 (6年前) 0MONA/0人
あとなぜか名誉名人さんとか他の人のsgminerは-Kが効かない（なぜか\User\AppData\Local\tempを参照する）ので
プロファイラにかけるのがめんどくさいのです
289 ：名無し名誉名人教士：2018/01/05 09:48:26 (6年前) 0MONA/0人
>>287
正直、CodeXLの使い方が分からんのよね…
ってか、元のカーネルすげーなｗｗｗ
290 ：光の名無し二段：2018/01/05 18:11:14 (6年前) 0MONA/0人
>>289
CLもってきても動かないから名人さんのやつをWindowsSDK10で動くように根性
search3へってますなあ
291 ：光の名無し二段：2018/01/05 18:15:14 (6年前) 0MONA/0人
さてどこをいじるか
292 ：光の名無し二段：2018/01/05 18:38:43 (6年前) 0MONA/0人
3ってVALU使用率100％なのにv/sALU Busyがやたら低い
LDSBankConflictがわずかにあるってことは跨いでしまっているのか
2と4はなんかちょいちょいストールしてますね
293 ：名無し名誉名人教士：2018/01/05 20:44:18 (6年前) 0MONA/0人
>>292
4スレッド単位で1つ目のスレッドのstate[0]を読みに行くため、
4WayBankConflictが4回働いています。これは不可避です。
なお、CUDAの方ではWarpShuffleがあるため、LDS(CUDAではSharedMemory)を経由せず、BankConflictが起こりません。
294 ：光の名無し二段：2018/01/05 22:33:41 (6年前) 0MONA/0人
>>293
APPSDK3.0とOpenCL2.0EXTEND AMD眺めてますが、
そもそもOpenCLに詳しくないのでホスト側もまだ眺めてます
RadeonSoftwareがAdrenalinになったときOpenCL2.1になったはずなんですがC++書くとコンパイルエラー吐かれて憤慨してます／(^o^)＼
GCN用APIに置き換えられそうなところ探してるんですがなかなかないですねえ
295 ：光の名無し二段：2018/01/05 23:11:04 (6年前) 0MONA/0人
どうでもいいバグ
applog(LOG_DEBUG, "lyra2REv2 buffer sizes: %lu RW, %lu RW", (unsigned long)bufsize, (unsigned long)buf1size);
でアサートしてました
手元では表示するものが無いので0にしてますが
296 ：光の名無し二段：2018/01/06 01:36:06 (6年前) 0MONA/0人
ttp://free.eol.cn/edu_net/edudown/AMDppt/OpenCL%20Programming%20and%20Optimization%20-%20Part%20I.pdf
GCNの勉強がはじまってしまった
テストコード走らせてると部屋が寒いです
297 ：光の名無し二段：2018/01/06 02:02:21 (6年前) 0MONA/0人
49ページ位からはじまるWavefrontの話がキモに見える
でも次のスレッドのデータを先読みする ってやってないよね多分
search3の部分だけもっと分解うーん
こういう悩みはハッテンさんが既にやって今に至ってる気がするので無駄な気がしてきた
298 ：名無し一級：2018/01/07 13:37:42 (6年前) 0MONA/0人
宜しくお願い致します(￣^￣)ゞ
299 ：光の名無し二段：2018/01/11 23:21:59 (6年前) 0MONA/0人
試しに data->compiler_options に-cl-opt-disable をいれて最適化オフにしてみたら26MH/s→300KH/s
300 ：光の名無し二段：2018/01/12 21:17:04 (6年前) 0MONA/0人
https://pastebin.com/YBAiDi4x kernelを非同期にしてしまうのはこうするらしい
リードをオーバーラップさせるには、
フェンス→なんか計算する→ローカルメモリにアトミック系命令→フェンス
ライトをオーバーラップさせるには
ローカルアドレスを取得する→フェンス→ひたすら書く
301 ：光の名無し二段：2018/01/12 21:22:56 (6年前) 0MONA/0人
https://pastebin.com/gLt4uV09
OpenCL/C++ templateふつくしい
今のCL/Cのコード、Defineまみれで読むのが辛い(´・ω:;.:...
全部書き直すとかちょっとありえ ぬ
302 ：名無し名誉名人教士：2018/01/22 14:38:44 (6年前) 0MONA/0人
現在、L1キャッシュを有効に使えないか考察中
1スレッド当たり、384バイト＋スレッド間共有32バイト＝416バイト
LocalMemoryのみの場合、
64kバイト÷416バイト/スレッド÷64コア≒2.46スレッド/コア となるため、2スレッド起動となる。
一方、L1キャッシュも含めると、
L1キャッシュ16kB＋LocalMemory64kB＝80kB
80kバイト÷416バイト/スレッド÷64コア≒3.08 となるため、3スレッド起動できそうだ。
まあ、上手くいけばね…
303 ：名無し名誉名人教士：2018/01/22 14:40:45 (6年前) 0MONA/0人
そして、L1キャッシュはCUDAでは(現状では)上手く使えていない…
グローバルメモリはL1にキャッシュされないので…
Radeonではどうかな…？
同様の理屈で無理っぽいけど、プライベートメモリならキャッシュされるかも？
304 ：名無し名誉名人教士：2018/01/22 15:33:42 (6年前) 0MONA/0人
1スレッドあたり、演算用に48個、スレッド間共有に4個のuint64_tが必要。
3スレッド起動を考えると…
LocalMemoryで確保できるのは42個(uint64_t換算)
L1キャッシュで確保できる(見込みな)のは10個(uint64_t換算)
なので、LocalMemoryで演算用に38個、スレッド間共有に4個。L1キャッシュで演算用に10個。
この場合、演算途中でif分岐が発生するのであまり望ましくない…(演算用を4の倍数にしたい…)
スレッド間共有を2回に分けて使えば、2個分余裕が生まれて、
LocalMemoryで演算用に40個、スレッド間共有に2個。L1キャッシュで演算用に8個。
これなら上手くできそうだ。(同期処理が1セット増えるが…)
今はCUDAの方を弄っているので、一段落したらコッチもやってみる。
305 ：光の名無し二段：2018/01/22 15:48:29 (6年前) 0MONA/0人
GCNの場合、WavefrontっていうSIMDへの先読み機構があってこいつがフル回転しているとメモリのレイテンシが隠蔽されてALUがフル回転できるようになるみたいです
LDSはL1の8倍ほど帯域があり、L2はL1の半分以下…ですがWavefrontが先読みを行いこれを隠蔽する、だそう
306 ：光の名無し二段：2018/01/22 15:52:15 (6年前) 0MONA/0人
またLDSのデータは2vector型(32bit)単位で行われるので粒度に注意
L1は32bit wordもしくは4vector
307 ：光の名無し二段：2018/01/22 15:56:49 (6年前) 0MONA/0人
WavefrontはCU単位に設置されていて(256/#使用レジスタ数)*4
ここで__private（レジスタ）との匙加減が決まるみたい
WFの数が上手くあえばグローバルメモリ使い放題（…なのか？）
308 ：光の名無し二段：2018/01/22 16:09:19 (6年前) 0MONA/0人
kernelを出入りするデータは可能な限りvector型
このへんをOpenCL/Cコンパイラが最適化してるみたいなのでオフにしたらめっちゃ遅くなったと
309 ：光の名無し二段：2018/01/27 01:03:32 (6年前) 0MONA/0人
悲報 290Xさんついに壊れる
RX580選手を手配中です
310 ：あ熱帯います七段教士：2018/01/27 01:13:01 (6年前) 0MONA/0人
290X 2台壊れたな 1080も1台
311 ：まつあき六段：2018/01/27 02:45:24 (6年前) 0MONA/0人
壊れたGPUをヒートガンで炙るのが趣味(´=ω=`)
312 ：あ熱帯います七段教士：2018/01/27 04:00:52 (6年前) 0MONA/0人
290Xのガラスっぽい表面にヒビが入っているので、あっためても無理(^^;
1080は端子掃除すると生き返るかも。ガが死んでいた。
313 ：鰯二段：2018/01/27 16:04:02 (6年前) 0MONA/0人
295X2に置き換えよう（提案）
314 ：名無しさん：2018/01/27 19:22:55 (6年前) 0MONA/0人
あいぽんもオーブンチンで治るつーかハンダクラックならサラマンダー通すみたいなもんなんで表面温度250度ぐらいまでならどんどん炙っておｋ
315 ：名無し名誉名人教士：2018/01/28 08:10:07 (6年前) 0MONA/0人
ccminerと同じアルゴリズムを適用しようとしたんだけど、プライベートメモリがうまく確保できない…
どうやら、確保数が少なすぎてコンパイラがレジスタに置き換えるらしい。
state[0]の下位2bitでアクセスするアドレスが分岐するから、コンパイルに失敗する…
これって、コンパイラの欠陥だろ…
316 ：光の名無し二段：2018/01/28 08:57:44 (6年前) 0MONA/0人
RX580にむけて資料探してるけどPolalisのあんまりみつからんな…
317 ：光の名無し二段：2018/01/28 09:17:51 (6年前) 0MONA/0人
https://pc.watch.impress.co.jp/docs/column/kaigai/1009584.html
https://news.mynavi.jp/photo/article/20160629-polaris/images/005l.jpg
これもうOpenCL/Cで意図的に狙って書くの難しすぎない…
asmいるんかなあ
318 ：光の名無し二段：2018/01/28 09:40:44 (6年前) 0MONA/0人
https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/private.html
うーん？
319 ：光の名無し二段：2018/01/28 14:38:05 (6年前) 0MONA/0人
290X VC+50mV PL+50% 1140MHz 38MH/s FAN100% 7000rpm 70℃
RX580 PT+12% TT+5℃ 1360MHz 38MH/s FAN40%~ 1330rpm 62℃
どっちもmkxV1 どっちも38MH/s
sgminer改は全然数字でないっすねーTargetTempもっとあげないとちっともまわらないな
290Xは極限までOCしてました＾ｐ＾燃え尽きた
320 ：名無し名誉名人教士：2018/01/30 20:03:18 (6年前) 0MONA/0人
ccminerの改造を取り込んでみたが、思うように取り込めなかったので、この形にした。
sgminer-5.2.0-r5
https://1drv.ms/u/s!Aud1FauQ46vHh0EfOL0Uqaj_Ks4D
前バージョン(sgminer-5.2.0-r4)からの差分はこちら
https://1drv.ms/u/s!Aud1FauQ46vHh0AfOL0Uqaj_Ks4D
プライベートメモリが確保できないなら、レジスタを使えばいいじゃない！
…ってなノリでif文多用で無茶してみました。
321 ：光の名無し二段：2018/01/30 22:33:15 (6年前) 0MONA/0人
うーむ
290Xのときは名誉名人さんのsgminer改でそこそこ上がってたんだけど
RX580だとさっぱり…
AfterburnerのGPUpower見ての通り全然回ってないみたい
タスクマネージャのComputeは100％なってんだけどねえ
322 ：名無し名誉名人教士：2018/01/31 06:27:24 (6年前) 0MONA/0人
Radeon RX 560にて測定
sgminer-5.2.0-r4
sgminer-5.2.0-r5
323 ：名無し名誉名人教士：2018/01/31 06:29:15 (6年前) 0MONA/0人
>>322
おっと、ユーザー名消し忘れたｗｗｗ
324 ：名無し名誉名人教士：2018/01/31 06:32:34 (6年前) 0MONA/0人
再うｐ
Radeon RX 560にて測定
sgminer-5.2.0-r4
sgminer-5.2.0-r5
325 ：名無し名誉名人教士：2018/02/05 10:20:09 (6年前) 0MONA/0人
現状の確認
・CUDAにて、プライベートメモリ(CUDAではローカルメモリと呼ぶ)を少量確保して、L1、L2キャッシュを有効利用できることを確認
・OpenCLでプライベートメモリを確保して同様のことをやってみる
⇒コンパイラの最適化が邪魔して(だと思う)プライベートメモリをレジスタで確保しようとして失敗する(for文中の配列の添字を可変にできない)
・仕方ないので、if文で分岐することで、添字を固定する。
⇒とりあえず上手く行った。速くなったので公開。(sgminer-5.2.0-r5)
・元に戻し、プライベートメモリの確保量を増やして再度やってみる
⇒今度はコンパイルが通った(ちゃんとプライベートメモリを確保した)が、速度は出なかった。プライベートメモリの容量が大きく、L1、L2から溢れたと思われる。
さて、どうしたものか…
326 ：光の名無し二段：2018/02/05 12:48:24 (6年前) 0MONA/0人
327 ：光の名無し二段：2018/02/05 12:55:33 (6年前) 0MONA/0人
290Xで高速化が実感できたr3をRX580で回しなおしてみました
r5よりはやい…？
アフターバーナーの監視はポーリング100msにしてあります
Compute1と2が動きますね ビデオエンコーダはAVCとHEVCで別なのが確認できていますが… HighPriorityComputeとかはなんぞや
328 ：光の名無し二段：2018/02/05 13:08:35 (6年前) 0MONA/0人
>>325
コンパイラの最適化
前に-cl-opt-disableをつけてclBuildProgramを実行したことがあるので
r5の手法が効果があるかどうかはそれで確認できるかも？
329 ：名無し＃三段：2018/05/30 12:36:27 (6年前) 0MONA/0人
mkxminer3.1.0のsuprnova SLL:3001でマイニングできなくなってますね・・・
他にRadeonでモナ掘れる所ないかな？
330 ：だれかさん七段教士：2018/05/30 18:05:22 (6年前) 0MONA/0人
SSLで掘る必要あるの・・・？
他のポートでやってみては？
331 ：名無し＃三段：2018/05/31 06:52:18 (6年前) 0MONA/0人
ありがとう
tls→tcpで解決しました
お気に入り
新規登録してMONAをもらえた
本サイトはAsk Mona 3.0に移行しましたが、登録すると昔のAsk Monaで遊ぶことができます。