sgminer を高速化サセようぜ

採掘 sgminer Lyra2REv2 OpenCL

141 Res. 621.39777484 MONA 31 Fav.

1 :ハッテン場五段:2016/04/10 02:01:05 (1年前)  512.34MONA/5人

Lyra2REv2 での採掘は現状 Radeon は圧倒的不利な状況です。K6 時代からの AMD 信者としては悔しくてたまらない。なんとかして高速化サセねば!

言うばっかではなんなので、グラボでマワしたら8%ほど高速化できたものを公開します。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 (1年前)  0.00000184MONA/1人

去年の9月頃に、いろいろと遊んでいたのだけど、あれから状況が変わっていなければ、-g 1から -g 2にしてスレッドを1から2にしても、性能が10%以上、高速化される。(R7 250XEによる測定結果)

そのときに考えた、高速化される、理由は

Lyra2REv2は、いろいろなハッシュ関数を7個ならべたもので、真ん中の1つだけが、メモリ帯域を大食いする関数。他のハッシュ関数は、メモリ帯域をあまり食わない。
2スレッドにすれば、1スレッド目がメモリ帯域を大食いする関数で演算器がアイドルになっている状況で、2スレッド目が、メモリ帯域を食わない関数を実行できれば、-g 2で性能が向上する理由になる。

>>1 の高速化版は、-g オプションを使うと、性能は、どれくらいになるのでしょうか?

3 :名無し名誉名人教士:2016/04/10 08:10:47 (1年前)  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 (1年前)  1.1MONA/2人

オレのR7 250が火を噴くぜ!

5 :izuna五段:2016/04/10 10:01:22 (1年前)  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 (1年前)  0MONA/0人

>>4

織田信長の3段撃ち、をするのだ!

7 :ハッテン場五段:2016/04/12 10:29:22 (1年前)  0.1MONA/1人

みなさま、投げモナとご意見ありがとです。
かなり濃いモナをぶつけて下さったようで恐縮です。
色々とヤル気がでてきました。

現実的なのはメモリ再配置か、グローバルメモリ、プログラム中で言えば DMatrix 配列への読み書きをいかに減らすか、はたまたその両方かと思ってます。
1536 byte 全部読まなくていいんだよ!とか、全部書き込まなくてもいいんだよ!とかそういう状況を見つけ出せたら。
いろいろと試しちゃいるんですが、結果がぜんぜんでません。
仕事もモンハンも射精活動もかわいい子の物色もしないといけなくて、なかなか進まない。
とりあえず休みたい

8 :hururamu六段:2016/04/12 20:45:11 (1年前)  0.2MONA/1人

>全部読まなくていいんだよ!とか、全部書き込まなくてもいいんだよ!
>仕事もモンハンも射精活動もかわいい子の物色もしないといけなくて、なかなか進まない。

全部やらなくてもいいんだよ!

9 :ハッテン場五段:2016/04/13 04:01:10 (1年前)  0.1MONA/1人

はーい 当社比12%性能アップ版が作れましたよー イクイクー
http://kernel.hattenba.xyz/kernel-20160413.zip

OpenCL kernel のほうじゃなくて sgminer のソース見てて気づいたんですけど、lyra2v2 のとこで使ってる 1536 byte のデータって、ほかのところでも使ってるわけじゃなくて、lyra2v2 のとこだけでしか使ってないみたいですね。
ということはデータの整合性がとれる範囲であれば、なぶっちゃっても問題なし、と。
そこんとこ省いたら性能が4%ほどアップ致しました。

10 :ハッテン場五段:2016/04/13 04:05:29 (1年前)  0MONA/0人

今回のは、DMatrix への最後の方の書き込みを一定の条件下であればヤラないことにしました。OpenCL の苦手な比較演算が激遅グローバルメモリへのアクセスに打ち勝ったのです。

>>8
オスとして・・その衝動は止められない!!

11 :ハッテン場五段:2016/04/13 14:58:53 (1年前)  0.1MONA/1人

あ、すみません、アップロードするファイルを間違えてたっぽいです。
差し替えてますんで、コンパイルが通らなかった方は今一度ダウンロードお願いします。このファイルは 17402 バイトです。間違えてたのは 17403 バイトです。
http://kernel.hattenba.xyz/kernel-20160413.zip

12 :リキプロマン六段:2016/04/13 18:49:37 (1年前)  0MONA/0人

reduceDuplexRowf を3回ループで回す際に、最後の1回の時に別のメソッドを用いて条件分岐させて、DMatrixの書き込みを減らしているのですね。
ところで、s1のこの条件はどうやって見つけたのですか?自分もLyra2v2の最適化に興味があるので良ければ教えて下さい。

13 :ハッテン場五段:2016/04/14 04:08:59 (1年前)  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 (1年前)  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 (1年前)  0MONA/0人

>>14 で書いた企ては見事に失敗しました。動くけど遅くなる!

16 :リキプロマン六段:2016/04/14 20:09:03 (1年前)  0MONA/0人

とても詳しく解説して頂きありがとうございました。
そのshiftのとりうる値に合わせてDmatrixの読み込む大きさを変化させるわけですね。

逆に遅くなるのはなぜなんでしょうか。
自分もいろいろ考えていますが決定的な案が浮かびませんね…

17 :名無し名誉名人教士:2016/04/14 21:19:48 (1年前)  0.1MONA/1人

>>16
GPUは分岐が苦手だから…

18 :ハッテン場五段:2016/04/18 07:05:36 (1年前)  0MONA/0人

>>17
その一言に尽きますな

19 :ハッテン場五段:2016/04/21 07:50:10 (1年前)  0MONA/0人

条件分岐をなくそうと bitselect 関数を使おうとするとなぜか OS が固まるこの頃、皆様イカがお過ごしでしょうか。なかなかうまくいきません。つか地震ばっかでそれどころじゃないです。すみません。

20 :ハッテン場五段:2016/04/21 19:19:26 (1年前)  0MONA/0人

当社比17%アップ版ができましたよー イクイクー
http://kernel.hattenba.xyz/kernel-20160421.zip

でもこれ開発マシン(OpenCL2.0)とは別マシン(OpenCL1.2)で動かしてみたら計算速度が4割ほど落ちました。環境を選びます。ドライバが新しくないとダメなのかな?別マシンは Vista なんで古いドライバを未だに使ってます。

21 :ハッテン場五段:2016/04/21 19:26:41 (1年前)  0MONA/0人

今回の変更は大きく分けて2つ。

一つはちょこっとお手入れして、掛け算の回数を減らしました。掛け算や割り算の回数を減らすのは計算量減らしの基本っすもんね。昔学校で習った。でもこれじゃ全然早くなってないんすよね。おまけ状態です。ププ

もう一つは激遅グローバルメモリに居座る DMatrix 配列を関数でカプセル化して、データの一部をプライベートメモリに移したことです。条件分岐の回数が格段に増えましたが、それでも以前よりも更に5%速度が向上しました。

本当は DMatrix 配列のデータは全部プライベートメモリに移したいのですが、皆様御存知の通り、Radeon ではコアあたりのメモリ量が全然足りません。なんで一部だけ。

22 :siv三段:2016/04/21 22:31:21 (1年前)  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 (1年前)  0MONA/0人

>>22
ご報告ありがとうございます。もしかしたら環境にすごく左右されることばっか俺はやってるのかもですね。

Lyra2v2.cl の 71 行めにある DBUFLEN の大きさを変えると多少違うかもしれません。 0 から 48 の間で選べます。うちの開発マシンでは 4 が最適でした。

24 :ハッテン場五段:2016/04/22 18:50:36 (1年前)  0MONA/0人

R9 290x + OpenCL1.2 なマシンで動かしてみたら、DBUFLEN はどうも 32 が最適みたいな感じでした。環境によってかなり数字が変わるなー

25 :名前はまだ無い四段:2016/04/24 17:12:23 (1年前)  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 (1年前)  0MONA/0人

>>25
うーむ、すみませんが、ulong4 を ulong の配列に変えることと、4スレッドで動かすことの関連性がよくわからないです。申し訳ありません。
APU はともかく、GPU はベクトル型の扱いは苦手ですから ulong4 を ulong の配列に変えてしまえば多少早くなると思います。4スレッドで動かせばメモリアクセスの空白期間が少なくなってムダが減るのもわかります。でもその2つの関連性がよくわからないです。
ちなみに俺はいつも -g 4 で動かしてまして、当社比の数字も -g 4 を指定した時の数字を使っています。

27 :izuna五段:2016/04/24 19:51:15 (1年前)  0MONA/0人

>>25

ハッシュの計算って逐次で並列性の少ない計算だと思うのだが
並列性のないハッシュの計算を4スレッドで計算できるの?

28 :名前はまだ無い四段:2016/04/25 23:41:47 (1年前)  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 (1年前)  0MONA/0人

>>28
> ハッシュの計算は並列性が低いことが多いと思いますが…

XORや加算が多いからといって、並列性が高いとは限らない。
4スレッドで2倍の性能がでても、最初から性能が半分になっている。

もう少し、具体的に実装を考えてみないと、どうとも言えない。

Lyra2REv2で使われるハッシュ関数の、ほとんどは、SHA-3候補のなんだけど、これらを4スレッドで実行するの?

30 :ハッテン場五段:2016/04/26 02:48:06 (1年前)  0MONA/0人

>>28
よくわかんないんで実際に sgminer の OpenCL カーネルいじって見本みせて頂けますか?

31 :ハッテン場五段:2016/04/28 08:03:05 (1年前)  0MONA/0人

前に配ったカーネルの Lyra2v2.cl の reduceDuplexRowf2 関数内で変数 max を使った条件分岐がいくつかあります。が、手で動きをトレースしたところ、それら条件分岐は必ず真または必ず偽にしかなりようがないことが分かりました。
で、条件文を取り去って動かしたのですが、速さはまったく変わりませんでした。というわけで訂正したものの配布は必要なさそうなので行いませぬ。あしからず。

32 :ハッテン場五段:2016/05/01 05:59:51 (1年前)  1.114MONA/2人

当社比 25% アップ版ができました。イクイクー
http://kernel.hattenba.xyz/kernel-20160501.zip

・グローバルメモリへのアクセス回数を減らしました。
・ムダな条件分岐を取り払いました。
・あと何か変えたような気がします。

うちの R9 380 では 25% アップしましたが、R9 290x では 14% しか上がりませんでした。あしからず。

33 :なむやん七段教士:2016/05/01 22:37:22 (1年前)  0MONA/0人

10%以上アップとかすごすぎ
R7 250無印だとどんだけ上がるんだろう?

34 :ハッテン場五段:2016/05/02 04:28:02 (1年前)  0MONA/0人

>>33
もし試す機会がきたら動作報告よろしくお願いします!

35 :名無し名誉名人教士:2016/05/02 16:43:20 (1年前)  0.1MONA/1人

Radeon R9 290Xでみると、
1CUあたり、(スカラユニット×16)×4基に対して、レジスタファイル(Vector Refisters)が64kB×4基…
スカラユニット1つあたり4kBも確保できるから、ローカル変数で1536バイト確保できちゃう?

素人考えですまぬ…

36 :ハッテン場五段:2016/05/03 04:51:33 (1年前)  0MONA/0人

>>35
ローカルメモリ!
sgminer の本体のほうを手入れしないとできませんが、幸いにも何ヶ月か前に自分でコンパイルしてたりしたのが残ってまして再利用すれば試せるかも?です。

37 :ハッテン場五段:2016/05/03 04:53:29 (1年前)  0MONA/0人

問題は、./autogen.sh とか ./configure に色々オプションをつけて動かしてたのは覚えてるんですが、それらをすっかり完全に忘れてしまってることです。どこかにも書いてもないみたい。ぐぬぬ

38 :名無し名誉名人教士:2016/05/03 06:10:49 (1年前)  0MONA/0人

>>37
シェアードメモリでも速くなるかな?
(OpenCLにシェアードメモリってあったっけ?)

39 :名無し名誉名人教士:2016/05/03 06:15:21 (1年前)  0MONA/0人

ちなみに、ccminerのほうでreduceDuplexRowf2を試してみたけど、逆に遅くなった…
if分で相当遅くなるのかな?
(GTX970で改造前9.5MH/s、改造後8.5MH/s)

いま、if(rioneq2)無しで試してみる。

40 :名無し名誉名人教士:2016/05/03 07:13:38 (1年前)  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 (1年前)  0MONA/0人

sgminer をビルドしなおして、グローバルメモリからローカルメモリに移行して動かしてみたんですが・・・ダメダメでした。どうもメモリが足りなすぎるようです。

-w 16 くらいなら動くんですが、すごく遅い。メモリ割り当て量や割り当て方を再検討してみます。

42 :ハッテン場五段:2016/05/03 07:36:03 (1年前)  0MONA/0人

>>40
あの分岐は鬼門だったんですなw

43 :ハッテン場五段:2016/05/03 09:43:52 (1年前)  0MONA/0人

色んなサイト見てローカルメモリについて分かったこと

・ローカルメモリは L2 キャッシュに置かれる
・びりあるさんが twitter にて 2015 年 10 月 13 日に「1.5KB くらいしか使わないので、普通に L2 キャッシュあたりに載っちゃう量」だと発言している

てことはやっぱりローカルメモリを使うこと自体は正解で、俺のやり方が間違ってるのかな。再検討!

44 :ハッテン場五段:2016/05/03 10:33:31 (1年前)  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 (1年前)  0.1MONA/1人

>>32
ccminerでは、dbufによるメモリ分岐の最適化は逆効果でした…
GTX970で、改造前10.4MH/s、改造後8.5MH/s
Radeonでは効果はありましたか?

>>43
1536バイトを全部バラの変数にして(レジスタのみ使用しメモリを使用しないで)演算、関数も(引数はメモリを使用するから)inline化すれば…

かなり面倒くさいな…

46 :ハッテン場五段:2016/05/04 04:04:55 (1年前)  0MONA/0人

>>45
うちの Radeon どもでは数%ですが効果ありましたよー どうも Radeon では条件分岐にかかるコストよりもグローバルメモリへのアクセスにかかるコストのほうが若干大きいようです。
ゲフォはメモリにもっと余裕がありそうなんで、DBUFLEN を増やしてみてもいいかもです。

1536 バイトを全部バラしますかw 手間が恐ろしいですw
でも実際いじってみて感じたんですが、同じ量のメモリを確保するのでも、でかい配列を一度に確保するよりは、小出しに確保していくほうがプログラムがすんなり動いてくれますね。まるで尿道結石のようです。

47 :ハッテン場五段:2016/05/04 04:05:57 (1年前)  0MONA/0人

1ワークアイテムで使えるローカルメモリ量について。

ローカルメモリが x バイトあって ulong4 一つで 32 バイト使うから、1ワークグループで ulong4 が x/32 個使える。ということは sgminer のオプション -w で指定したワークサイズを n とすると、1ワークアイテムで使えるローカルメモリ上の ulong4 は x/(32*n) 個。

うちの R9 290x も R9 380 も APU どももどいつもこいつもローカルメモリの量は 32KB という表示が。ならば1ワークアイテムで使えるローカルメモリ上の 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 (1年前)  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 (1年前)  0.1MONA/1人

トピの内容は難しくてさっぱりですが、速度向上等々ありましたのでご報告とお礼。
今のところ >>32 版が一番あってる、うちの7950君(1.1M → 1.4M)
>>9 のときには 1.25M位でした。
>>48 は速度低下(1.4M → 1.2M)してしまいましたので戻しました。

こういった最適化による改善の余地があるなんて、そしてそれを施しているトピ主さまはすごい!
ありがとうございます~

50 :名無し名誉名人教士:2016/05/04 23:43:27 (1年前)  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 (1年前)  0MONA/0人

>>49
テストご協力ありがとうございます!俺は信者ですから広めるのは当然なのです。

>>50
勇者現る!
速度減は残念ですが、何か別のいいことにつながるといいですね
uint2 をお使いのようですが、そちらのデバイスでは uint では 2 が最適なベクトル長なんでしょうか?

52 :名無し名誉名人教士:2016/05/05 05:21:45 (1年前)  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 (1年前)  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 (1年前)  0MONA/0人

>>53の方法でやってみたら、思いのほか速くなったでござる…
(ccminerでやってみた)
GTX970:改造前10.4MH/s→改造後21MH/s

…ぱねぇ!
でも、コードミスったのかエラーがチョイチョイ出る。ちょっと見直し中。

55 :なむやん七段教士:2016/05/05 17:08:36 (1年前)  0.00000758MONA/1人

ふぁ!?倍!?打ち間違い?

56 :名無し名誉名人教士:2016/05/05 23:08:36 (1年前)  0MONA/0人

>>54の最適化結果
数値がよかったのはエラーのせいだったようだ…
方針はいいと思ったんだけどな…

57 :名無し名誉名人教士:2016/05/06 01:21:35 (1年前)  10MONA/1人

シェアードメモリで768バイト、グローバルメモリで768バイト
このパターンで作ってみた。

(ccminerでやってみた)
GTX970:改造前10.4MH/s



グローバルメモリ768バイトならL2キャッシュに収まる!
ここまで引き出せれば十分でしょ!

ハッテン場 様、参考までにソースをうpしておきます。(CUDAですが)
https://onedrive.live.com/redir?resid=C7ABE390AB1575E7!782&authkey=!AI84vQzotFa0_bQ&ithint=file,zip

58 :tomotomo9696七段教士:2016/05/06 02:12:58 (1年前)  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 (1年前)  0MONA/0人

>>57
素晴らしきプリケツ! もうちょっとして手が空いたらじっくり鑑賞します!

>>58
うちはまだ 15.12 使ってます。つーか 16.3.2 に一度アップグレードしたんですが、sgminer がうまく動いてくれなくて戻しちゃいました。新しいドライバだと前のやつのほうがうまく動くんですね。ご報告感謝です!

60 :ハッテン場五段:2016/05/06 08:46:20 (1年前)  0MONA/0人

>>57
じっくりねっとりと拝見しました。素晴らしいプリケツです。俺が書いた、関数を使ったバッファ?とは比べようもないエレガントさでした。

ただこちらの L2 キャッシュはどうも 32KB くらいしか確保ができないようで、1ワークアイテムあたり 768 バイトも確保するのは非常に難しそうです。せめて 384 バイトでも確保できればいいのですが。256 バイトがせいぜいかな。頭がいたい

61 :名無し名誉名人教士:2016/05/06 17:23:34 (1年前)  0MONA/0人

>>60
RadeonはTonga、Hawaii、FijiにはLocalDataShareMemoryが64kBあるようですが、使えないのでしょうか?

62 :名無し名誉名人教士:2016/05/06 23:15:02 (1年前)  10MONA/1人

>>57をさらに最適化してみた。もう、ゴールしてもいいよね…



シェアードメモリのバンクコンフリクトが面倒くせぇ…

ハッテン場 様、参考までにソースをうpしておきます。(シェアードメモリの参考に)
https://onedrive.live.com/redir?resid=C7ABE390AB1575E7!783&authkey=!ACvnu6myOBoE8ME&ithint=file,zip

63 :名無し名誉名人教士:2016/05/06 23:21:18 (1年前)  0MONA/0人

ちなみに>>53の最適化案はL2キャッシュがない場合は有効だと思うが、
シェアードメモリを併用したため、L2に入りきるサイズになったから、効果はないかもね…

64 :名無し名誉名人教士:2016/05/06 23:25:21 (1年前)  0MONA/0人

冷静に考えると、sgminerを最適化するなら、
同じやり方でccminerも最適化できるから、差が縮まらないんじゃね?

ってか、ccminerはCUDAのインラインアセンブラまで手を出しているから、そう簡単に差は縮まらないと思われる…

65 :名無し名誉名人教士:2016/05/07 03:19:16 (1年前)  0.1MONA/1人

>>62の最適化で1Block掘り当てた!!

66 :名無し名誉名人教士:2016/05/07 03:22:49 (1年前)  0.1MONA/1人

ついさっき、もう1block掘れた!!

67 :ハッテン場五段:2016/05/07 07:14:42 (1年前)  0MONA/0人

>>61
Tonga、Hawaii ともに試しましたが、32KB が限度のようです。むむむ。64KB 確保しようとすると sgminer 止まっちゃいます。

>>62
ありがとうございます! よく分からん部分が結構ありますが、努力します。

>>63
OpenCL で言えばスレッド化はサブデバイス化にあたるのかなー サブデバイスはノータッチだったんでもっと勉強しないと

>>64
こないだ見せていただいたやつは、俺にとって初めて見る CUDA ソースだったんですが、アセンブラっぽいの見て「あ、これずるい!」と思いましたw
確かに差は縮まらないかもしれませんが、ケツを貸しあった者同士の距離は縮まったんじゃないかなと思いますぜ

68 :名無し名誉名人教士:2016/05/07 16:49:11 (1年前)  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 (1年前)  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 (1年前)  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 (1年前)  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 (1年前)  0MONA/0人

>>71
これで解決できればいいけど…(結局原因がわからない…)
http://d.hatena.ne.jp/aont/20110504/1304498046

ちなみに、http://askmona.org/3853にて、ソロマイニング用ではありますが、Windows向けバイナリをアップしています。

73 :リキプロマン六段:2016/05/07 22:35:52 (1年前)  0MONA/0人

>>72
このエラーが発生している箇所がデバイス関数なので、cppには書くことが出来ないので無理ですね・・・
Windows+VisualStudio環境も用意しているので、こちらでコンパイルしてみますね。

74 :tomotomo9696七段教士:2016/05/07 23:41:48 (1年前)  1MONA/1人

>>59
ちなみに、今更ですが
当方もバージョンあげて動かなくなりましたが、
%TEMP% に使用.clファイルコピペで解決しました

75 :ハッテン場五段:2016/05/08 05:28:35 (1年前)  0MONA/0人

>>68
たしかに大きすぎるんですけど、64 とか 128 とか 256 じゃないとまともな速度で動いてくれんのです。

>>74
なんとそんな方法が!見るディレクトリが変わっちゃったのかな

76 :ハッテン場五段:2016/05/08 05:32:39 (1年前)  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 (1年前)  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 (1年前)  0.1MONA/1人

>>30
遅くなりましたが、今はテストに良さそうな環境が空いていませんし、テスト環境無しでさらっと書く自信は無いです。
(かなり手間がかかりそうで気が向かないというのが本音かもしれません)
流れとしてはlocal_idを取得してそれを元にどの入力のハッシュを処理するか
とulong4のどの次元を担当するかを決定する感じでしょうか。

ただ、AMDの資料を見てみるとwavefrontのCUでの実行のされ方などの関係もあり
1つのworkgroupでLDSを使い過ぎると効率が下がりそうですし、グローバルメモリへのアクセスが避けられそうになく
グローバルメモリでのデータ格納方法を見直す方が優先順位は高そうです。

64bitのデータではコアレスアクセスにならないという記述があったり、
分岐粒度のこともあったりとRadeonにとってはかなり不利なアルゴリズムの気がしてきました。

79 :名無し名誉名人教士:2016/05/09 21:34:16 (1年前)  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 (1年前)  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 (1年前)  1MONA/1人

>>80 つづき
ちなみに「32バイトで読み込む」のではなく、「8バイトで4回読み込む」にすると、
バンクコンフリクトは1/4となって同時アクセスは少なくなりますが、読み込み回数が4倍になるため、
結局かわらないという微妙な状況になってしまいます。
(OpenCLの方で32バイトを分割読み込みしている場合を除く。CUDAでは「16バイト2回読み込み」のため、バンクコンフリクトは1/4、読み込み回数2倍、となり、一応メリットはありました)

82 :ハッテン場五段:2016/05/12 03:57:32 (1年前)  0MONA/0人

>>78
テストならすぐできます。sgminer に添付されてる cl ファイルをテキストエディタでいじるだけですよ。

>>80
サンクスであります!
何かしらの単位ごとにメモリ配置をバラしていくんですね。やはりそこに手をつけずに済むことはありえない・・か!

ここ2,3日は SVM が使えないか奮闘しております。これが使えれば激速になるはず!でも SVM 使おうが使うまいが手元のローカルメモリの扱いも避けて通れそうにないので両方ともぼちぼちやっていきます。
GW 明けて仕事忙しくなってきて死にそうです。あと鯖たべたい

83 :ハッテン場五段:2016/05/15 11:06:24 (1年前)  6.114MONA/3人

はーい 当社比 100 %アップ版ができましたよー イクイクー
http://kernel.hattenba.xyz/kernel-20160515.zip

改造版の sgminer でしか動きませんが、ローカルメモリは使ってません。とりあえずグローバルメモリをどうにかするのが先じゃね?っと思いまして、名人たちからご教示いただいた方法を使ってみたら、2倍近くの速度がでるようになりました。ただし r9 290x だと 1.5 倍くらいにとどまってました。

実はそういう方法があるのは分かってはいたんですが、それをやるとどうして早くなるのか理解できず、納得できなくて着手していませんでした。なるほどメモリのバンク絡みの話だったんですな。

ローカルメモリの使い方は十分注意がいるようなので、今回はとりあえずローカルメモリはまったく使ってません。下手に使うと速度が落ちちゃう。

84 :なむやん七段教士:2016/05/15 11:08:40 (1年前)  0MONA/0人

ふぁーww
凄まじい性能アップですな!

85 :ハッテン場五段:2016/05/15 11:08:41 (1年前)  0MONA/0人

SVM は sgminer を 64bit 化しないと使えないのでコンパイルを通すべく奮闘してたのですが、うまくいかないのでとりあえず既存の CL ファイル周りをいじったらうまくいったので公開に至りました。

86 :名無し名誉名人教士:2016/05/18 11:33:18 (1年前)  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 (1年前)  0MONA/0人

>>86
ありがとうございます。情報サンクスです。

OpenCL でスレッド化ってどうするのか、方法がないのかいまいちよく分からないです。サブデバイス化というのはカーネルの局所局所をスレッド分岐するという技術ではなさそうです。今試しにサブデバイス化実験をやってるんですが clGetDeviceInfo でコア吐き出すという無間地獄中です。デイバス分割しても使えるメモリの量が減らないんなら分割し甲斐もあるのですが、それを調べる方法を使うと固まっててはどうしようもない。頭が痛いです。

あっちのスレは賑やかですなー こちらは寂しい限りです (´Д⊂グスン 無理も無いか

88 :名無し名誉名人教士:2016/05/19 08:36:28 (1年前)  0MONA/0人

>>87
私は単純に、スレッド数4倍にして動かしています。サブデバイス化(?)ではないはずです。

89 :ハッテン場五段:2016/05/20 03:03:18 (1年前)  0MONA/0人

>>88
コードを30回くらい読みなおしてようやく話が見えてきました!

raund_lyra した後に state を8バイト分ずらして state2 に xor する処理について、別スレッドからデータをもってきたりしてるんですね。

元のコードとどうして等価になるのかまだよく分からないのでもうちょっと読みなおしてきます。

90 :名無し名誉名人教士:2016/05/20 04:36:12 (1年前)  0MONA/0人

>>89
raund_lyraも別スレッドから持ってきていますよ~

91 :ハッテン場五段:2016/05/21 05:35:28 (1年前)  0MONA/0人

>>90
! 見落としまくりでした。
でやっぱりどうして以前のコードとスレッド化させたコードが等価になるのかよくわからないです。

元の中身がちがう、別スレッドの state の中身を手元にもってきてどうして計算できてしまうのだろうかとずっと考えています。

もしや 全てのスレッドに渡すもともとのデータはすべて同じなのでしょうか。いやいやもしそうだったらスレッド間でのデータやりとりするどころかまったく同じ計算をたくさんのスレッドでやらせる意味自体がないハズ。なんですが

92 :ハッテン場五段:2016/05/21 05:56:42 (1年前)  0MONA/0人

いや! もしやワークグループでは、元々与えられるデータは同一なのかな!? もしそうであれば話は通るし、何も不自然でないし、これまで見てきた OpenCL や CUDA との解説とも矛盾しない!

93 :名無し名誉名人教士:2016/05/21 10:10:20 (1年前)  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 (1年前)  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 (1年前)  0MONA/0人

>>93
ありがとうございます!
段々と頭がすっきりしてきました。朝日が上る日も近い!と感じてきます。

96 :名前はまだ無い四段:2016/05/21 23:46:01 (1年前)  0MONA/0人

>>82
そのテストに良さそうな環境が空いていないのです。

>>88
4並列で行っている部分はブロック数を4倍にしているのですね。

Gfunc_v35ではuint64_tのデータをuint2として処理しているようですが、
加算を上位32ビットと下位32ビットに分けて行うことになり問題が発生しないのでしょうか?

97 :名無し名誉名人教士:2016/05/22 01:00:21 (1年前)  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 (1年前)  0MONA/0人

>>97
桁あふれが反映されるのですか。
知らないとはまりそうですが、ulongをuint2として処理する場合は便利ですね。

99 :名無し名誉名人教士:2016/05/22 18:10:10 (1年前)  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 (1年前)  0MONA/0人

Monacoinウォレットがバージョンアップしたようです。
プールのウォレット更新がちょっと心配。
ソロマイニングはエラー報告が来ていたので、採掘方法が何か変わっている?

101 :ハッテン場五段:2016/05/24 06:28:12 (1年前)  0MONA/0人

>>100
え、ソロまたなんか採掘やりにくくなったんですか?

うちのプールの monacoind を更新しました。今んとこ問題ないようですが、ブロック見つけてモナを配分できるのを確認するまでは気が抜けませんです

102 :名無し名誉名人教士:2016/05/24 08:15:16 (1年前)  0MONA/0人

>>101
プールでは大丈夫って報告もあったんですが…どうなんでしょう?

103 :ハッテン場五段:2016/05/25 06:25:02 (1年前)  0MONA/0人

>>102
うちのプールはちゃんとブロック発見できて、ちゃんと分配できたようです。分配つか俺一人しかいませんけどw

104 :siv三段:2016/05/26 17:56:49 (1年前)  0MONA/0人

そういえば当初のハッシュレートと比べてどの程度向上したんですか?
具体的な値があるといいんですが...

105 :ハッテン場五段:2016/05/29 04:36:36 (1年前)  0.1MONA/1人

>>104
現状うちのマシンの R9 380 で言えば
3.755MH/s → 7.650MH/s
くらいにはなっとります。

106 :siv三段:2016/05/29 05:25:36 (1年前)  0MONA/0人

>>105
なるほど
初期と比べて劇的に改善されてますね!

107 :ハッテン場五段:2016/05/29 06:49:47 (1年前)  0MONA/0人

>>106
独力ではここまでは絶対にムリでした。
皆さんから頂いたアイデアとモナとご声援の賜物です。

108 :ハッテン場五段:2016/05/29 06:51:56 (1年前)  0MONA/0人

最近の技術的課題は、いかにスレッド間でデータをやりとりするか?ってところでつまずいてます。

CUDA には便利な命令があるようですが、OpenCL でとなると、ローカルメモリやグローバルメモリを使ってやり取りするほかないかな?と思ってます。でもこれがまたなかなかうまくいかないです (´・ω・`)

109 :名無し名誉名人教士:2016/05/29 09:37:06 (1年前)  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 (1年前)  0MONA/0人

>>108
ハッテン場さん…高速化、難航してます?
なかなかこっちでの発言が無いようなので…

111 :名無し四段:2016/06/13 04:14:31 (1年前)  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 (1年前)  0.00114114MONA/1人

久々にスレ一覧を見てたら素晴らしいスレで凄い高速化されてる・・・!!
自分も無駄な分岐を除くだけしてHD7950で5Mhs(15%向上?)くらいで回してましたが、もうとっくにそういった次元ではなさそうですね
rx480/rx470もそろそろ出るので期待してます!!
(Cはあんまり出来ないクソ雑魚ですが、もし役に立てそうであれば自分も協力します!)

113 :きさらぎ八段:2016/06/20 21:44:29 (1年前)  0MONA/0人

関係ない話ですがEthereumとDecredを同時に採掘できるマイナーは何で同時に採掘できるんでしょう?
monaとkumaを同時に採掘、とか出来ないものでしょうか

114 :なむやん七段教士:2016/06/20 22:23:54 (1年前)  0MONA/0人

kuma!
クマはオススメですよ(意味深

115 :名無し初段:2016/06/21 13:52:25 (1年前)  0MONA/0人

久々に参加してみたら
6950で1.7M うーむ

116 :siv三段:2016/06/28 20:16:14 (1年前)  0MONA/0人

そろそろRadeon RX4xx発売されますね
ハッシュレートどれぐらい出るんでしょうか

117 :ハッテン場五段:2016/06/29 06:44:30 (1年前)  59.63114114MONA/2人

ご無沙汰です。ハッテンです。難航というか、開発に取り掛かれない状況に陥っていましてみなさま申し訳ないです。5月末から時系列で申しますと・・

Vista を 7 にアップグレード!ドライバも更新して R9 290X が 7.5MH/s から 10.6MH/s に!ヒャッハー!
 ↓
地震保険おりた!保険金で monacoin 全力買い!
 ↓
不動産屋「そのアパート熊本地震で倒壊の危険があるから引っ越してね
 ↓
アパート片付け&引っ越す金もないけどモナは売りたくないので100円単位で金策
 ↓
ア゛ア゛ア゛ア゛ア゛ア゛ア゛ア゛! ←今ココ

ただでさえ仕事忙しい時期なのに死にそうです すみません うちの汚部屋が手強いです

118 :ハッテン場五段:2016/06/29 06:45:09 (1年前)  0MONA/0人

正直 RX 480 ほしいけどこんな状況じゃ買えないですね アハハ

119 :名無し四段:2016/06/29 23:07:19 (1年前)  0MONA/0人

>>117
熊本っすかw近所かもw
こっちはあまり被害無かったけど

120 :名無し四段:2016/07/03 00:57:23 (1年前)  0MONA/0人

RADEON RX460が75Wらしいので
常時起動して掘れるようになりたいなぁ。
ハッシュレート次第だけど

121 :なむやん七段教士:2016/07/03 03:22:06 (1年前)  0MONA/0人

あの爆熱御殿常連のRADEONが75Wだと・・?

122 :名無し四段:2016/07/03 09:09:48 (1年前)  0MONA/0人

GFの14nm FinFETだからね。
半分のサイズになれば、電力効率なら
2倍程度にはなるらしい。
実測で1.9倍だったみたいだが

123 :siv三段:2016/07/05 19:31:55 (1年前)  0MONA/0人

そういやいつかしらにNicehashでRadeonのLyra2rev2ハッシュレート大幅に改善されたんだっけ
6月2日だったかしら
自分のHD7970のハッシュレートは現在12.5MH/s

124 :リキプロマン六段:2016/07/09 12:59:19 (1年前)  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 (1年前)  0MONA/0人

Nicehashのバイナリ使ったらこんな感じですよ。
R9 390 一枚
i.imgur.com/zzNg11W

126 :名無し四段:2016/07/18 02:40:22 (1年前)  0MONA/0人

>>123
なんかエラーでちゃう・・・。

127 :名無し四段:2016/07/30 21:30:54 (1年前)  0MONA/0人

16.7.3にしたら、またハッシュレート上がった。

128 :名無しさん:2016/08/01 03:25:31 (1年前)  0MONA/0人

マイニングでさ AWSのGPUは採算どうなんだろう?

129 :名無し四段:2016/08/03 11:18:36 (1年前)  0.114114MONA/1人

Anniversary Updateしたら
コマンドプロンプトのフォントサイズ小さくされたし
Intel HD graphicsとRADEONの順番が逆になったのか
batのgpu-platform 1を0にしたら直ったよ。

130 :もがみん七段教士:2016/08/25 21:40:28 (1年前)  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 (1年前)  3.9MONA/1人

>>130
デフォルトではintensityが8とかなり低くなっているので、「-I」オプションで環境にあった値を設定する必要があります。
実行中でもメイン画面から「g」「i」で変更可能です。

132 :もがみん七段教士:2016/08/28 23:12:48 (1年前)  0MONA/0人

>>131
ありがとうございます!

133 :kenichiA380初段:2016/10/22 04:31:53  0MONA/0人

Sapphire R9 390X Nitro使ってます
Iは19以上にすると不安定
OCはコア1200 メモリー1700 電圧+75mV 電力制限は+50%
でハッシュレートは10.64MH/sなんですがこれは正常なんですかね

134 :kenichiA380初段:2016/11/01 02:35:56  0MONA/0人

Radeon Crimson 16.10.3では掘れないみたいですね...

135 :kenichiA380初段:2016/11/01 03:33:55  0MONA/0人

>>134
掘れました

136 :ハッテン場五段:2016/11/25 04:12:04  0MONA/0人

久しぶりに閲覧

グラボ片っ端から売っちゃったり切手買っては転売したりしてて開発環境まるごとふっとんじゃいました。今あるのは R9 290x 一枚だけ。電気代もこれ以上増やせないくらいに生活が厳しいです。ごめんなさいです。

オークモナーの商品が売れてくれれば・・まあメシ代くらいにはなるか

137 :名無し名誉名人教士:2016/11/25 08:42:42  1.14MONA/1人

>>136
こっちの最適化にも手を出してみたいけど、開発環境がイマイチ作れない…
コンパイルするにも一部ソースが欠けているみたい…
私にはsgminerは難易度が高いのだろうか…
(そもそも改造しかできない私はプログラマーとは呼べない半端者なわけだが…)

138 :ハッテン場五段:2016/11/25 08:52:43  0MONA/0人

>>137
メモリ管理システムを根本から変えるならともかく、メモリへのデータの出し入れや計算順序を最適化するだけならコンパイラは不要ですぜ。

kernel ディレクトリにある cl ファイルをテキストエディタでいじって、sgminer を起動するだけです。起動時に cl ファイルをコンパイルしてくれます。

139 :名無し名誉名人教士:2016/11/25 12:42:39  0.00114114MONA/1人

>>138
いや、三段打ちメソッド(命名、izuna氏?>>5>>6)とか、ソロマイニングへの対応とか…いろいろやってみたいじゃない?
…私の技量でできるかはともかくとして…

140 :Defekt名人四段錬士:2016/11/29 23:49:47  0MONA/0人

俺のGTX960Mは一晩1モナ

141 :名無しさん:2017/05/02 04:43:58  0MONA/0人

ハッテン場様、

初めまして。突然のお願いで大変恐縮なのですが、ハッテン場様が
lyra2rev2カーネルに追加されたコードのライセンス販売、もしくは権利の
販売にご興味がおありでしたら以下のメールアドレスにご連絡を
いただけないでしょうか。

zawawawawawawawawa@gmail.com

以上、よろしくお願いいたします。スレ汚し大変失礼しました。

zawawa @ bitcointalk.org

お気に入り

新規登録してMONAをもらおう

登録すると、投稿したり、MONAをもらったりすることができます。質問したり、答えたりしてMONAを手に入れてください。

新規登録ログイン