16 Commits

Author SHA1 Message Date
Yutaka Sawada
6dd7949030 Release note of version 1.3.3.1 2023-11-11 13:11:15 +09:00
Yutaka Sawada
a27f8221cd Update installer usage 2023-11-11 12:50:47 +09:00
Yutaka Sawada
8cff776c5e Update installer usage 2023-11-11 12:49:45 +09:00
Yutaka Sawada
bae9e8a0d8 Add files via upload 2023-11-11 12:49:05 +09:00
Yutaka Sawada
bbfad5b9df Add files via upload 2023-11-11 12:48:13 +09:00
Yutaka Sawada
fd24693c6b Update installer usage 2023-11-11 11:54:12 +09:00
Yutaka Sawada
cc9d3595bd Update installer usage 2023-11-11 11:53:34 +09:00
Yutaka Sawada
1b397d8976 Update version number 2023-11-07 12:53:49 +09:00
Yutaka Sawada
8c06ad76b6 Change max buffer size 2023-10-29 16:57:15 +09:00
Yutaka Sawada
50b735d3a5 Update PAR2 clients 2023-10-29 16:56:04 +09:00
Yutaka Sawada
5660fcf7c5 Improve GPU function a little 2023-10-27 13:31:56 +09:00
Yutaka Sawada
e979c07600 Update PAR2 clients 2023-10-27 13:31:18 +09:00
Yutaka Sawada
cd7d1f9450 Fixed a bug in GPU function 2023-10-25 14:06:06 +09:00
Yutaka Sawada
cdaceef840 Fixed a bug 2023-10-25 14:05:16 +09:00
Yutaka Sawada
6ce606977b Update PAR2 clients 2023-10-23 10:57:51 +09:00
Yutaka Sawada
fb72e811d0 Improve GPU function 2023-10-23 10:54:28 +09:00
32 changed files with 1150 additions and 605 deletions

View File

@@ -1,6 +1,6 @@
# MultiPar # MultiPar
### v1.3.3.0 is public ### v1.3.3.1 is public
  This is a testing version to improve speed of PAR2 calculation.   This is a testing version to improve speed of PAR2 calculation.
Because the new method isn't tested so much, there may be a bug, failure, or mistake. Because the new method isn't tested so much, there may be a bug, failure, or mistake.
@@ -9,60 +9,55 @@ When you don't want to test by yourself, you should not use this yet.
If you see a problem, please report the incident. If you see a problem, please report the incident.
I will try to solve as possible as I can. I will try to solve as possible as I can.
  The PAR2 calculation speed may be 10% ~ 50% faster than old version.   CPU's L3 cache optimization depends on hardware environment.
The optimization depends on hardware environment. It's difficult to guess the best setting for unknown type.
I don't know what is the best setting on which PC. It seems to work well on Intel and AMD 's most CPUs.
From [many tests of debug versions](https://github.com/Yutaka-Sawada/MultiPar/issues/99), Thanks Anime Tosho and MikeSW17 for long tests.
it will select maybe better setting automatically. But, I'm not sure the perfomance of rare strange kind CPUs.
Thanks testers for many trials. If you want to compare speed of different settings on your CPU,
If you want to compare speed of different settings on your PC, you may try those debug versions. you may try samples (TestBlock_2023-08-31.zip) in "MultiPar_sample" folder
on [OneDrive](https://1drv.ms/u/s!AtGhNMUyvbWOg0cF2UHcs709Icv4).
  I changed GPU implementation largely, too.   I improved GPU implementation very much.
To adopt CPU optimization, it will process smaller tasks on GPU. Thanks [Slava46 and K2M74 for many tests](https://github.com/Yutaka-Sawada/MultiPar/issues/99).
Because GPU don't use CPU's cache, it's inefficient for GPU's task. While I almost gave up to increase speed, their effort encouraged me to try many ways.
I don't know that new method is faster than old version or not. Without their aid, I could not implement this GPU function.
OpenCL perfomance is varied in every graphics boards.
If you have a fast graphics board, enabling "GPU acceleration" would be faster.
If it's not so fast (or is slow) on your PC, just un-check the feature.
Threshold to use GPU:   I saw a new feature of Inno Setup 6, which changes install mode.
- Data size must be larger than 200 MB. It shows a dialog to ask which install mode.
- Block size must be larger than 64 KB. Then, a user can install MultiPar in "Program Files" directory by selecting "Install for all users".
- Number of source blocks must be more than 192. This method may be easier than starting installer by "Run as administrator".
- Number of recovery blocks must be more than 8. I test the selection dialog at this version.
If there is no problem nor complaint from users, I use this style in later versions, too.
  Because [a user requested](https://github.com/Yutaka-Sawada/MultiPar/issues/102),
I implemented a way to add 5th item in "Media size" on Create window.
Write this line `MediaList4=name:size` under `[Option]` section in `MultiPar.ini`.
Currently, you cannot change the item on Option window.
[ Changes from 1.3.2.9 to 1.3.3.0 ] [ Changes from 1.3.3.0 to 1.3.3.1 ]
GUI update Installer update
- Change - It shows dialog to select "per user" or "per machine" installation.
- Option adapted to new "lc" settings.
- It's possible to add 5th item in "Media size" on Create window.
PAR2 client update PAR2 client update
- Change - Change
- Max number of using threads is increased to 32. - Max number of threads to read files on SSD was increased to 6.
- Threshold to use GPU was decreased.
- Improvement - Improvement
- Matrix inversion may use more threads. - GPU acceleration would become faster.
- L3 cache optimization was improved for recent CPUs.
[ Hash value ] [ Hash value ]
MultiPar1330.zip MultiPar1331.zip
MD5: 79570F84B74ECF8E5100561F7AAC3803 MD5: ECFC1570C839DD30A2492A7B05C2AD6E
SHA1: ACF7F164001708789C5D94003ED6B5C172235D54 SHA1: 5E0E4CC38DAA995294A93ECA10AEB3AE84596170
MultiPar1330_setup.exe MultiPar1331_setup.exe
MD5: D1F1A5A4DF1C9EDD698C9A017AF31039 MD5: A55E6FA5A6853CB42E3410F35706BAD9
SHA1: 4C3314B909572A303EBBE8E015A2E813841CFA33 SHA1: 8D46BD6702E82ABA9ACCFA5223B2763B4DCEFE9E
  To install under "Program Files" or "Program Files (x86)" directory,   To install under "Program Files" or "Program Files (x86)" directory,
you must start the installer with administrative privileges by selecting you must select "Install for all users" at the first dialog.
"Run as administrator" on right-click menu.
  Old versions and source code packages are available at   Old versions and source code packages are available at
[GitHub](https://github.com/Yutaka-Sawada/MultiPar/releases) or [GitHub](https://github.com/Yutaka-Sawada/MultiPar/releases) or

Binary file not shown.

View File

@@ -71,14 +71,13 @@ Don't send current PAR3 files to others, who may not have the same version.
[ How to install or uninstall with installer package ] [ How to install or uninstall with installer package ]
Double click setup file ( MultiPar131_setup.exe or something like this name ), Double click setup file ( MultiPar133_setup.exe or something like this name ),
and follow the installer dialog. and follow the installer dialog.
At version up, if you want to use previous setting, overwrite install is possible. At version up, if you want to use previous setting, overwrite install is possible.
Before overwrite install, you should un-check "Integrate MultiPar into Shell". Before overwrite install, you should un-check "Integrate MultiPar into Shell".
You may need to re-start OS after overwrite install or uninstall rarely. You may need to re-start OS after overwrite install or uninstall rarely.
To install under "Program Files" or "Program Files (x86)" directory, To install under "Program Files" or "Program Files (x86)" directory,
you must start the installer with administrative privileges by selecting you must select "Install for all users" at the first dialog.
"Run as administrator" on right-click menu.
You can uninstall through the Windows OS's Control Panel, You can uninstall through the Windows OS's Control Panel,
or double click unins000.exe in a folder which MultiPar was installed. or double click unins000.exe in a folder which MultiPar was installed.
@@ -100,7 +99,7 @@ In either case, user made icons and association are available for the user only.
[ How to install with archive version ] [ How to install with archive version ]
Unpack compressed file ( MultiPar131.zip or something like this name ) in a folder. Unpack compressed file ( MultiPar133.zip or something like this name ) in a folder.
MultiPar.exe is the interface of MultiPar. MultiPar.exe is the interface of MultiPar.
You can create short-cut icon or send-to link at Option window later. You can create short-cut icon or send-to link at Option window later.

View File

@@ -107,12 +107,14 @@ PAR 3.0 仕様のフォーマットは細部が流動的で最終版との互換
[ インストーラー版のインストールとアンインストール ] [ インストーラー版のインストールとアンインストール ]
 インストーラー ( MultiPar131_setup.exe みたいな名前 ) をダブル・クリックすると、  インストーラー ( MultiPar133_setup.exe みたいな名前 ) をダブル・クリックすると、
インストール画面が表示されるので、その指示に従ってください。 インストール画面が表示されるので、その指示に従ってください。
バージョン・アップ時に、設定項目をそのまま使いたい時は上書きインストールしてもいいです。 バージョン・アップ時に、設定項目をそのまま使いたい時は上書きインストールしてもいいです。
上書きインストールする前に、「MultiPar をシェルに統合する」のチェックを外してください。 上書きインストールする前に、「MultiPar をシェルに統合する」のチェックを外してください。
上書きインストールやアンインストール後に OS の再起動を求められるかもしれません。 上書きインストールやアンインストール後に OS の再起動を求められるかもしれません。
「Program Files」や「Program Files (x86)」内にインストールするには、 「Program Files」や「Program Files (x86)」内にインストールするには、
最初のダイアログで「すべてのユーザー用にインストール」を選んでください。
右クリック・メニューの「管理者として実行」を選んで 右クリック・メニューの「管理者として実行」を選んで
管理者権限でインストーラーを開始する必要があります。 管理者権限でインストーラーを開始する必要があります。
@@ -137,7 +139,7 @@ MultiPar をインストールしたフォルダ内の unins000.exe をダブル
[ アーカイブ版のインストール ] [ アーカイブ版のインストール ]
 配布されてる圧縮ファイル ( MultiPar131.zip みたいな名前 ) を解凍してできたファイルを  配布されてる圧縮ファイル ( MultiPar133.zip みたいな名前 ) を解凍してできたファイルを
どこか適当なフォルダに全て入れてください。 どこか適当なフォルダに全て入れてください。
その中の MultiPar.exe というのが MultiPar の実行ファイルです。 その中の MultiPar.exe というのが MultiPar の実行ファイルです。
それをダブル・クリックすると MultiPar が起動します。 それをダブル・クリックすると MultiPar が起動します。

View File

@@ -1,5 +1,19 @@
Release note of v1.3.3 tree Release note of v1.3.3 tree
[ Changes from 1.3.3.0 to 1.3.3.1 ] (2023/11/11)
Installer update
It shows dialog to select "per user" or "per machine" installation.
PAR2 client update
Change
Max number of threads to read files on SSD was increased to 6.
Improvement
GPU acceleration would become faster.
/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
[ Changes from 1.3.2.9 to 1.3.3.0 ] (2023/10/10) [ Changes from 1.3.2.9 to 1.3.3.0 ] (2023/10/10)
GUI update GUI update

View File

@@ -1,4 +1,4 @@
v1.3.3 の更新情報 (2023/10/10) v1.3.3 の更新情報 (2023/11/11)
 まだ動作実験中ですので、不安な人は前のバージョンを使ってください。  まだ動作実験中ですので、不安な人は前のバージョンを使ってください。
@@ -6,6 +6,7 @@
・クライアントの変更点 ・クライアントの変更点
 CPU Cache の利用方法を改善して速くなりました。  CPU Cache の利用方法を改善して速くなりました。
 GPU による高速化も速くなりました。
/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */

View File

@@ -51,7 +51,7 @@ There are command-line manuals in "<tt>help</tt>" folder.
</table> </table>
<hr> <hr>
<small>last update 2023/06/13 for version 1.3.2.9</small> <small>last update 2023/11/11 for version 1.3.3.1</small>
</body> </body>
</html> </html>

View File

@@ -8,14 +8,13 @@
<h3>Install or uninstall with installer package</h3> <h3>Install or uninstall with installer package</h3>
<p>&nbsp <p>&nbsp
Double click setup file ( <tt>MultiPar131_setup.exe</tt> or something like this name ), Double click setup file ( <tt>MultiPar133_setup.exe</tt> or something like this name ),
and follow the installer dialog. and follow the installer dialog.
At version up, if you want to use previous setting, overwrite install is possible. At version up, if you want to use previous setting, overwrite install is possible.
Before overwrite install, you should un-check "Integrate MultiPar into Shell". Before overwrite install, you should un-check "Integrate MultiPar into Shell".
You may need to re-start OS after overwrite install or uninstall rarely. You may need to re-start OS after overwrite install or uninstall rarely.
To install under "<tt>Program Files</tt>" or "<tt>Program Files (x86)</tt>" directory, To install under "<tt>Program Files</tt>" or "<tt>Program Files (x86)</tt>" directory,
you must start the installer with administrative privileges by selecting you must select "Install for all users" at the first dialog.
"Run as administrator" on right-click menu.
</p> </p>
<p>&nbsp <p>&nbsp
You can uninstall through the Windows OS's Control Panel, You can uninstall through the Windows OS's Control Panel,
@@ -42,7 +41,7 @@ In either case, user made icons and association are available for the user only.
<h3>Install with archive version</h3> <h3>Install with archive version</h3>
<p>&nbsp <p>&nbsp
Unpack compressed file ( <tt>MultiPar131.zip</tt> or something like this name ) in a folder. Unpack compressed file ( <tt>MultiPar133.zip</tt> or something like this name ) in a folder.
<tt>MultiPar.exe</tt> is the interface of MultiPar. <tt>MultiPar.exe</tt> is the interface of MultiPar.
</p> </p>
<p>&nbsp <p>&nbsp

View File

@@ -51,7 +51,7 @@
</table> </table>
<hr> <hr>
<small><EFBFBD>ŏI<EFBFBD>X<EFBFBD>V 2023/02/27 (<28>o<EFBFBD>[<5B>W<EFBFBD><57><EFBFBD><EFBFBD> 1.3.2.8)</small> <small><EFBFBD>ŏI<EFBFBD>X<EFBFBD>V 2023/11/11 (<28>o<EFBFBD>[<5B>W<EFBFBD><57><EFBFBD><EFBFBD> 1.3.3.1)</small>
</body> </body>
</html> </html>

View File

@@ -8,14 +8,13 @@
<h3><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>[<5B>ł̃C<CC83><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>ƃA<C683><41><EFBFBD>C<EFBFBD><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><></h3> <h3><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>[<5B>ł̃C<CC83><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>ƃA<C683><41><EFBFBD>C<EFBFBD><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><></h3>
<p>&nbsp <p>&nbsp
<EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>[ ( <tt>MultiPar131_setup.exe</tt> <20>݂<EFBFBD><DD82><EFBFBD><EFBFBD>Ȗ<EFBFBD><C896>O ) <20><><EFBFBD>_<EFBFBD>u<EFBFBD><75><EFBFBD>E<EFBFBD>N<EFBFBD><4E><EFBFBD>b<EFBFBD>N<EFBFBD><4E><EFBFBD><EFBFBD><EFBFBD>ƁA <EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>[ ( <tt>MultiPar133_setup.exe</tt> <20>݂<EFBFBD><DD82><EFBFBD><EFBFBD>Ȗ<EFBFBD><C896>O ) <20><><EFBFBD>_<EFBFBD>u<EFBFBD><75><EFBFBD>E<EFBFBD>N<EFBFBD><4E><EFBFBD>b<EFBFBD>N<EFBFBD><4E><EFBFBD><EFBFBD><EFBFBD>ƁA
<EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD>ʂ<EFBFBD><CA82>\<5C><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>̂ŁA<C581><41><EFBFBD>̎w<CC8E><77><EFBFBD>ɏ]<5D><><EFBFBD>Ă<EFBFBD><C482><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B <EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD>ʂ<EFBFBD><CA82>\<5C><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>̂ŁA<C581><41><EFBFBD>̎w<CC8E><77><EFBFBD>ɏ]<5D><><EFBFBD>Ă<EFBFBD><C482><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B
<EFBFBD>o<EFBFBD>[<5B>W<EFBFBD><57><EFBFBD><EFBFBD><EFBFBD>E<EFBFBD>A<EFBFBD>b<EFBFBD>v<EFBFBD><76><EFBFBD>ɁA<C981>ݒ荀<DD92>ڂ<EFBFBD><DA82><EFBFBD><EFBFBD>̂܂܎g<DC8E><67><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>͏<CD8F><E38F91><EFBFBD>C<EFBFBD><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD>Ă<EFBFBD><C482><EFBFBD><EFBFBD><EFBFBD><EFBFBD>ł<EFBFBD><C582>B <EFBFBD>o<EFBFBD>[<5B>W<EFBFBD><57><EFBFBD><EFBFBD><EFBFBD>E<EFBFBD>A<EFBFBD>b<EFBFBD>v<EFBFBD><76><EFBFBD>ɁA<C981>ݒ荀<DD92>ڂ<EFBFBD><DA82><EFBFBD><EFBFBD>̂܂܎g<DC8E><67><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>͏<CD8F><E38F91><EFBFBD>C<EFBFBD><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD>Ă<EFBFBD><C482><EFBFBD><EFBFBD><EFBFBD><EFBFBD>ł<EFBFBD><C582>B
<EFBFBD><EFBFBD><EFBFBD><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>O<EFBFBD>ɁA<C981>uMultiPar <20><><EFBFBD>V<EFBFBD>F<EFBFBD><46><EFBFBD>ɓ<EFBFBD><C993><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>v<EFBFBD>̃`<60>F<EFBFBD>b<EFBFBD>N<EFBFBD><4E><EFBFBD>O<EFBFBD><4F><EFBFBD>Ă<EFBFBD><C482><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B <EFBFBD><EFBFBD><EFBFBD><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>O<EFBFBD>ɁA<C981>uMultiPar <20><><EFBFBD>V<EFBFBD>F<EFBFBD><46><EFBFBD>ɓ<EFBFBD><C993><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>v<EFBFBD>̃`<60>F<EFBFBD>b<EFBFBD>N<EFBFBD><4E><EFBFBD>O<EFBFBD><4F><EFBFBD>Ă<EFBFBD><C482><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B
<EFBFBD><EFBFBD><EFBFBD><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD>A<EFBFBD><41><EFBFBD>C<EFBFBD><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD><EFBFBD> OS <20>̍ċN<C48B><4E><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>߂<EFBFBD><DF82><EFBFBD><EFBFBD><EFBFBD><E982A9><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>܂<EFBFBD><DC82><EFBFBD><EFBFBD>B <EFBFBD><EFBFBD><EFBFBD><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD>A<EFBFBD><41><EFBFBD>C<EFBFBD><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD><EFBFBD> OS <20>̍ċN<C48B><4E><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>߂<EFBFBD><DF82><EFBFBD><EFBFBD><EFBFBD><E982A9><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>܂<EFBFBD><DC82><EFBFBD><EFBFBD>B
<EFBFBD>u<tt>Program Files</tt><EFBFBD>v<EFBFBD><EFBFBD><EFBFBD>u<tt>Program Files (x86)</tt><EFBFBD>v<EFBFBD><EFBFBD><EFBFBD>ɃC<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>ɂ́A <EFBFBD>u<tt>Program Files</tt><EFBFBD>v<EFBFBD><EFBFBD><EFBFBD>u<tt>Program Files (x86)</tt><EFBFBD>v<EFBFBD><EFBFBD><EFBFBD>ɃC<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>ɂ́A
<EFBFBD>E<EFBFBD>N<EFBFBD><EFBFBD><EFBFBD>b<EFBFBD>N<EFBFBD>E<EFBFBD><EFBFBD><EFBFBD>j<EFBFBD><EFBFBD><EFBFBD>[<5B>́u<CC81>Ǘ<EFBFBD><C797>҂Ƃ<D282><C682>Ď<EFBFBD><C48E>s<EFBFBD>v<EFBFBD><76><EFBFBD>I<EFBFBD><49><EFBFBD><EFBFBD> <EFBFBD>ŏ<EFBFBD><EFBFBD>̃_<EFBFBD>C<EFBFBD>A<EFBFBD><EFBFBD><EFBFBD>O<EFBFBD>Łu<EFBFBD><EFBFBD><EFBFBD>ׂẴ<EFBFBD><EFBFBD>[<5B>U<EFBFBD>[<5B>p<EFBFBD>ɃC<C983><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>v<EFBFBD><76><EFBFBD>I<EFBFBD><49><EFBFBD>ł<EFBFBD><C582><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B
<EFBFBD>Ǘ<EFBFBD><EFBFBD>Ҍ<EFBFBD><EFBFBD><EFBFBD><EFBFBD>ŃC<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD>[<5B><><EFBFBD>J<EFBFBD>n<EFBFBD><6E><EFBFBD><EFBFBD><EFBFBD>K<EFBFBD>v<EFBFBD><76><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>܂<EFBFBD><DC82>B
</p> </p>
<p>&nbsp <p>&nbsp
<EFBFBD>A<EFBFBD><EFBFBD><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD> Windows OS <20>̃R<CC83><52><EFBFBD>g<EFBFBD><67><EFBFBD>[<5B><><EFBFBD>E<EFBFBD>p<EFBFBD>l<EFBFBD><6C><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>s<EFBFBD><73><EFBFBD><EFBFBD><EFBFBD>A <EFBFBD>A<EFBFBD><EFBFBD><EFBFBD>C<EFBFBD><EFBFBD><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><><EFBFBD><EFBFBD> Windows OS <20>̃R<CC83><52><EFBFBD>g<EFBFBD><67><EFBFBD>[<5B><><EFBFBD>E<EFBFBD>p<EFBFBD>l<EFBFBD><6C><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>s<EFBFBD><73><EFBFBD><EFBFBD><EFBFBD>A
@@ -43,7 +42,7 @@ MultiPar
<h3><EFBFBD>A<EFBFBD>[<5B>J<EFBFBD>C<EFBFBD>u<EFBFBD>ł̃C<CC83><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><></h3> <h3><EFBFBD>A<EFBFBD>[<5B>J<EFBFBD>C<EFBFBD>u<EFBFBD>ł̃C<CC83><43><EFBFBD>X<EFBFBD>g<EFBFBD>[<5B><></h3>
<p>&nbsp <p>&nbsp
<EFBFBD>z<EFBFBD>z<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>Ă鈳<EFBFBD>k<EFBFBD>t<EFBFBD>@<40>C<EFBFBD><43> ( <tt>MultiPar131.zip</tt> <20>݂<EFBFBD><DD82><EFBFBD><EFBFBD>Ȗ<EFBFBD><C896>O ) <20><><EFBFBD>𓀂<EFBFBD><F0938082>Ăł<C482><C582><EFBFBD><EFBFBD>t<EFBFBD>@<40>C<EFBFBD><43><EFBFBD><EFBFBD> <EFBFBD>z<EFBFBD>z<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>Ă鈳<EFBFBD>k<EFBFBD>t<EFBFBD>@<40>C<EFBFBD><43> ( <tt>MultiPar133.zip</tt> <20>݂<EFBFBD><DD82><EFBFBD><EFBFBD>Ȗ<EFBFBD><C896>O ) <20><><EFBFBD>𓀂<EFBFBD><F0938082>Ăł<C482><C582><EFBFBD><EFBFBD>t<EFBFBD>@<40>C<EFBFBD><43><EFBFBD><EFBFBD>
<EFBFBD>ǂ<EFBFBD><EFBFBD><EFBFBD><EFBFBD>K<EFBFBD><EFBFBD><EFBFBD>ȃt<EFBFBD>H<EFBFBD><EFBFBD><EFBFBD>_<EFBFBD>ɑS<EFBFBD>ē<EFBFBD><EFBFBD><EFBFBD><EFBFBD>Ă<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B <EFBFBD>ǂ<EFBFBD><EFBFBD><EFBFBD><EFBFBD>K<EFBFBD><EFBFBD><EFBFBD>ȃt<EFBFBD>H<EFBFBD><EFBFBD><EFBFBD>_<EFBFBD>ɑS<EFBFBD>ē<EFBFBD><EFBFBD><EFBFBD><EFBFBD>Ă<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>B
<EFBFBD><EFBFBD><EFBFBD>̒<EFBFBD><EFBFBD><EFBFBD> <tt>MultiPar.exe</tt> <20>Ƃ<EFBFBD><C682><EFBFBD><EFBFBD>̂<EFBFBD> MultiPar <20>̎<EFBFBD><CC8E>s<EFBFBD>t<EFBFBD>@<40>C<EFBFBD><43><EFBFBD>ł<EFBFBD><C582>B <EFBFBD><EFBFBD><EFBFBD>̒<EFBFBD><EFBFBD><EFBFBD> <tt>MultiPar.exe</tt> <20>Ƃ<EFBFBD><C682><EFBFBD><EFBFBD>̂<EFBFBD> MultiPar <20>̎<EFBFBD><CC8E>s<EFBFBD>t<EFBFBD>@<40>C<EFBFBD><43><EFBFBD>ł<EFBFBD><C582>B
<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>_<EFBFBD>u<EFBFBD><EFBFBD><EFBFBD>E<EFBFBD>N<EFBFBD><EFBFBD><EFBFBD>b<EFBFBD>N<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD> MultiPar <20><><EFBFBD>N<EFBFBD><4E><EFBFBD><EFBFBD><EFBFBD>܂<EFBFBD><DC82>B <EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD>_<EFBFBD>u<EFBFBD><EFBFBD><EFBFBD>E<EFBFBD>N<EFBFBD><EFBFBD><EFBFBD>b<EFBFBD>N<EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD><EFBFBD> MultiPar <20><><EFBFBD>N<EFBFBD><4E><EFBFBD><EFBFBD><EFBFBD>܂<EFBFBD><DC82>B

View File

@@ -50,7 +50,7 @@
</table> </table>
<hr> <hr>
<small>最新更新于2023年2月27适用于1.3.2.8版本,简体中文化 Deng Shiqing</small> <small>最新更新于2023年11月11适用于1.3.3.1版本,简体中文化 Deng Shiqing</small>
</body> </body>
</html> </html>

View File

@@ -8,7 +8,7 @@
<h3>使用安装包安装或卸载</h3> <h3>使用安装包安装或卸载</h3>
<p>&nbsp <p>&nbsp
双击安装文件(<tt>MultiPar131_setup.exe</tt>或类似名称文件),然后按照安装程序对话框进行操作。在版本升级时,如果要使用先前的设置,可以进行覆盖安装。在覆盖安装之前, 应取消勾选“将MultiPar整合到右键菜单”。在写入安装或卸载之后您可能需要重新启动操作系统。“<tt>Program Files</tt>”或“<tt>Program Files (x86)</tt>”目录下,必须在右键菜单上选择“以管理员身份运行”,用管理员权限启动安装程序 双击安装文件(<tt>MultiPar133_setup.exe</tt>或类似名称文件),然后按照安装程序对话框进行操作。在版本升级时,如果要使用先前的设置,可以进行覆盖安装。在覆盖安装之前, 应取消勾选“将MultiPar整合到右键菜单”。在写入安装或卸载之后您可能需要重新启动操作系统。要在<tt>Program Files</tt>”或“<tt>Program Files (x86)</tt>”目录下安装,您必须在第一个对话框中选择“为所有用户安装”
</p> </p>
<p>&nbsp <p>&nbsp
您可以通过Windows操作系统的控制面板卸载程序或双击MultiPar安装文件夹中的<tt>unins000.exe</tt>。由于卸载程序不会删除设置文件或安装后新添加的文件,因此您可以自行删除它们。 您可以通过Windows操作系统的控制面板卸载程序或双击MultiPar安装文件夹中的<tt>unins000.exe</tt>。由于卸载程序不会删除设置文件或安装后新添加的文件,因此您可以自行删除它们。
@@ -26,7 +26,7 @@
<h3>使用压缩包安装</h3> <h3>使用压缩包安装</h3>
<p>&nbsp <p>&nbsp
在文件夹中解压压缩文件(<tt>MultiPar131.zip</tt>或类似名称文件)。 在文件夹中解压压缩文件(<tt>MultiPar133.zip</tt>或类似名称文件)。
<tt>MultiPar.exe</tt>是MultiPar的启动程序。 <tt>MultiPar.exe</tt>是MultiPar的启动程序。
</p> </p>
<p>&nbsp <p>&nbsp

Binary file not shown.

Binary file not shown.

View File

@@ -1,4 +1,4 @@
[ par2j.exe - version 1.3.3.0 or later ] [ par2j.exe - version 1.3.3.1 or later ]
Type "par2j.exe" to see version, test integrity, and show usage below. Type "par2j.exe" to see version, test integrity, and show usage below.
@@ -367,10 +367,10 @@ The format is "/lc#" (# is from 1 to 32 as the number of using threads).
253: It uses 3/4 number of physical Cores. 253: It uses 3/4 number of physical Cores.
254: It uses one less threads than number of physical Cores. 254: It uses one less threads than number of physical Cores.
0: It uses the number of physical Cores. 0: It uses the number of physical Cores.
255: It uses one more threads than number of physical Cores. 255: It tries to use more threads than number of physical Cores.
You may set additional combinations; You may set additional combinations;
+1024 to disable CLMUL (and use old SSSE3 code), +1024 to disable CLMUL (and use slower SSSE3 code),
+2048 to disable JIT (for SSE2), +2048 to disable JIT (for SSE2),
+4096 to disable SSSE3, +4096 to disable SSSE3,
+8192 to disable AVX2, +8192 to disable AVX2,

View File

@@ -1,5 +1,5 @@
// common2.c // common2.c
// Copyright : 2023-09-23 Yutaka Sawada // Copyright : 2023-10-13 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -1848,7 +1848,7 @@ int sqrt32(int num)
/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
int cpu_num = 1; // CPU/Core 個数が制限されてる場合は、上位に本来の数を置く int cpu_num = 1; // CPU/Core 個数が制限されてる場合は、上位に本来の数を置く
// /arch:SSE2, +1=SSSE3, +2=SSE4.1, +4=SSE4.2, +8=CLMUL, +16=AVX2, +128=JIT(SSE2), +256=Old // /arch:SSE2, +1=SSSE3, +2=SSE4.1, +4=SSE4.2, +8=CLMUL, +16=AVX2, +128=JIT(SSE2), +256=ALTMAPなし
// 上位 16-bit = L2 cache サイズから計算した制限サイズ // 上位 16-bit = L2 cache サイズから計算した制限サイズ
unsigned int cpu_flag = 0; unsigned int cpu_flag = 0;
unsigned int cpu_cache = 0; // 上位 16-bit = L3 cache の制限サイズ, 下位 16-bit = 同時処理数 unsigned int cpu_cache = 0; // 上位 16-bit = L3 cache の制限サイズ, 下位 16-bit = 同時処理数

View File

@@ -1,5 +1,5 @@
// create.c // create.c
// Copyright : 2023-09-23 Yutaka Sawada // Copyright : 2023-10-22 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -320,7 +320,7 @@ error_end:
return off; return off;
} }
#define MAX_MULTI_READ 4 // SSDで同時に読み込む最大ファイル数 #define MAX_MULTI_READ 6 // SSDで同時に読み込む最大ファイル数
// SSD 上で複数ファイルのハッシュ値を同時に求めるバージョン // SSD 上で複数ファイルのハッシュ値を同時に求めるバージョン
int set_common_packet_multi( int set_common_packet_multi(
@@ -348,11 +348,9 @@ unsigned int time_start = GetTickCount();
memset(th, 0, sizeof(FILE_HASH_TH) * MAX_MULTI_READ); memset(th, 0, sizeof(FILE_HASH_TH) * MAX_MULTI_READ);
// Core数に応じてスレッド数を増やす // Core数に応じてスレッド数を増やす
if ((memory_use & 32) != 0){ // NVMe SSD if ((memory_use & 32) != 0){ // NVMe SSD
if (cpu_num >= 8){ // 8 ~ 16 Cores multi_read = (cpu_num + 2) / 3 + 1; // 3=2, 4~6=3, 7~9=4, 10~12=5, 13~=6
multi_read = 4; if (multi_read > MAX_MULTI_READ)
} else { // 3 Cores + Hyper-threading, or 4 ~ 7 Cores multi_read = MAX_MULTI_READ;
multi_read = 3;
}
} else { // SATA SSD } else { // SATA SSD
multi_read = 2; multi_read = 2;
} }
@@ -1282,6 +1280,7 @@ int create_recovery_file_1pass(
int footer_size, // 末尾パケットのバッファー・サイズ int footer_size, // 末尾パケットのバッファー・サイズ
HANDLE *rcv_hFile, // 各リカバリ・ファイルのハンドル HANDLE *rcv_hFile, // 各リカバリ・ファイルのハンドル
unsigned char *p_buf, // 計算済みのパリティ・ブロック unsigned char *p_buf, // 計算済みのパリティ・ブロック
unsigned char *g_buf, // GPU用 (GPUを使わない場合は NULLにすること)
unsigned int unit_size) unsigned int unit_size)
{ {
unsigned char *packet_header, hash[HASH_SIZE]; unsigned char *packet_header, hash[HASH_SIZE];
@@ -1438,6 +1437,10 @@ int create_recovery_file_1pass(
// Recovery Slice packet は後から書き込む // Recovery Slice packet は後から書き込む
for (j = block_start; j < block_start + block_count; j++){ for (j = block_start; j < block_start + block_count; j++){
if (g_buf != NULL){ // GPUを使った場合
// CPUスレッドと GPUスレッドの計算結果を合わせる
galois_align_xor(g_buf + (size_t)unit_size * j, p_buf, unit_size);
}
// パリティ・ブロックのチェックサムを検証する // パリティ・ブロックのチェックサムを検証する
checksum16_return(p_buf, hash, unit_size - HASH_SIZE); checksum16_return(p_buf, hash, unit_size - HASH_SIZE);
if (memcmp(p_buf + unit_size - HASH_SIZE, hash, HASH_SIZE) != 0){ if (memcmp(p_buf + unit_size - HASH_SIZE, hash, HASH_SIZE) != 0){

View File

@@ -82,6 +82,7 @@ int create_recovery_file_1pass(
int footer_size, // 末尾パケットのバッファー・サイズ int footer_size, // 末尾パケットのバッファー・サイズ
HANDLE *rcv_hFile, // 各リカバリ・ファイルのハンドル HANDLE *rcv_hFile, // 各リカバリ・ファイルのハンドル
unsigned char *p_buf, // 計算済みのパリティ・ブロック unsigned char *p_buf, // 計算済みのパリティ・ブロック
unsigned char *g_buf, // GPU用 (GPUを使わない場合は NULLにすること)
unsigned int unit_size); unsigned int unit_size);
// 作成中のリカバリ・ファイルを削除する // 作成中のリカバリ・ファイルを削除する

View File

@@ -71,7 +71,6 @@ extern unsigned int cpu_flag; // declared in common2.h
// CPU によって使う関数を変更する際の仮宣言 // CPU によって使う関数を変更する際の仮宣言
//#define NO_SIMD // SIMD を使わない場合 //#define NO_SIMD // SIMD を使わない場合
//#define NO_ALTMAP // SSSE3 や JIT(SSE2) の並び替えを使わない場合 (CLMULや32バイト単位は有効)
int sse_unit; int sse_unit;
@@ -134,8 +133,11 @@ int galois_create_table(void)
checksum16_altmap = checksum16; checksum16_altmap = checksum16;
checksum16_return = checksum16; checksum16_return = checksum16;
#ifndef NO_SIMD #ifndef NO_SIMD
#ifndef NO_ALTMAP if (cpu_flag & 256){ // AVX2, SSSE3, JIT(SSE2) の並び替えを使わない場合
if (cpu_flag & 16){ // AVX2 対応なら // 将来的には AVX-512 などの命令に対応してもいい
//printf("\nWithout ALTMAP\n");
//sse_unit = 32;
} else if (cpu_flag & 16){ // AVX2 対応なら
//printf("\nUse AVX2 & ALTMAP\n"); //printf("\nUse AVX2 & ALTMAP\n");
sse_unit = 32; // 32, 64, 128 のどれでもいい sse_unit = 32; // 32, 64, 128 のどれでもいい
galois_align_multiply = galois_align32avx_multiply; galois_align_multiply = galois_align32avx_multiply;
@@ -145,7 +147,6 @@ int galois_create_table(void)
checksum16_altmap = checksum16_altmap32; checksum16_altmap = checksum16_altmap32;
checksum16_return = checksum16_return32; checksum16_return = checksum16_return32;
} else if (cpu_flag & 1){ // SSSE3 対応なら } else if (cpu_flag & 1){ // SSSE3 対応なら
if ((cpu_flag & 256) == 0){ // SSSE3 & ALTMAP を使う
//printf("\nUse SSSE3 & ALTMAP\n"); //printf("\nUse SSSE3 & ALTMAP\n");
sse_unit = 32; // 32, 64, 128 のどれでもいい sse_unit = 32; // 32, 64, 128 のどれでもいい
galois_align_multiply = galois_align32_multiply; galois_align_multiply = galois_align32_multiply;
@@ -154,7 +155,6 @@ int galois_create_table(void)
galois_altmap_return = galois_altmap32_return; galois_altmap_return = galois_altmap32_return;
checksum16_altmap = checksum16_altmap32; checksum16_altmap = checksum16_altmap32;
checksum16_return = checksum16_return32; checksum16_return = checksum16_return32;
}
} else { // SSSE3 が利用できない場合 } else { // SSSE3 が利用できない場合
if ((cpu_flag & 128) && (jit_alloc() == 0)){ // JIT(SSE2) を使う if ((cpu_flag & 128) && (jit_alloc() == 0)){ // JIT(SSE2) を使う
//printf("\nUse JIT(SSE2) & ALTMAP\n"); //printf("\nUse JIT(SSE2) & ALTMAP\n");
@@ -167,7 +167,6 @@ int galois_create_table(void)
checksum16_return = checksum16_return256; checksum16_return = checksum16_return256;
} }
} }
#endif
#endif #endif
return 0; return 0;
@@ -2792,7 +2791,7 @@ void galois_align_xor(
#endif #endif
} }
// 16バイト境界のバッファー専用の掛け算 // 16バイト境界のバッファー専用の掛け算 (ALTMAP しない)
void galois_align16_multiply( void galois_align16_multiply(
unsigned char *r1, // Region to multiply (must be aligned by 16) unsigned char *r1, // Region to multiply (must be aligned by 16)
unsigned char *r2, // Products go here unsigned char *r2, // Products go here
@@ -2826,6 +2825,16 @@ void galois_align16_multiply(
// 掛け算用のテーブルを常に作成する (32バイトだと少し遅くなる) // 掛け算用のテーブルを常に作成する (32バイトだと少し遅くなる)
#ifndef NO_SIMD #ifndef NO_SIMD
/*
// sse_unit が 32の倍数な時だけ
} else if (cpu_flag & 16){ // AVX2 対応なら
__declspec( align(32) ) unsigned char small_table[128];
create_eight_table_avx2(small_table, factor);
gf16_avx2_block32u(r1, r2, len, small_table);
*/
} else if (cpu_flag & 1){ // SSSE3 対応なら } else if (cpu_flag & 1){ // SSSE3 対応なら
__declspec( align(16) ) unsigned char small_table[128]; __declspec( align(16) ) unsigned char small_table[128];
@@ -2869,7 +2878,6 @@ void galois_align16_multiply(
len -= 8; len -= 8;
} }
#endif #endif
} }
} }

View File

@@ -1,5 +1,5 @@
// lib_opencl.c // lib_opencl.c
// Copyright : 2023-09-23 Yutaka Sawada // Copyright : 2023-10-22 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _WIN32_WINNT #ifndef _WIN32_WINNT
@@ -102,7 +102,6 @@ API_clEnqueueNDRangeKernel gfn_clEnqueueNDRangeKernel;
入力 入力
OpenCL_method : どのデバイスを選ぶか OpenCL_method : どのデバイスを選ぶか
unit_size : ブロックの単位サイズ unit_size : ブロックの単位サイズ
chunk_size: 分割された断片サイズ
src_max : ソース・ブロック個数 src_max : ソース・ブロック個数
出力 出力
@@ -112,7 +111,7 @@ OpenCL_method : 動作フラグいろいろ
*/ */
// 0=成功, 1エラー番号 // 0=成功, 1エラー番号
int init_OpenCL(int unit_size, int chunk_size, int *src_max) int init_OpenCL(int unit_size, int *src_max)
{ {
char buf[2048], *p_source; char buf[2048], *p_source;
int err = 0, i, j; int err = 0, i, j;
@@ -283,23 +282,12 @@ int init_OpenCL(int unit_size, int chunk_size, int *src_max)
// printf("Shared Virtual Memory = 0x%I64X\n", param_value8); // printf("Shared Virtual Memory = 0x%I64X\n", param_value8);
#endif #endif
ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint), &param_value, NULL);
if (ret != CL_SUCCESS)
continue;
ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &param_value8, NULL); ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &param_value8, NULL);
if (ret != CL_SUCCESS) if (ret != CL_SUCCESS)
continue; continue;
#ifdef DEBUG_OUTPUT #ifdef DEBUG_OUTPUT
printf("ADDRESS_BITS = %d\n", param_value);
printf("MAX_MEM_ALLOC_SIZE = %I64d MB\n", param_value8 >> 20); printf("MAX_MEM_ALLOC_SIZE = %I64d MB\n", param_value8 >> 20);
#endif #endif
if (param_value == 32){ // CL_DEVICE_ADDRESS_BITS によって確保するメモリー領域の上限を変える
if (param_value8 > 0x30000000) // 768MB までにする
param_value8 = 0x30000000;
} else { // 64-bit OS でも 2GB までにする
if (param_value8 > 0x80000000)
param_value8 = 0x80000000;
}
ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &num_groups, NULL); ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &num_groups, NULL);
if (ret != CL_SUCCESS) if (ret != CL_SUCCESS)
@@ -325,13 +313,13 @@ int init_OpenCL(int unit_size, int chunk_size, int *src_max)
OpenCL_group_num = num_groups; // ワークグループ数は COMPUTE_UNITS 数にする OpenCL_group_num = num_groups; // ワークグループ数は COMPUTE_UNITS 数にする
alloc_max = (size_t)param_value8; alloc_max = (size_t)param_value8;
// AMD Radeon ではメモリー領域が全体の 1/4 とは限らない // AMD や Intel の GPU ではメモリー領域が全体の 1/4 とは限らない
ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &param_value8, NULL); ret = fn_clGetDeviceInfo(device_id[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &param_value8, NULL);
if (ret == CL_SUCCESS){ if (ret == CL_SUCCESS){
#ifdef DEBUG_OUTPUT #ifdef DEBUG_OUTPUT
printf("GLOBAL_MEM_SIZE = %I64d MB\n", param_value8 >> 20); printf("GLOBAL_MEM_SIZE = %I64d MB\n", param_value8 >> 20);
#endif #endif
// 領域一個あたりのサイズは全体の 1/4 までにする // 領域一個あたりのサイズは全体の 1/4 までにするVRAMを使いすぎると不安定になる
param_value8 /= 4; param_value8 /= 4;
if ((cl_ulong)alloc_max > param_value8) if ((cl_ulong)alloc_max > param_value8)
alloc_max = (size_t)param_value8; alloc_max = (size_t)param_value8;
@@ -366,7 +354,7 @@ int init_OpenCL(int unit_size, int chunk_size, int *src_max)
return (ret << 8) | 12; return (ret << 8) | 12;
// 計算方式を選択する // 計算方式を選択する
if ((((cpu_flag & 0x101) == 1) || ((cpu_flag & 16) != 0)) && (sse_unit == 32)){ if ((((cpu_flag & 0x101) == 1) || ((cpu_flag & 0x110) == 0x10)) && (sse_unit == 32)){
OpenCL_method = 2; // SSSE3 & ALTMAP または AVX2 ならデータの並び替え対応版を使う OpenCL_method = 2; // SSSE3 & ALTMAP または AVX2 ならデータの並び替え対応版を使う
} else if (((cpu_flag & 128) != 0) && (sse_unit == 256)){ } else if (((cpu_flag & 128) != 0) && (sse_unit == 256)){
OpenCL_method = 4; // JIT(SSE2) は bit ごとに上位から 16バイトずつ並ぶ OpenCL_method = 4; // JIT(SSE2) は bit ごとに上位から 16バイトずつ並ぶ
@@ -380,10 +368,10 @@ int init_OpenCL(int unit_size, int chunk_size, int *src_max)
// work group 数が必要以上に多い場合は減らす // work group 数が必要以上に多い場合は減らす
if (OpenCL_method == 2){ if (OpenCL_method == 2){
// work item 一個が 8バイトずつ計算する、256個なら work group ごとに 2KB 担当する // work item 一個が 8バイトずつ計算する、256個なら work group ごとに 2KB 担当する
data_size = chunk_size / 2048; data_size = unit_size / 2048;
} else { } else {
// work item 一個が 4バイトずつ計算する、256個なら work group ごとに 1KB 担当する // work item 一個が 4バイトずつ計算する、256個なら work group ごとに 1KB 担当する
data_size = chunk_size / 1024; data_size = unit_size / 1024;
} }
if (OpenCL_group_num > data_size){ if (OpenCL_group_num > data_size){
OpenCL_group_num = data_size; OpenCL_group_num = data_size;
@@ -401,9 +389,9 @@ int init_OpenCL(int unit_size, int chunk_size, int *src_max)
printf("src buf : %zd KB (%d blocks), possible\n", data_size >> 10, count); printf("src buf : %zd KB (%d blocks), possible\n", data_size >> 10, count);
#endif #endif
// 出力先はchunk 1個分だけあればいい // 出力先は1ブロック分だけあればいい
// CL_MEM_ALLOC_HOST_PTRを使えばpinned memoryになるらしい // CL_MEM_ALLOC_HOST_PTRを使えばpinned memoryになるらしい
data_size = (chunk_size + 63) & ~63; // cache line sizes (64 bytes) の倍数にする data_size = unit_size;
OpenCL_dst = gfn_clCreateBuffer(OpenCL_context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, data_size, NULL, &ret); OpenCL_dst = gfn_clCreateBuffer(OpenCL_context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, data_size, NULL, &ret);
if (ret != CL_SUCCESS) if (ret != CL_SUCCESS)
return (ret << 8) | 13; return (ret << 8) | 13;
@@ -700,12 +688,11 @@ int gpu_copy_blocks(
} }
// ソース・ブロックを掛け算する // ソース・ブロックを掛け算する
int gpu_multiply_chunks( int gpu_multiply_blocks(
int src_num, // Number of multiplying source blocks int src_num, // Number of multiplying source blocks
unsigned short *mat, // Matrix of numbers to multiply by unsigned short *mat, // Matrix of numbers to multiply by
unsigned char *buf, // Products go here unsigned char *buf, // Products go here
int offset, // Offset in each block int len) // Byte length
int length) // Byte length
{ {
unsigned __int64 *vram, *src, *dst; unsigned __int64 *vram, *src, *dst;
size_t global_size, local_size; size_t global_size, local_size;
@@ -720,14 +707,6 @@ int gpu_multiply_chunks(
ret = gfn_clSetKernelArg(OpenCL_kernel, 3, sizeof(int), &src_num); ret = gfn_clSetKernelArg(OpenCL_kernel, 3, sizeof(int), &src_num);
if (ret != CL_SUCCESS) if (ret != CL_SUCCESS)
return (ret << 8) | 103; return (ret << 8) | 103;
offset /= 4; // 4バイト整数単位にする
ret = gfn_clSetKernelArg(OpenCL_kernel, 4, sizeof(int), &offset);
if (ret != CL_SUCCESS)
return (ret << 8) | 104;
length /= 4; // 4バイト整数単位にする
ret = gfn_clSetKernelArg(OpenCL_kernel, 5, sizeof(int), &length);
if (ret != CL_SUCCESS)
return (ret << 8) | 105;
// カーネル並列実行 // カーネル並列実行
local_size = 256; // テーブルやキャッシュのため、work item 数は 256に固定する local_size = 256; // テーブルやキャッシュのため、work item 数は 256に固定する
@@ -738,18 +717,18 @@ int gpu_multiply_chunks(
return (ret << 8) | 11; return (ret << 8) | 11;
// 出力内容をホスト側に反映させる // 出力内容をホスト側に反映させる
vram = gfn_clEnqueueMapBuffer(OpenCL_command, OpenCL_dst, CL_TRUE, CL_MAP_READ, 0, length * 4, 0, NULL, NULL, &ret); vram = gfn_clEnqueueMapBuffer(OpenCL_command, OpenCL_dst, CL_TRUE, CL_MAP_READ, 0, len, 0, NULL, NULL, &ret);
if (ret != CL_SUCCESS) if (ret != CL_SUCCESS)
return (ret << 8) | 12; return (ret << 8) | 12;
// 8バイトごとに XOR する (SSE2 で XOR しても速くならず) // 8バイトごとに XOR する (SSE2 で XOR しても速くならず)
src = vram; src = vram;
dst = (unsigned __int64 *)buf; dst = (unsigned __int64 *)buf;
while (length > 0){ while (len > 0){
*dst ^= *src; *dst ^= *src;
dst++; dst++;
src++; src++;
length -= 2; len -= 8;
} }
// ホスト側でデータを変更しなくても、clEnqueueMapBufferと対で呼び出さないといけない // ホスト側でデータを変更しなくても、clEnqueueMapBufferと対で呼び出さないといけない

View File

@@ -10,7 +10,7 @@ extern "C" {
extern int OpenCL_method; extern int OpenCL_method;
int init_OpenCL(int unit_size, int chunk_size, int *src_max); int init_OpenCL(int unit_size, int *src_max);
int free_OpenCL(void); int free_OpenCL(void);
void info_OpenCL(char *buf, int buf_size); void info_OpenCL(char *buf, int buf_size);
@@ -19,12 +19,11 @@ int gpu_copy_blocks(
int unit_size, int unit_size,
int src_num); int src_num);
int gpu_multiply_chunks( int gpu_multiply_blocks(
int src_num, // Number of multiplying source blocks int src_num, // Number of multiplying source blocks
unsigned short *mat, // Matrix of numbers to multiply by unsigned short *mat, // Matrix of numbers to multiply by
unsigned char *buf, // Products go here unsigned char *buf, // Products go here
int offset, // Offset in each block int len); // Byte length
int length); // Byte length
int gpu_finish(void); int gpu_finish(void);

View File

@@ -1,5 +1,5 @@
// list.c // list.c
// Copyright : 2022-10-14 Yutaka Sawada // Copyright : 2023-10-15 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -348,7 +348,7 @@ if (time_start > 0){
// SSD 上で複数ファイルを同時に検査する // SSD 上で複数ファイルを同時に検査する
// MAX_MULTI_READ の2倍ぐらいにする? // MAX_MULTI_READ の2倍ぐらいにする?
#define MAX_READ_NUM 10 #define MAX_READ_NUM 12
int check_file_complete_multi( int check_file_complete_multi(
char *ascii_buf, char *ascii_buf,
@@ -370,11 +370,9 @@ unsigned int time_start = GetTickCount();
memset(hSub, 0, sizeof(HANDLE) * MAX_READ_NUM); memset(hSub, 0, sizeof(HANDLE) * MAX_READ_NUM);
// Core数に応じてスレッド数を増やす // Core数に応じてスレッド数を増やす
if ((memory_use & 32) != 0){ // NVMe SSD if ((memory_use & 32) != 0){ // NVMe SSD
if (cpu_num >= 8){ // 8 ~ 16 Cores multi_read = (cpu_num + 2) / 3 + 1; // 3=2, 4~6=3, 7~9=4, 10~12=5, 13~=6
multi_read = 4; if (multi_read > MAX_READ_NUM / 2)
} else { // 3 Cores + Hyper-threading, or 4 ~ 7 Cores multi_read = MAX_READ_NUM / 2;
multi_read = 3;
}
} else { // SATA SSD } else { // SATA SSD
multi_read = 2; multi_read = 2;
} }

View File

@@ -1,5 +1,5 @@
// md5_crc.c // md5_crc.c
// Copyright : 2023-08-28 Yutaka Sawada // Copyright : 2023-10-29 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -204,6 +204,8 @@ static unsigned int time_start, time1_start;
static unsigned int time_total = 0, time2_total = 0, time3_total = 0; static unsigned int time_total = 0, time2_total = 0, time3_total = 0;
#endif #endif
#define MAX_BUF_SIZE 2097152 // ヒープ領域を使う場合の最大サイズ
// ファイルのハッシュ値と各スライスのチェックサムを同時に計算する // ファイルのハッシュ値と各スライスのチェックサムを同時に計算する
int file_hash_crc( int file_hash_crc(
wchar_t *file_name, // ハッシュ値を求めるファイル wchar_t *file_name, // ハッシュ値を求めるファイル
@@ -671,7 +673,7 @@ time1_start = GetTickCount();
} }
// バッファー・サイズが大きいのでヒープ領域を使う // バッファー・サイズが大きいのでヒープ領域を使う
for (io_size = IO_SIZE; io_size < 1048576; io_size += IO_SIZE){ // 1 MB までにする for (io_size = IO_SIZE; io_size <= MAX_BUF_SIZE; io_size += IO_SIZE){ // IO_SIZE の倍数にする
if ((io_size + IO_SIZE > (cpu_cache & 0xFFFE0000)) || ((__int64)(io_size + IO_SIZE) * 4 > file_left)) if ((io_size + IO_SIZE > (cpu_cache & 0xFFFE0000)) || ((__int64)(io_size + IO_SIZE) * 4 > file_left))
break; break;
} }
@@ -866,7 +868,7 @@ DWORD WINAPI file_hash_crc2(LPVOID lpParameter)
// バッファー・サイズが大きいのでヒープ領域を使う // バッファー・サイズが大きいのでヒープ領域を使う
prog_tick = 1; prog_tick = 1;
for (io_size = IO_SIZE; io_size < 1048576; io_size += IO_SIZE){ // IO_SIZE の倍数で 1 MB までにする for (io_size = IO_SIZE; io_size <= MAX_BUF_SIZE; io_size += IO_SIZE){ // IO_SIZE の倍数にする
if ((io_size + IO_SIZE > (cpu_cache & 0xFFFE0000)) || ((__int64)(io_size + IO_SIZE) * 4 > file_left)) if ((io_size + IO_SIZE > (cpu_cache & 0xFFFE0000)) || ((__int64)(io_size + IO_SIZE) * 4 > file_left))
break; break;
prog_tick++; prog_tick++;
@@ -1303,7 +1305,7 @@ DWORD WINAPI file_hash_background(LPVOID lpParameter)
find_next = files[num].b_off; // 先頭ブロックの番号 find_next = files[num].b_off; // 先頭ブロックの番号
// バッファー・サイズが大きいのでヒープ領域を使う // バッファー・サイズが大きいのでヒープ領域を使う
for (io_size = IO_SIZE; io_size < 1048576; io_size += IO_SIZE){ // IO_SIZE の倍数で 1 MB までにする for (io_size = IO_SIZE; io_size <= MAX_BUF_SIZE; io_size += IO_SIZE){ // IO_SIZE の倍数にする
if ((io_size + IO_SIZE > (cpu_cache & 0xFFFE0000)) || ((__int64)(io_size + IO_SIZE) * 4 > file_size)) if ((io_size + IO_SIZE > (cpu_cache & 0xFFFE0000)) || ((__int64)(io_size + IO_SIZE) * 4 > file_size))
break; break;
} }

View File

@@ -1,5 +1,5 @@
// par2.c // par2.c
// Copyright : 2023-09-21 Yutaka Sawada // Copyright : 2023-10-15 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -181,7 +181,7 @@ int par2_create(
} }
} else { } else {
// 共通パケットを作成する // 共通パケットを作成する
if ((memory_use & 16) && (cpu_num >= 4) && (entity_num >= 2)){ // SSDなら複数ファイルを同時に処理する if ((memory_use & 16) && (cpu_num >= 3) && (entity_num >= 2)){ // SSDなら複数ファイルを同時に処理する
common_size = set_common_packet_multi(common_buf, &packet_num, (switch_p & 2) >> 1, files); common_size = set_common_packet_multi(common_buf, &packet_num, (switch_p & 2) >> 1, files);
} else { } else {
common_size = set_common_packet(common_buf, &packet_num, (switch_p & 2) >> 1, files); common_size = set_common_packet(common_buf, &packet_num, (switch_p & 2) >> 1, files);
@@ -529,7 +529,7 @@ int par2_verify(
// ソース・ファイルが完全かどうかを調べる // ソース・ファイルが完全かどうかを調べる
// ファイルの状態は 完全、消失、追加、破損(完全なブロックの数) の4種類 // ファイルの状態は 完全、消失、追加、破損(完全なブロックの数) の4種類
if ((memory_use & 16) && (cpu_num >= 4) && (entity_num >= 2)){ // SSDなら複数ファイルを同時に処理する if ((memory_use & 16) && (cpu_num >= 3) && (entity_num >= 2)){ // SSDなら複数ファイルを同時に処理する
err = check_file_complete_multi(ascii_buf, uni_buf, files, s_blk); err = check_file_complete_multi(ascii_buf, uni_buf, files, s_blk);
} else { } else {
err = check_file_complete(ascii_buf, uni_buf, files, s_blk); err = check_file_complete(ascii_buf, uni_buf, files, s_blk);
@@ -741,7 +741,7 @@ int par2_repair(
// ソース・ファイルが完全かどうかを一覧表示する // ソース・ファイルが完全かどうかを一覧表示する
// ファイルの状態は 完全、消失、追加、破損(完全なブロックの数) の4種類 // ファイルの状態は 完全、消失、追加、破損(完全なブロックの数) の4種類
if ((memory_use & 16) && (cpu_num >= 4) && (entity_num >= 2)){ // SSDなら複数ファイルを同時に処理する if ((memory_use & 16) && (cpu_num >= 3) && (entity_num >= 2)){ // SSDなら複数ファイルを同時に処理する
err = check_file_complete_multi(ascii_buf, uni_buf, files, s_blk); err = check_file_complete_multi(ascii_buf, uni_buf, files, s_blk);
} else { } else {
err = check_file_complete(ascii_buf, uni_buf, files, s_blk); err = check_file_complete(ascii_buf, uni_buf, files, s_blk);

View File

@@ -1,5 +1,5 @@
// par2_cmd.c // par2_cmd.c
// Copyright : 2023-09-28 Yutaka Sawada // Copyright : 2023-10-15 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -87,35 +87,37 @@ static void print_environment(void)
printf("CPU thread\t: %d / %d\n", cpu_num & 0xFFFF, cpu_num >> 24); printf("CPU thread\t: %d / %d\n", cpu_num & 0xFFFF, cpu_num >> 24);
cpu_num &= 0xFFFF; // 利用するコア数だけにしておく cpu_num &= 0xFFFF; // 利用するコア数だけにしておく
printf("CPU cache limit : %d KB, %d KB\n", (cpu_flag & 0xFFFF0000) >> 10, (cpu_cache & 0xFFFE0000) >> 10); printf("CPU cache limit : %d KB, %d KB\n", (cpu_flag & 0xFFFF0000) >> 10, (cpu_cache & 0xFFFE0000) >> 10);
#ifndef _WIN64 // 32-bit 版は MMX, SSE2, SSSE3 のどれかを表示する #ifndef _WIN64 // 32-bit 版は MMX, SSE2, SSSE3, AVX2 のどれかを表示する
printf("CPU extra\t:"); printf("CPU extra\t:");
if (cpu_flag & 1){ if (((cpu_flag & 16) != 0) && ((cpu_flag & 256) == 0)){
printf(" AVX2");
} else if (cpu_flag & 1){
if (cpu_flag & 256){ if (cpu_flag & 256){
printf(" SSSE3(old)"); printf(" SSSE3(slow)");
} else { } else {
printf(" SSSE3"); printf(" SSSE3");
} }
} else if (cpu_flag & 128){ } else if (((cpu_flag & 128) != 0) && ((cpu_flag & 256) == 0)){
printf(" SSE2"); printf(" SSE2");
} else { } else {
printf(" MMX"); printf(" MMX");
} }
#else // 64-bit 版は SSE2, SSSE3 を表示する #else // 64-bit 版は SSE2, SSSE3, AVX2 を表示する
printf("CPU extra\t: x64"); printf("CPU extra\t: x64");
if (cpu_flag & 1){ if (((cpu_flag & 16) != 0) && ((cpu_flag & 256) == 0)){
printf(" AVX2");
} else if (cpu_flag & 1){
if (cpu_flag & 256){ if (cpu_flag & 256){
printf(" SSSE3(old)"); printf(" SSSE3(slow)");
} else { } else {
printf(" SSSE3"); printf(" SSSE3");
} }
} else if (cpu_flag & 128){ } else if (((cpu_flag & 128) != 0) && ((cpu_flag & 256) == 0)){
printf(" SSE2"); printf(" SSE2");
} }
#endif #endif
if (cpu_flag & 8) if (cpu_flag & 8)
printf(" CLMUL"); printf(" CLMUL");
if (cpu_flag & 16)
printf(" AVX2");
printf("\nMemory usage\t: "); printf("\nMemory usage\t: ");
if (memory_use & 7){ if (memory_use & 7){
printf("%d/8", memory_use & 7); printf("%d/8", memory_use & 7);
@@ -1486,8 +1488,8 @@ ri= switch_set & 0x00040000
} else if (k & 512){ } else if (k & 512){
OpenCL_method = -1; // Slower GPU OpenCL_method = -1; // Slower GPU
} }
if (k & 1024) // CLMUL を使わない、SSSE3 の古いエンコーダーを使う if (k & 1024) // CLMUL と ALTMAP を使わない
cpu_flag = (cpu_flag & 0xFFFFFFF7) | 0x100; cpu_flag = (cpu_flag & 0xFFFFFFF7) | 256;
if (k & 2048) // JIT(SSE2) を使わない if (k & 2048) // JIT(SSE2) を使わない
cpu_flag &= 0xFFFFFF7F; cpu_flag &= 0xFFFFFF7F;
if (k & 4096) // SSSE3 を使わない if (k & 4096) // SSSE3 を使わない
@@ -1506,10 +1508,10 @@ ri= switch_set & 0x00040000
} else if (k == 254){ // 物理コア数より減らす } else if (k == 254){ // 物理コア数より減らす
k = ((cpu_num & 0x00FF0000) >> 16) - 1; k = ((cpu_num & 0x00FF0000) >> 16) - 1;
} else if (k == 255){ // 物理コア数より増やす } else if (k == 255){ // 物理コア数より増やす
k = ((cpu_num & 0x00FF0000) >> 16) + 1; k = cpu_num >> 16;
//k = cpu_num >> 16; k = ((k & 0xFF) + (k >> 8)) / 2; // 物理コア数と論理コア数の中間にする?
//k = ((k & 0xFF) + (k >> 8)) / 2; // 物理コア数と論理コア数の中間にする?
// タスクマネージャーにおける CPU使用率は 100%になるけど、速くはならない・・・ // タスクマネージャーにおける CPU使用率は 100%になるけど、速くはならない・・・
// k = (k & 0xFF) + ((k >> 8) - (k & 0xFF)) / 4; // 物理コア数の 5/4 にする?
} }
if (k > MAX_CPU){ if (k > MAX_CPU){
k = MAX_CPU; k = MAX_CPU;

View File

@@ -1,5 +1,5 @@
// reedsolomon.c // reedsolomon.c
// Copyright : 2023-09-28 Yutaka Sawada // Copyright : 2023-10-26 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -204,6 +204,48 @@ int read_block_num(
return buf_num; return buf_num;
} }
// 1st encode, decode を何スレッドで実行するか決める
int calc_thread_num1(int max_num)
{
int i, num;
// 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする
num = 0;
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
num++;
i *= 2;
}
if (num > max_num)
num = max_num;
return num;
}
// 1st & 2nd encode, decode を何スレッドで実行するか決める
int calc_thread_num2(int max_num, int *cpu_num2)
{
int i, num1, num2;
// 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする
num1 = 0;
i = 2;
while (i <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
num1++;
i *= 2;
}
if (num1 > max_num)
num1 = max_num;
// CPU と GPU で必ず2スレッド使う
num2 = cpu_num;
if (num2 < 2)
num2 = 2;
*cpu_num2 = num2;
return num1;
}
/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
// 戸川 隼人 の「演習と応用FORTRAN77」の逆行列の計算方法を参考にして // 戸川 隼人 の「演習と応用FORTRAN77」の逆行列の計算方法を参考にして
// Gaussian Elimination を少し修正して行列の数を一つにしてみた // Gaussian Elimination を少し修正して行列の数を一つにしてみた
@@ -960,8 +1002,7 @@ time_matrix = GetTickCount() - time_matrix;
// ブロック数が多いなら、ブロックごとにスレッドを割り当てる (GPU を使う) // ブロック数が多いなら、ブロックごとにスレッドを割り当てる (GPU を使う)
if (memory_use & 16){ if (memory_use & 16){
err = -4; // SSD なら Read all 方式でブロックが断片化しても速い err = -4; // SSD なら Read all 方式でブロックが断片化しても速い
} else } else if (read_block_num(block_lost * 2, 0, MEM_UNIT) != 0){
if (read_block_num(block_lost, 0, MEM_UNIT) != 0){
err = -5; // HDD でメモリーが足りてるなら Read some 方式を使う err = -5; // HDD でメモリーが足りてるなら Read some 方式を使う
} else { } else {
err = -4; // メモリー不足なら Read all 方式でブロックを断片化させる err = -4; // メモリー不足なら Read all 方式でブロックを断片化させる
@@ -970,8 +1011,7 @@ time_matrix = GetTickCount() - time_matrix;
// ソース・ブロックを全て断片的に読み込むか、いくつかを丸ごと読み込むかを決める // ソース・ブロックを全て断片的に読み込むか、いくつかを丸ごと読み込むかを決める
if (memory_use & 16){ if (memory_use & 16){
err = -2; // SSD なら Read all 方式でブロックが断片化しても速い err = -2; // SSD なら Read all 方式でブロックが断片化しても速い
} else } else if (read_block_num(block_lost, 0, sse_unit) != 0){
if (read_block_num(block_lost, 0, sse_unit) != 0){
err = -3; // HDD でメモリーが足りてるなら Read some 方式を使う err = -3; // HDD でメモリーが足りてるなら Read some 方式を使う
} else { } else {
err = -2; // メモリー不足なら Read all 方式でブロックを断片化させる err = -2; // メモリー不足なら Read all 方式でブロックを断片化させる

View File

@@ -17,6 +17,10 @@ extern "C" {
#define READ_MIN_RATE 1 // 保持するブロック数の 1/2 = 50% #define READ_MIN_RATE 1 // 保持するブロック数の 1/2 = 50%
#define READ_MIN_NUM 16 #define READ_MIN_NUM 16
// CPU cache 最適化のため、同時に処理するブロック数を制限する
#define CACHE_MIN_NUM 8
#define CACHE_MAX_NUM 128
/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
// Cache Blocking を試みる // Cache Blocking を試みる
@@ -35,6 +39,12 @@ int read_block_num(
size_t trial_alloc, // 確保できるか確認するのか size_t trial_alloc, // 確保できるか確認するのか
int alloc_unit); // メモリー単位の境界 (sse_unit か MEM_UNIT) int alloc_unit); // メモリー単位の境界 (sse_unit か MEM_UNIT)
// 1st encode, decode を何スレッドで実行するか決める
int calc_thread_num1(int max_num);
// 1st & 2nd encode, decode を何スレッドで実行するか決める
int calc_thread_num2(int max_num, int *cpu_num2);
/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
// リード・ソロモン符号を使ってエンコードする // リード・ソロモン符号を使ってエンコードする

View File

@@ -1,7 +1,7 @@
1 RT_STRING ".\\source.cl" 1 RT_STRING ".\\source.cl"
1 VERSIONINFO 1 VERSIONINFO
FILEVERSION 1,3,3,0 FILEVERSION 1,3,3,1
PRODUCTVERSION 1,3,3,0 PRODUCTVERSION 1,3,3,0
FILEOS 0x40004 FILEOS 0x40004
FILETYPE 0x1 FILETYPE 0x1
@@ -13,7 +13,7 @@ BLOCK "StringFileInfo"
VALUE "FileDescription", "PAR2 client" VALUE "FileDescription", "PAR2 client"
VALUE "LegalCopyright", "Copyright (C) 2023 Yutaka Sawada" VALUE "LegalCopyright", "Copyright (C) 2023 Yutaka Sawada"
VALUE "ProductName", "par2j" VALUE "ProductName", "par2j"
VALUE "FileVersion", "1.3.3.0" VALUE "FileVersion", "1.3.3.1"
VALUE "ProductVersion", "1.3.3.0" VALUE "ProductVersion", "1.3.3.0"
} }
} }

View File

@@ -1,5 +1,5 @@
// rs_decode.c // rs_decode.c
// Copyright : 2023-09-21 Yutaka Sawada // Copyright : 2023-10-29 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -284,11 +284,11 @@ printf(" 2nd decode %d.%03d sec, %d loop, %d MB/s\n", time_encode2b / 1000, time
// GPU 対応のサブ・スレッド (最後のスレッドなので、1st decode では呼ばれない) // GPU 対応のサブ・スレッド (最後のスレッドなので、1st decode では呼ばれない)
static DWORD WINAPI thread_decode_gpu(LPVOID lpParameter) static DWORD WINAPI thread_decode_gpu(LPVOID lpParameter)
{ {
unsigned char *s_buf, *p_buf; unsigned char *s_buf, *g_buf;
unsigned short *factor; unsigned short *factor;
int i, j, block_lost, max_num, chunk_num; int i, j, block_lost;
int src_num; int src_num;
unsigned int unit_size, len, off, chunk_size; unsigned int unit_size;
HANDLE hRun, hEnd; HANDLE hRun, hEnd;
RS_TH *th; RS_TH *th;
#ifdef TIMER #ifdef TIMER
@@ -296,18 +296,14 @@ unsigned int time_start2, time_encode2 = 0, loop_count2 = 0;
#endif #endif
th = (RS_TH *)lpParameter; th = (RS_TH *)lpParameter;
p_buf = th->buf; g_buf = th->buf;
unit_size = th->size; unit_size = th->size;
chunk_size = th->len;
block_lost = th->count; block_lost = th->count;
hRun = th->run; hRun = th->run;
hEnd = th->end; hEnd = th->end;
//_mm_sfence(); //_mm_sfence();
SetEvent(hEnd); // 設定完了を通知する SetEvent(hEnd); // 設定完了を通知する
chunk_num = (unit_size + chunk_size - 1) / chunk_size;
max_num = chunk_num * block_lost;
WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ
while (th->now < INT_MAX / 2){ while (th->now < INT_MAX / 2){
#ifdef TIMER #ifdef TIMER
@@ -325,17 +321,10 @@ time_start2 = GetTickCount();
InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する
} }
// スレッドごとに復元する消失ブロックの chunk を変える // スレッドごとに復元する消失ブロックを変える
len = chunk_size; while ((j = InterlockedIncrement(&(th->now))) < block_lost){ // j = ++th_now
while ((j = InterlockedIncrement(&(th->now))) < max_num){ // j = ++th_now // 倍率は逆行列から部分的にコピーする
off = j / block_lost; // chunk の番号 i = gpu_multiply_blocks(src_num, factor + source_num * j, g_buf + (size_t)unit_size * j, unit_size);
j = j % block_lost; // lost block の番号
off *= chunk_size; // chunk の位置
if (off + len > unit_size)
len = unit_size - off; // 最後の chunk だけサイズが異なるかも
// VRAM上のソース・ブロックごとにパリティを追加していく
i = gpu_multiply_chunks(src_num, factor + source_num * j, p_buf + (size_t)unit_size * j + off, off, len);
if (i != 0){ if (i != 0){
th->len = i; th->len = i;
InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する
@@ -359,7 +348,6 @@ time_encode2 += GetTickCount() - time_start2;
WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ
} }
#ifdef TIMER #ifdef TIMER
loop_count2 /= chunk_num; // chunk数で割ってブロック数にする
printf("gpu-thread :\n"); printf("gpu-thread :\n");
if (time_encode2 > 0){ if (time_encode2 > 0){
i = (int)((__int64)loop_count2 * unit_size * 125 / ((__int64)time_encode2 * 131072)); i = (int)((__int64)loop_count2 * unit_size * 125 / ((__int64)time_encode2 * 131072));
@@ -575,16 +563,9 @@ int decode_method2( // ソース・データを全て読み込む場合
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
//len = ((len + 2) / 3 + (sse_unit - 1)) & ~(sse_unit - 1); // 1/3の実験用 //len = ((len + 2) / 3 + (sse_unit - 1)) & ~(sse_unit - 1); // 1/3の実験用
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num1(part_num); // 読み込み中はスレッド数を減らす
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > part_num)
cpu_num1 = part_num;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num == 1)) if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない
#ifdef TIMER #ifdef TIMER
printf("\n read all blocks, and keep some recovering blocks\n"); printf("\n read all blocks, and keep some recovering blocks\n");
@@ -1020,16 +1001,9 @@ int decode_method3( // 復元するブロックを全て保持できる場合
prog_base = (__int64)(source_num + prog_write) * block_lost + prog_read * source_num; // ブロックの合計掛け算個数 + 読み書き回数 prog_base = (__int64)(source_num + prog_write) * block_lost + prog_read * source_num; // ブロックの合計掛け算個数 + 読み書き回数
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num1(block_lost); // 読み込み中はスレッド数を減らす
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > block_lost)
cpu_num1 = block_lost;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num == 1)) if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない
#ifdef TIMER #ifdef TIMER
printf("\n read some blocks, and keep all recovering blocks\n"); printf("\n read some blocks, and keep all recovering blocks\n");
@@ -1364,27 +1338,29 @@ int decode_method4( // 全てのブロックを断片的に保持する場合 (G
parity_ctx_r *p_blk, // パリティ・ブロックの情報 parity_ctx_r *p_blk, // パリティ・ブロックの情報
unsigned short *mat) unsigned short *mat)
{ {
unsigned char *buf = NULL, *p_buf, *work_buf, *hash; unsigned char *buf = NULL, *p_buf, *g_buf, *work_buf, *hash;
unsigned short *id; unsigned short *id;
int err = 0, i, j, last_file, chunk_num, recv_now; int err = 0, i, j, last_file, chunk_num, recv_now;
int cpu_num1, src_off, src_num, src_max, vram_max; int cpu_num1, src_off, src_num, src_max;
int cpu_num2, vram_max, cpu_end, gpu_end, th_act;
unsigned int io_size, unit_size, len, block_off; unsigned int io_size, unit_size, len, block_off;
unsigned int time_last, prog_read, prog_write; unsigned int time_last, prog_read, prog_write;
__int64 file_off, prog_num = 0, prog_base; __int64 file_off, prog_num = 0, prog_base;
HANDLE hFile = NULL; HANDLE hFile = NULL;
HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU]; HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU];
RS_TH th[1]; RS_TH th[1], th2[1];
memset(hSub, 0, sizeof(HANDLE) * MAX_CPU); memset(hSub, 0, sizeof(HANDLE) * MAX_CPU);
id = mat + (block_lost * source_num); // 何番目の消失ソース・ブロックがどのパリティで代替されるか id = mat + (block_lost * source_num); // 何番目の消失ソース・ブロックがどのパリティで代替されるか
// 作業バッファーを確保する // 作業バッファーを確保する
// part_num を使わず、全てのブロックを保持する所がdecode_method2と異なることに注意 // part_num を使わず、全てのブロックを保持する所がdecode_method2と異なることに注意
io_size = get_io_size(source_num + block_lost, NULL, 1, MEM_UNIT); // CPU計算スレッドと GPU計算スレッドで保存先を別けるので、消失ブロック分を倍確保する
io_size = get_io_size(source_num + block_lost * 2, NULL, 1, MEM_UNIT);
//io_size = (((io_size + 1) / 2 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 2分割の実験用 //io_size = (((io_size + 1) / 2 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 2分割の実験用
//io_size = (((io_size + 2) / 3 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 3分割の実験用 //io_size = (((io_size + 2) / 3 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 3分割の実験用
unit_size = io_size + HASH_SIZE; // チェックサムの分だけ増やす unit_size = io_size + HASH_SIZE; // チェックサムの分だけ増やす
file_off = (source_num + block_lost) * (size_t)unit_size + HASH_SIZE; file_off = (source_num + block_lost * 2) * (size_t)unit_size + HASH_SIZE;
buf = _aligned_malloc((size_t)file_off, MEM_UNIT); // GPU 用の境界 buf = _aligned_malloc((size_t)file_off, MEM_UNIT); // GPU 用の境界
if (buf == NULL){ if (buf == NULL){
printf("malloc, %I64d\n", file_off); printf("malloc, %I64d\n", file_off);
@@ -1392,42 +1368,36 @@ int decode_method4( // 全てのブロックを断片的に保持する場合 (G
goto error_end; goto error_end;
} }
p_buf = buf + (size_t)unit_size * source_num; // 復元したブロックを記録する領域 p_buf = buf + (size_t)unit_size * source_num; // 復元したブロックを記録する領域
hash = p_buf + (size_t)unit_size * block_lost; g_buf = p_buf + (size_t)unit_size * block_lost; // GPUスレッド用
hash = g_buf + (size_t)unit_size * block_lost;
prog_base = (block_size + io_size - 1) / io_size; prog_base = (block_size + io_size - 1) / io_size;
prog_read = (block_lost + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする prog_read = (block_lost + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする
prog_write = (source_num + 31) / 32; prog_write = (source_num + 31) / 32;
prog_base *= (__int64)(source_num + prog_write) * block_lost + prog_read * source_num; // 全体の断片の個数 prog_base *= (__int64)(source_num + prog_write) * block_lost + prog_read * source_num; // 全体の断片の個数
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num2(block_lost, &cpu_num2); // 使用するスレッド数を調節する
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > block_lost)
cpu_num1 = block_lost;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num <= 2)) if ((src_max < CACHE_MIN_NUM) || (src_max > CACHE_MAX_NUM))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = CACHE_MAX_NUM; // 不明または極端な場合は、規定値にする
//cpu_num1 = 0; // 2nd decode の実験用に 1st decode を停止する
#ifdef TIMER #ifdef TIMER
printf("\n read all blocks, and keep all recovering blocks (GPU)\n"); printf("\n read all blocks, and keep all recovering blocks (GPU)\n");
printf("buffer size = %I64d MB, io_size = %d, split = %d\n", file_off >> 20, io_size, (block_size + io_size - 1) / io_size); printf("buffer size = %I64d MB, io_size = %d, split = %d\n", file_off >> 20, io_size, (block_size + io_size - 1) / io_size);
printf("cache: limit size = %d, chunk_size = %d, split = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num); printf("cache: limit size = %d, chunk_size = %d, chunk_num = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num);
printf("unit_size = %d, cpu_num1 = %d, src_max = %d\n", unit_size, cpu_num1, src_max); printf("unit_size = %d, cpu_num1 = %d, cpu_num2 = %d\n", unit_size, cpu_num1, cpu_num2);
#endif #endif
// OpenCL の初期化 // OpenCL の初期化
vram_max = source_num; vram_max = source_num;
i = init_OpenCL(unit_size, len, &vram_max); i = init_OpenCL(unit_size, &vram_max);
if (i != 0){ if (i != 0){
if (i != 3) // GPU が見つからなかった場合はエラー表示しない if (i != 3) // GPU が見つからなかった場合はエラー表示しない
printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8); printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
i = free_OpenCL(); i = free_OpenCL();
if (i != 0) if (i != 0)
printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8); printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
OpenCL_method = 0; // GPU を使わない設定にする OpenCL_method = 0; // GPU を使えなかった印
// GPU を使わずに計算を続行する場合は以下をコメントアウト
err = -2; // CPU だけの方式に切り替える err = -2; // CPU だけの方式に切り替える
goto error_end; goto error_end;
} }
@@ -1437,10 +1407,14 @@ int decode_method4( // 全てのブロックを断片的に保持する場合 (G
// マルチ・スレッドの準備をする // マルチ・スレッドの準備をする
th->buf = p_buf; th->buf = p_buf;
th2->buf = g_buf;
th->size = unit_size; th->size = unit_size;
th2->size = unit_size;
th->count = block_lost; th->count = block_lost;
th->len = len; // chunk size th2->count = block_lost;
for (j = 0; j < cpu_num; j++){ // サブ・スレッドごとに th->len = len ; // chunk size
th2->len = 0; // GPUのエラー通知用にする
for (j = 0; j < cpu_num2; j++){ // サブ・スレッドごとに
hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする
if (hRun[j] == NULL){ if (hRun[j] == NULL){
print_win32_err(); print_win32_err();
@@ -1457,12 +1431,13 @@ int decode_method4( // 全てのブロックを断片的に保持する場合 (G
goto error_end; goto error_end;
} }
// サブ・スレッドを起動する // サブ・スレッドを起動する
if (j == cpu_num2 - 1){ // 最後のスレッドを GPU 管理用にする
th2->run = hRun[j];
th2->end = hEnd[j];
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode_gpu, (LPVOID)th2, 0, NULL);
} else {
th->run = hRun[j]; th->run = hRun[j];
th->end = hEnd[j]; th->end = hEnd[j];
//_mm_sfence(); // メモリーへの書き込みを完了してからスレッドを起動する
if ((j == cpu_num - 1) && (OpenCL_method != 0)){ // 最後のスレッドを GPU 管理用にする
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode_gpu, (LPVOID)th, 0, NULL);
} else {
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode3, (LPVOID)th, 0, NULL); hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode3, (LPVOID)th, 0, NULL);
} }
if (hSub[j] == NULL){ if (hSub[j] == NULL){
@@ -1475,7 +1450,6 @@ int decode_method4( // 全てのブロックを断片的に保持する場合 (G
} }
WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない) WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない)
} }
th->len = 0; // GPUのエラー通知用にする
// ブロック断片を読み込んで、消失ブロック断片を復元する // ブロック断片を読み込んで、消失ブロック断片を復元する
print_progress_text(0, "Recovering slice"); print_progress_text(0, "Recovering slice");
@@ -1629,6 +1603,7 @@ skip_count++;
time_read += GetTickCount() - time_start; time_read += GetTickCount() - time_start;
#endif #endif
memset(g_buf, 0, (size_t)unit_size * block_lost); // 待機中に GPU用の領域をゼロ埋めしておく
WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ
src_off += 1; // 計算を開始するソース・ブロックの番号 src_off += 1; // 計算を開始するソース・ブロックの番号
if (src_off > 0){ // 計算不要なソース・ブロックはとばす if (src_off > 0){ // 計算不要なソース・ブロックはとばす
@@ -1647,74 +1622,164 @@ skip_count++;
j = (src_off * 1000) / source_num; j = (src_off * 1000) / source_num;
printf("partial decode = %d / %d (%d.%d%%), read = %d, skip = %d\n", src_off, source_num, j / 10, j % 10, read_count, skip_count); printf("partial decode = %d / %d (%d.%d%%), read = %d, skip = %d\n", src_off, source_num, j / 10, j % 10, read_count, skip_count);
#endif #endif
recv_now = -1; // 消失ブロックの本来のソース番号 recv_now = -1; // 消失ブロックの本来のソース番号
last_file = -1; last_file = -1;
th2->size = 0; // 計算前の状態にしておく (th->size は既に 0 になってる)
// GPU と CPU のどちらに最適化するかが難しい cpu_end = gpu_end = 0;
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_num > vram_max){ // VRAM に収まらない場合は、VRAM のサイズに応じて分割する
src_num = vram_max & ~1; // 減らして偶数にする(元が奇数なら分割数が増えるかも)
i = (source_num - src_off + src_num - 1) / src_num; // 何回に分けて処理するか
src_num = (source_num - src_off + i - 1) / i; // 一度に処理する量を平均化する
src_num = (src_num + 1) & ~1; // 増やして偶数にする
}
#ifdef TIMER #ifdef TIMER
printf("remain = %d, src_off = %d, src_num = %d\n", source_num - src_off, src_off, src_num); printf("remain = %d, src_off = %d, src_max = %d\n", source_num - src_off, src_off, src_max);
#endif #endif
while (src_off < source_num){ while (src_off < source_num){
// ソース・ブロックを何個ずつ処理するか // GPUスレッドと CPUスレッドのどちらかが待機中になるまで待つ
if (src_off + src_num > source_num){ do {
src_num = source_num - src_off; th_act = 0;
#ifdef TIMER // CPUスレッドの動作状況を調べる
printf("last1: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForMultipleObjects(cpu_num2 - 1, hEnd, TRUE, 0) == WAIT_TIMEOUT){
#endif th_act |= 1; // CPUスレッドが動作中
} else if (src_off + src_num * 2 - 1 >= source_num){ } else if (th->size > 0){ // CPUスレッドの計算量を加算する
src_num = source_num - src_off; prog_num += th->size * block_lost;
if (src_num > vram_max){ // VRAM のサイズまでにする th->size = 0;
src_num = (src_num + 1) / 2; // 半分にする
src_num = (src_num + 1) & ~1; // 偶数にする
} }
#ifdef TIMER // GPUスレッドの動作状況を調べる
printf("last2: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForSingleObject(hEnd[cpu_num2 - 1], 0) == WAIT_TIMEOUT){
#endif th_act |= 2; // GPUスレッドが動作中
} else if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
} }
prog_num += th2->size * block_lost;
// GPU と CPU がスレッドごとに消失ブロックを計算する th2->size = 0;
th->buf = buf + (size_t)unit_size * src_off;
th->mat = mat + src_off;
th->size = src_num;
th->now = -1; // 初期値 - 1
//_mm_sfence(); // メモリーへの書き込みを完了してからスレッドを再開する
for (j = 0; j < cpu_num; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
} }
if (th_act == 3){ // 両方が動作中なら
// サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待 // サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){ while (WaitForMultipleObjects(cpu_num2, hEnd, FALSE, UPDATE_TIME) == WAIT_TIMEOUT){
// th-now が最高値なので、計算が終わってるのは th-now + 1 - cpu_num 個となる // th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
j = th->now + 1 - cpu_num; i = th2->now;
if (j < 0) if (i < 0){
i = 0;
} else {
i *= th2->size;
}
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0; j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず) // 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + src_num * j) * 1000) / prog_base))){ if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
time_last = GetTickCount(); time_last = GetTickCount();
} }
if (th->len != 0){ // エラー発生 }
i = th->len; } while (th_act == 3);
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1; // どちらかのスレッドで消失ブロックを計算する
goto error_end; if ((th_act & 1) == 0){ // CPUスレッドを優先的に開始する
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_off + src_num * 2 - 1 >= source_num){
src_num = source_num - src_off;
#ifdef TIMER
printf("CPU last: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
cpu_end += src_num;
th->buf = buf + (size_t)unit_size * src_off;
th->mat = mat + src_off;
th->size = src_num;
th->now = -1; // CPUスレッドの初期値 - 1
//_mm_sfence();
for (j = 0; j < cpu_num2 - 1; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
}
} else { // CPUスレッドが動作中なら、GPUスレッドを開始する
src_num = (source_num - src_off) * gpu_end / (cpu_end + gpu_end); // 残りブロック数に対する割合
if (src_num < src_max){
if (gpu_end == 0){ // 最初に負担するブロック数は CPUスレッドの 2倍まで
src_num = (source_num - src_off) / (cpu_num2 + 2);
if (src_num < src_max){
src_num = src_max;
} else if (src_num > src_max * 2){
src_num = src_max * 2;
}
} else if (gpu_end * 2 < cpu_end){ // GPU が遅い場合は最低負担量も減らす
if (gpu_end * 4 < cpu_end){
if (src_num < src_max / 4)
src_num = src_max / 4;
} else if (src_num < src_max / 2){
src_num = src_max / 2;
}
} else {
src_num = src_max; // 最低でも CPUスレッドと同じ量を担当する
}
}
if (src_num > vram_max)
src_num = vram_max;
if (src_off + src_num >= source_num){
src_num = source_num - src_off;
#ifdef TIMER
printf("GPU last 1: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
} else if (src_off + src_num + src_max > source_num){
src_num = source_num - src_off - src_max;
if ((src_num < src_max) && (src_num + src_max <= vram_max) && (gpu_end * 2 > cpu_end)){
src_num += src_max; // GPU担当量が少なくて、余裕がある場合は、残りも全て任せる
#ifdef TIMER
printf("GPU last +: src_off = %d, src_num = %d + %d\n", src_off, src_num - src_max, src_max);
} else {
printf("GPU last 2: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
#ifdef TIMER
} else {
printf("GPU: remain = %d, src_off = %d, src_num = %d\n", source_num - src_off, src_off, src_num);
#endif
}
gpu_end += src_num;
th2->buf = buf + (size_t)unit_size * src_off;
th2->mat = mat + src_off;
th2->size = src_num;
th2->now = -1; // GPUスレッドの初期値 - 1
//_mm_sfence();
ResetEvent(hEnd[cpu_num2 - 1]); // リセットしておく
SetEvent(hRun[cpu_num2 - 1]); // サブ・スレッドに計算を開始させる
} }
// 経過表示 // 経過表示
prog_num += src_num * block_lost;
if (GetTickCount() - time_last >= UPDATE_TIME){ if (GetTickCount() - time_last >= UPDATE_TIME){
if (print_progress((int)((prog_num * 1000) / prog_base))){ if (th2->size == 0){
i = 0;
} else {
// th2-now がGPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
@@ -1724,6 +1789,50 @@ skip_count++;
src_off += src_num; src_off += src_num;
} }
// 全スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num2, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){
if (th2->size == 0){
i = 0;
} else {
// th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2;
goto error_end;
}
time_last = GetTickCount();
}
if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
}
prog_num += th2->size * block_lost;
}
if (th->size > 0) // CPUスレッドの計算量を加算する
prog_num += th->size * block_lost;
#ifdef TIMER #ifdef TIMER
time_start = GetTickCount(); time_start = GetTickCount();
#endif #endif
@@ -1738,6 +1847,8 @@ time_start = GetTickCount();
} }
//printf(" lost block[%d] = source block[%d]\n", i, recv_now); //printf(" lost block[%d] = source block[%d]\n", i, recv_now);
// CPUスレッドと GPUスレッドの計算結果を合わせる
galois_align_xor(g_buf + (size_t)unit_size * i, work_buf, unit_size);
// 復元されたソース・ブロックのチェックサムを検証する // 復元されたソース・ブロックのチェックサムを検証する
checksum16_return(work_buf, hash, io_size); checksum16_return(work_buf, hash, io_size);
if (memcmp(work_buf + io_size, hash, HASH_SIZE) != 0){ if (memcmp(work_buf + io_size, hash, HASH_SIZE) != 0){
@@ -1817,7 +1928,8 @@ if (prog_num != prog_base)
error_end: error_end:
InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する
for (j = 0; j < cpu_num; j++){ InterlockedExchange(&(th2->now), INT_MAX / 2);
for (j = 0; j < cpu_num2; j++){
if (hSub[j]){ // サブ・スレッドを終了させる if (hSub[j]){ // サブ・スレッドを終了させる
SetEvent(hRun[j]); SetEvent(hRun[j]);
WaitForSingleObject(hSub[j], INFINITE); WaitForSingleObject(hSub[j], INFINITE);
@@ -1843,31 +1955,33 @@ int decode_method5( // 復元するブロックだけ保持する場合 (GPU対
parity_ctx_r *p_blk, // パリティ・ブロックの情報 parity_ctx_r *p_blk, // パリティ・ブロックの情報
unsigned short *mat) unsigned short *mat)
{ {
unsigned char *buf = NULL, *p_buf, *work_buf, *hash; unsigned char *buf = NULL, *p_buf, *g_buf, *work_buf, *hash;
unsigned short *id; unsigned short *id;
int err = 0, i, j, last_file, chunk_num, recv_now; int err = 0, i, j, last_file, chunk_num, recv_now;
int source_off, read_num, parity_now; int source_off, read_num, parity_now;
int cpu_num1, src_off, src_num, src_max, vram_max; int cpu_num1, src_off, src_num, src_max;
int cpu_num2, vram_max, cpu_end, gpu_end, th_act;
unsigned int unit_size, len; unsigned int unit_size, len;
unsigned int time_last, prog_read, prog_write; unsigned int time_last, prog_read, prog_write;
__int64 file_off, prog_num = 0, prog_base; __int64 file_off, prog_num = 0, prog_base;
HANDLE hFile = NULL; HANDLE hFile = NULL;
HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU]; HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU];
RS_TH th[1]; RS_TH th[1], th2[1];
memset(hSub, 0, sizeof(HANDLE) * MAX_CPU); memset(hSub, 0, sizeof(HANDLE) * MAX_CPU);
id = mat + (block_lost * source_num); // 何番目の消失ソース・ブロックがどのパリティで代替されるか id = mat + (block_lost * source_num); // 何番目の消失ソース・ブロックがどのパリティで代替されるか
unit_size = (block_size + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1); // MEM_UNIT の倍数にする unit_size = (block_size + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1); // MEM_UNIT の倍数にする
// 作業バッファーを確保する // 作業バッファーを確保する
read_num = read_block_num(block_lost, 1, MEM_UNIT); // ソース・ブロックを何個読み込むか // CPU計算スレッドと GPU計算スレッドで保存先を別けるので、消失ブロック分を倍確保する
read_num = read_block_num(block_lost * 2, 1, MEM_UNIT); // ソース・ブロックを何個読み込むか
if (read_num == 0){ if (read_num == 0){
//printf("cannot keep enough blocks, use another method\n"); //printf("cannot keep enough blocks, use another method\n");
return -4; // スライスを分割して処理しないと無理 return -4; // スライスを分割して処理しないと無理
} }
//read_num = (read_num + 1) / 2 + 1; // 2分割の実験用 //read_num = (read_num + 1) / 2 + 1; // 2分割の実験用
//read_num = (read_num + 2) / 3 + 1; // 3分割の実験用 //read_num = (read_num + 2) / 3 + 1; // 3分割の実験用
file_off = (read_num + block_lost) * (size_t)unit_size + HASH_SIZE; file_off = (read_num + block_lost * 2) * (size_t)unit_size + HASH_SIZE;
buf = _aligned_malloc((size_t)file_off, MEM_UNIT); // GPU 用の境界 buf = _aligned_malloc((size_t)file_off, MEM_UNIT); // GPU 用の境界
if (buf == NULL){ if (buf == NULL){
printf("malloc, %I64d\n", file_off); printf("malloc, %I64d\n", file_off);
@@ -1875,41 +1989,35 @@ int decode_method5( // 復元するブロックだけ保持する場合 (GPU対
goto error_end; goto error_end;
} }
p_buf = buf + (size_t)unit_size * read_num; // パリティ・ブロックを記録する領域 p_buf = buf + (size_t)unit_size * read_num; // パリティ・ブロックを記録する領域
hash = p_buf + (size_t)unit_size * block_lost; g_buf = p_buf + (size_t)unit_size * block_lost; // GPUスレッド用
hash = g_buf + (size_t)unit_size * block_lost;
prog_read = (block_lost + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする prog_read = (block_lost + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする
prog_write = (source_num + 31) / 32; prog_write = (source_num + 31) / 32;
prog_base = (__int64)(source_num + prog_write) * block_lost + prog_read * source_num; // ブロックの合計掛け算個数 + 書き込み回数 prog_base = (__int64)(source_num + prog_write) * block_lost + prog_read * source_num; // ブロックの合計掛け算個数 + 書き込み回数
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num2(block_lost, &cpu_num2); // 使用するスレッド数を調節する
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > block_lost)
cpu_num1 = block_lost;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num <= 2)) if ((src_max < CACHE_MIN_NUM) || (src_max > CACHE_MAX_NUM))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = CACHE_MAX_NUM; // 不明または極端な場合は、規定値にする
//cpu_num1 = 0; // 2nd decode の実験用に 1st decode を停止する
#ifdef TIMER #ifdef TIMER
printf("\n read some blocks, and keep all recovering blocks (GPU)\n"); printf("\n read some blocks, and keep all recovering blocks (GPU)\n");
printf("buffer size = %I64d MB, read_num = %d, round = %d\n", file_off >> 20, read_num, (source_num + read_num - 1) / read_num); printf("buffer size = %I64d MB, read_num = %d, round = %d\n", file_off >> 20, read_num, (source_num + read_num - 1) / read_num);
printf("cache: limit size = %d, chunk_size = %d, split = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num); printf("cache: limit size = %d, chunk_size = %d, chunk_num = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num);
printf("unit_size = %d, cpu_num1 = %d, src_max = %d\n", unit_size, cpu_num1, src_max); printf("unit_size = %d, cpu_num1 = %d, cpu_num2 = %d\n", unit_size, cpu_num1, cpu_num2);
#endif #endif
// OpenCL の初期化 // OpenCL の初期化
vram_max = read_num; // 読み込める分だけにする vram_max = read_num; // 読み込める分だけにする
i = init_OpenCL(unit_size, len, &vram_max); i = init_OpenCL(unit_size, &vram_max);
if (i != 0){ if (i != 0){
if (i != 3) // GPU が見つからなかった場合はエラー表示しない if (i != 3) // GPU が見つからなかった場合はエラー表示しない
printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8); printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
i = free_OpenCL(); i = free_OpenCL();
if (i != 0) if (i != 0)
printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8); printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
OpenCL_method = 0; // GPU を使わない設定にする OpenCL_method = 0; // GPU を使えなかった印
// GPU を使わずに計算を続行する場合は以下をコメントアウト
err = -3; // CPU だけの方式に切り替える err = -3; // CPU だけの方式に切り替える
goto error_end; goto error_end;
} }
@@ -1919,10 +2027,14 @@ int decode_method5( // 復元するブロックだけ保持する場合 (GPU対
// マルチ・スレッドの準備をする // マルチ・スレッドの準備をする
th->buf = p_buf; th->buf = p_buf;
th2->buf = g_buf;
th->size = unit_size; th->size = unit_size;
th2->size = unit_size;
th->count = block_lost; th->count = block_lost;
th->len = len; // chunk size th2->count = block_lost;
for (j = 0; j < cpu_num; j++){ // サブ・スレッドごとに th->len = len ; // chunk size
th2->len = 0; // GPUのエラー通知用にする
for (j = 0; j < cpu_num2; j++){ // サブ・スレッドごとに
hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする
if (hRun[j] == NULL){ if (hRun[j] == NULL){
print_win32_err(); print_win32_err();
@@ -1939,12 +2051,13 @@ int decode_method5( // 復元するブロックだけ保持する場合 (GPU対
goto error_end; goto error_end;
} }
// サブ・スレッドを起動する // サブ・スレッドを起動する
if (j == cpu_num2 - 1){ // 最後のスレッドを GPU 管理用にする
th2->run = hRun[j];
th2->end = hEnd[j];
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode_gpu, (LPVOID)th2, 0, NULL);
} else {
th->run = hRun[j]; th->run = hRun[j];
th->end = hEnd[j]; th->end = hEnd[j];
//_mm_sfence(); // メモリーへの書き込みを完了してからスレッドを起動する
if ((j == cpu_num - 1) && (OpenCL_method != 0)){ // 最後のスレッドを GPU 管理用にする
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode_gpu, (LPVOID)th, 0, NULL);
} else {
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode3, (LPVOID)th, 0, NULL); hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_decode3, (LPVOID)th, 0, NULL);
} }
if (hSub[j] == NULL){ if (hSub[j] == NULL){
@@ -1957,7 +2070,6 @@ int decode_method5( // 復元するブロックだけ保持する場合 (GPU対
} }
WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない) WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない)
} }
th->len = 0; // GPUのエラー通知用にする
// 何回かに別けてブロックを読み込んで、消失ブロックを少しずつ復元する // 何回かに別けてブロックを読み込んで、消失ブロックを少しずつ復元する
print_progress_text(0, "Recovering slice"); print_progress_text(0, "Recovering slice");
@@ -2086,6 +2198,8 @@ read_count++;
time_read += GetTickCount() - time_start; time_read += GetTickCount() - time_start;
#endif #endif
if (source_off == 0)
memset(g_buf, 0, (size_t)unit_size * block_lost); // 待機中に GPU用の領域をゼロ埋めしておく
WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ
src_off += 1; // 計算を開始するソース・ブロックの番号 src_off += 1; // 計算を開始するソース・ブロックの番号
if (src_off == 0) // 1st decode しなかった場合src_off = 0は、消失ブロックをゼロ埋めする if (src_off == 0) // 1st decode しなかった場合src_off = 0は、消失ブロックをゼロ埋めする
@@ -2094,75 +2208,165 @@ time_read += GetTickCount() - time_start;
j = (src_off - source_off) * 1000 / read_num; j = (src_off - source_off) * 1000 / read_num;
printf("partial decode = %d / %d (%d.%d%%), source_off = %d, read = %d\n", src_off - source_off, read_num, j / 10, j % 10, source_off, read_count); printf("partial decode = %d / %d (%d.%d%%), source_off = %d, read = %d\n", src_off - source_off, read_num, j / 10, j % 10, source_off, read_count);
#endif #endif
recv_now = -1; // 消失ブロックの本来のソース番号 recv_now = -1; // 消失ブロックの本来のソース番号
last_file = -1; last_file = -1;
th2->size = 0; // 計算前の状態にしておく (th->size は既に 0 になってる)
// GPU と CPU のどちらに最適化するかが難しい cpu_end = gpu_end = 0;
src_off -= source_off; // バッファー内でのソース・ブロックの位置にする src_off -= source_off; // バッファー内でのソース・ブロックの位置にする
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_num > vram_max){ // VRAM に収まらない場合は、VRAM のサイズに応じて分割する
src_num = vram_max & ~1; // 減らして偶数にする(元が奇数なら分割数が増えるかも)
i = (read_num - src_off + src_num - 1) / src_num; // 何回に分けて処理するか
src_num = (read_num - src_off + i - 1) / i; // 一度に処理する量を平均化する
src_num = (src_num + 1) & ~1; // 増やして偶数にする
}
#ifdef TIMER #ifdef TIMER
printf("remain = %d, src_off = %d, src_num = %d\n", read_num - src_off, src_off, src_num); printf("remain = %d, src_off = %d, src_max = %d\n", read_num - src_off, src_off, src_max);
#endif #endif
while (src_off < read_num){ while (src_off < read_num){
// ソース・ブロックを何個ずつ処理するか // GPUスレッドと CPUスレッドのどちらかが待機中になるまで待つ
if (src_off + src_num > read_num){ do {
src_num = read_num - src_off; th_act = 0;
#ifdef TIMER // CPUスレッドの動作状況を調べる
printf("last1: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForMultipleObjects(cpu_num2 - 1, hEnd, TRUE, 0) == WAIT_TIMEOUT){
#endif th_act |= 1; // CPUスレッドが動作中
} else if (src_off + src_num * 2 - 1 >= read_num){ } else if (th->size > 0){ // CPUスレッドの計算量を加算する
src_num = read_num - src_off; prog_num += th->size * block_lost;
if (src_num > vram_max){ // VRAM のサイズまでにする th->size = 0;
src_num = (src_num + 1) / 2; // 半分にする
src_num = (src_num + 1) & ~1; // 偶数にする
} }
#ifdef TIMER // GPUスレッドの動作状況を調べる
printf("last2: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForSingleObject(hEnd[cpu_num2 - 1], 0) == WAIT_TIMEOUT){
#endif th_act |= 2; // GPUスレッドが動作中
} else if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
} }
prog_num += th2->size * block_lost;
// GPU と CPU がスレッドごとに消失ブロックを計算する th2->size = 0;
th->buf = buf + (size_t)unit_size * src_off;
th->mat = mat + (source_off + src_off); // ソース・ブロックの番号にする
th->size = src_num;
th->now = -1; // 初期値 - 1
//_mm_sfence(); // メモリーへの書き込みを完了してからスレッドを再開する
for (j = 0; j < cpu_num; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
} }
if (th_act == 3){ // 両方が動作中なら
// サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待 // サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){ while (WaitForMultipleObjects(cpu_num2, hEnd, FALSE, UPDATE_TIME) == WAIT_TIMEOUT){
// th-now が最高値なので、計算が終わってるのは th-now + 1 - cpu_num 個となる // th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
j = th->now + 1 - cpu_num; i = th2->now;
if (j < 0) if (i < 0){
i = 0;
} else {
i *= th2->size;
}
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0; j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず) // 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + src_num * j) * 1000) / prog_base))){ if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
time_last = GetTickCount(); time_last = GetTickCount();
} }
if (th->len != 0){ // エラー発生 }
i = th->len; } while (th_act == 3);
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1; // どちらかのスレッドで消失ブロックを計算する
goto error_end; if ((th_act & 1) == 0){ // CPUスレッドを優先的に開始する
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_off + src_num * 2 - 1 >= read_num){
src_num = read_num - src_off;
#ifdef TIMER
printf("CPU last: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
cpu_end += src_num;
th->buf = buf + (size_t)unit_size * src_off;
th->mat = mat + (source_off + src_off); // ソース・ブロックの番号にする
th->size = src_num;
th->now = -1; // CPUスレッドの初期値 - 1
//_mm_sfence();
for (j = 0; j < cpu_num2 - 1; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
}
} else { // CPUスレッドが動作中なら、GPUスレッドを開始する
src_num = (read_num - src_off) * gpu_end / (cpu_end + gpu_end); // 残りブロック数に対する割合
if (src_num < src_max){
if (gpu_end == 0){ // 最初に負担するブロック数は CPUスレッドの 2倍まで
src_num = (read_num - src_off) / (cpu_num2 + 2);
if (src_num < src_max){
src_num = src_max;
} else if (src_num > src_max * 2){
src_num = src_max * 2;
}
} else if (gpu_end * 2 < cpu_end){ // GPU が遅い場合は最低負担量も減らす
if (gpu_end * 4 < cpu_end){
if (src_num < src_max / 4)
src_num = src_max / 4;
} else if (src_num < src_max / 2){
src_num = src_max / 2;
}
} else {
src_num = src_max; // 最低でも CPUスレッドと同じ量を担当する
}
}
if (src_num > vram_max)
src_num = vram_max;
if (src_off + src_num >= read_num){
src_num = read_num - src_off;
#ifdef TIMER
printf("GPU last 1: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
} else if (src_off + src_num + src_max > read_num){
src_num = read_num - src_off - src_max;
if ((src_num < src_max) && (src_num + src_max <= vram_max) && (gpu_end * 2 > cpu_end)){
src_num += src_max; // GPU担当量が少なくて、余裕がある場合は、残りも全て任せる
#ifdef TIMER
printf("GPU last +: src_off = %d, src_num = %d + %d\n", src_off, src_num - src_max, src_max);
} else {
printf("GPU last 2: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
#ifdef TIMER
} else {
printf("GPU: remain = %d, src_off = %d, src_num = %d\n", read_num - src_off, src_off, src_num);
#endif
}
gpu_end += src_num;
th2->buf = buf + (size_t)unit_size * src_off;
th2->mat = mat + (source_off + src_off); // ソース・ブロックの番号にする
th2->size = src_num;
th2->now = -1; // GPUスレッドの初期値 - 1
//_mm_sfence();
ResetEvent(hEnd[cpu_num2 - 1]); // リセットしておく
SetEvent(hRun[cpu_num2 - 1]); // サブ・スレッドに計算を開始させる
} }
// 経過表示 // 経過表示
prog_num += src_num * block_lost;
if (GetTickCount() - time_last >= UPDATE_TIME){ if (GetTickCount() - time_last >= UPDATE_TIME){
if (print_progress((int)((prog_num * 1000) / prog_base))){ if (th2->size == 0){
i = 0;
} else {
// th2-now がGPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
@@ -2172,6 +2376,50 @@ time_read += GetTickCount() - time_start;
src_off += src_num; src_off += src_num;
} }
// 全スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num2, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){
if (th2->size == 0){
i = 0;
} else {
// th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2;
goto error_end;
}
time_last = GetTickCount();
}
if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
}
prog_num += th2->size * block_lost;
}
if (th->size > 0) // CPUスレッドの計算量を加算する
prog_num += th->size * block_lost;
source_off += read_num; source_off += read_num;
} }
@@ -2189,6 +2437,8 @@ time_start = GetTickCount();
} }
//printf(" lost block[%d] = source block[%d]\n", i, recv_now); //printf(" lost block[%d] = source block[%d]\n", i, recv_now);
// CPUスレッドと GPUスレッドの計算結果を合わせる
galois_align_xor(g_buf + (size_t)unit_size * i, work_buf, unit_size);
// 復元されたソース・ブロックのチェックサムを検証する // 復元されたソース・ブロックのチェックサムを検証する
checksum16_return(work_buf, hash, unit_size - HASH_SIZE); checksum16_return(work_buf, hash, unit_size - HASH_SIZE);
if (memcmp(work_buf + unit_size - HASH_SIZE, hash, HASH_SIZE) != 0){ if (memcmp(work_buf + unit_size - HASH_SIZE, hash, HASH_SIZE) != 0){
@@ -2252,7 +2502,8 @@ if (prog_num != prog_base)
error_end: error_end:
InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する
for (j = 0; j < cpu_num; j++){ InterlockedExchange(&(th2->now), INT_MAX / 2);
for (j = 0; j < cpu_num2; j++){
if (hSub[j]){ // サブ・スレッドを終了させる if (hSub[j]){ // サブ・スレッドを終了させる
SetEvent(hRun[j]); SetEvent(hRun[j]);
WaitForSingleObject(hSub[j], INFINITE); WaitForSingleObject(hSub[j], INFINITE);

View File

@@ -1,5 +1,5 @@
// rs_encode.c // rs_encode.c
// Copyright : 2023-09-21 Yutaka Sawada // Copyright : 2023-10-29 Yutaka Sawada
// License : GPL // License : GPL
#ifndef _UNICODE #ifndef _UNICODE
@@ -295,11 +295,11 @@ printf(" 2nd encode %d.%03d sec, %d loop, %d MB/s\n", time_encode2b / 1000, time
// GPU 対応のサブ・スレッド (最後のスレッドなので、1st encode では呼ばれない) // GPU 対応のサブ・スレッド (最後のスレッドなので、1st encode では呼ばれない)
static DWORD WINAPI thread_encode_gpu(LPVOID lpParameter) static DWORD WINAPI thread_encode_gpu(LPVOID lpParameter)
{ {
unsigned char *s_buf, *p_buf; unsigned char *s_buf, *g_buf;
unsigned short *constant, *factor; unsigned short *constant, *factor;
int i, j, max_num, chunk_num; int i, j;
int src_off, src_num; int src_off, src_num;
unsigned int unit_size, len, off, chunk_size; unsigned int unit_size;
HANDLE hRun, hEnd; HANDLE hRun, hEnd;
RS_TH *th; RS_TH *th;
#ifdef TIMER #ifdef TIMER
@@ -308,17 +308,14 @@ unsigned int time_start2, time_encode2 = 0, loop_count2 = 0;
th = (RS_TH *)lpParameter; th = (RS_TH *)lpParameter;
constant = th->mat; constant = th->mat;
p_buf = th->buf; g_buf = th->buf;
unit_size = th->size; unit_size = th->size;
chunk_size = th->len;
hRun = th->run; hRun = th->run;
hEnd = th->end; hEnd = th->end;
//_mm_sfence(); //_mm_sfence();
SetEvent(hEnd); // 設定完了を通知する SetEvent(hEnd); // 設定完了を通知する
factor = constant + source_num; factor = constant + source_num;
chunk_num = (unit_size + chunk_size - 1) / chunk_size;
max_num = chunk_num * parity_num;
WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ
while (th->now < INT_MAX / 2){ while (th->now < INT_MAX / 2){
@@ -337,27 +334,19 @@ time_start2 = GetTickCount();
InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する
} }
// スレッドごとに作成するパリティ・ブロックの chunk を変え // 一つの GPUスレッドが全てのパリティ・ブロックを処理す
len = chunk_size; while ((j = InterlockedIncrement(&(th->now))) < parity_num){ // j = ++th_now
while ((j = InterlockedIncrement(&(th->now))) < max_num){ // j = ++th_now
off = j / parity_num; // chunk の番号
j = j % parity_num; // parity の番号
off *= chunk_size; // chunk の位置
if (off + len > unit_size)
len = unit_size - off; // 最後の chunk だけサイズが異なるかも
// factor は定数行列の乗数になる // factor は定数行列の乗数になる
for (i = 0; i < src_num; i++) for (i = 0; i < src_num; i++)
factor[i] = galois_power(constant[src_off + i], first_num + j); factor[i] = galois_power(constant[src_off + i], first_num + j);
// VRAM上のソース・ブロックごとにパリティを追加していく // VRAM上のソース・ブロックごとにパリティを追加していく
i = gpu_multiply_chunks(src_num, factor, p_buf + (size_t)unit_size * j + off, off, len); i = gpu_multiply_blocks(src_num, factor, g_buf + (size_t)unit_size * j, unit_size);
if (i != 0){ if (i != 0){
th->len = i; th->len = i;
InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 3); // サブ・スレッドの計算を中断する
break; break;
} }
#ifdef TIMER #ifdef TIMER
loop_count2 += src_num; loop_count2 += src_num;
#endif #endif
@@ -375,7 +364,6 @@ time_encode2 += GetTickCount() - time_start2;
WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ WaitForSingleObject(hRun, INFINITE); // 計算開始の合図を待つ
} }
#ifdef TIMER #ifdef TIMER
loop_count2 /= chunk_num; // chunk数で割ってブロック数にする
printf("gpu-thread :\n"); printf("gpu-thread :\n");
if (time_encode2 > 0){ if (time_encode2 > 0){
i = (int)((__int64)loop_count2 * unit_size * 125 / ((__int64)time_encode2 * 131072)); i = (int)((__int64)loop_count2 * unit_size * 125 / ((__int64)time_encode2 * 131072));
@@ -664,16 +652,9 @@ int encode_method2( // ソース・データを全て読み込む場合
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
//len = ((len + 2) / 3 + (sse_unit - 1)) & ~(sse_unit - 1); // 1/3の実験用 //len = ((len + 2) / 3 + (sse_unit - 1)) & ~(sse_unit - 1); // 1/3の実験用
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num1(part_num); // 読み込み中はスレッド数を減らす
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > part_num)
cpu_num1 = part_num;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num == 1)) if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない
#ifdef TIMER #ifdef TIMER
printf("\n read all source blocks, and keep some parity blocks\n"); printf("\n read all source blocks, and keep some parity blocks\n");
@@ -1145,16 +1126,9 @@ int encode_method3( // パリティ・ブロックを全て保持して、一度
prog_base = (__int64)(source_num + prog_write) * parity_num + prog_read * source_num; // ブロックの合計掛け算個数 + 読み書き回数 prog_base = (__int64)(source_num + prog_write) * parity_num + prog_read * source_num; // ブロックの合計掛け算個数 + 読み書き回数
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num1(parity_num); // 読み込み中はスレッド数を減らす
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > parity_num)
cpu_num1 = parity_num;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num == 1)) if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない
#ifdef TIMER #ifdef TIMER
printf("\n read some source blocks, and keep all parity blocks\n"); printf("\n read some source blocks, and keep all parity blocks\n");
@@ -1424,7 +1398,7 @@ time_start = GetTickCount();
memcpy(common_buf + common_size, common_buf, common_size); // 後の半分に前半のをコピーする memcpy(common_buf + common_size, common_buf, common_size); // 後の半分に前半のをコピーする
// 最後にパリティ・ブロックのチェックサムを検証して、リカバリ・ファイルに書き込む // 最後にパリティ・ブロックのチェックサムを検証して、リカバリ・ファイルに書き込む
err = create_recovery_file_1pass(file_path, recovery_path, packet_limit, block_distri, err = create_recovery_file_1pass(file_path, recovery_path, packet_limit, block_distri,
packet_num, common_buf, common_size, footer_buf, footer_size, rcv_hFile, p_buf, unit_size); packet_num, common_buf, common_size, footer_buf, footer_size, rcv_hFile, p_buf, NULL, unit_size);
#ifdef TIMER #ifdef TIMER
time_write = GetTickCount() - time_start; time_write = GetTickCount() - time_start;
#endif #endif
@@ -1463,26 +1437,28 @@ int encode_method4( // 全てのブロックを断片的に保持する場合 (G
parity_ctx_c *p_blk, // パリティ・ブロックの情報 parity_ctx_c *p_blk, // パリティ・ブロックの情報
unsigned short *constant) // 複数ブロック分の領域を確保しておく? unsigned short *constant) // 複数ブロック分の領域を確保しておく?
{ {
unsigned char *buf = NULL, *p_buf, *work_buf, *hash; unsigned char *buf = NULL, *p_buf, *g_buf, *work_buf, *hash;
int err = 0, i, j, last_file, chunk_num; int err = 0, i, j, last_file, chunk_num;
int cpu_num1, src_off, src_num, src_max, vram_max; int cpu_num1, src_off, src_num, src_max;
int cpu_num2, vram_max, cpu_end, gpu_end, th_act;
unsigned int io_size, unit_size, len, block_off; unsigned int io_size, unit_size, len, block_off;
unsigned int time_last, prog_read, prog_write; unsigned int time_last, prog_read, prog_write;
__int64 file_off, prog_num = 0, prog_base; __int64 file_off, prog_num = 0, prog_base;
HANDLE hFile = NULL; HANDLE hFile = NULL;
HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU]; HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU];
RS_TH th[1]; RS_TH th[1], th2[1];
PHMD5 md_ctx, *md_ptr = NULL; PHMD5 md_ctx, *md_ptr = NULL;
memset(hSub, 0, sizeof(HANDLE) * MAX_CPU); memset(hSub, 0, sizeof(HANDLE) * MAX_CPU);
// 作業バッファーを確保する // 作業バッファーを確保する
// part_num を使わず、全てのブロックを保持する所がencode_method2と異なることに注意 // part_num を使わず、全てのブロックを保持する所がencode_method2と異なることに注意
io_size = get_io_size(source_num + parity_num, NULL, 1, MEM_UNIT); // CPU計算スレッドと GPU計算スレッドで保存先を別けるので、パリティ・ブロック分を倍確保する
io_size = get_io_size(source_num + parity_num * 2, NULL, 1, MEM_UNIT);
//io_size = (((io_size + 1) / 2 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 2分割の実験用 //io_size = (((io_size + 1) / 2 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 2分割の実験用
//io_size = (((io_size + 2) / 3 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 3分割の実験用 //io_size = (((io_size + 2) / 3 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE; // 3分割の実験用
unit_size = io_size + HASH_SIZE; // チェックサムの分だけ増やす unit_size = io_size + HASH_SIZE; // チェックサムの分だけ増やす
file_off = (source_num + parity_num) * (size_t)unit_size + HASH_SIZE; file_off = (source_num + parity_num * 2) * (size_t)unit_size + HASH_SIZE;
buf = _aligned_malloc((size_t)file_off, MEM_UNIT); // GPU 用の境界 buf = _aligned_malloc((size_t)file_off, MEM_UNIT); // GPU 用の境界
if (buf == NULL){ if (buf == NULL){
printf("malloc, %I64d\n", file_off); printf("malloc, %I64d\n", file_off);
@@ -1490,30 +1466,24 @@ int encode_method4( // 全てのブロックを断片的に保持する場合 (G
goto error_end; goto error_end;
} }
p_buf = buf + (size_t)unit_size * source_num; // パリティ・ブロックを記録する領域 p_buf = buf + (size_t)unit_size * source_num; // パリティ・ブロックを記録する領域
hash = p_buf + (size_t)unit_size * parity_num; g_buf = p_buf + (size_t)unit_size * parity_num; // GPUスレッド用
hash = g_buf + (size_t)unit_size * parity_num;
prog_base = (block_size + io_size - 1) / io_size; prog_base = (block_size + io_size - 1) / io_size;
prog_read = (parity_num + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする prog_read = (parity_num + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする
prog_write = (source_num + 31) / 32; prog_write = (source_num + 31) / 32;
prog_base *= (__int64)(source_num + prog_write) * parity_num + prog_read * source_num; // 全体の断片の個数 prog_base *= (__int64)(source_num + prog_write) * parity_num + prog_read * source_num; // 全体の断片の個数
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num2(parity_num, &cpu_num2); // 使用するスレッド数を調節する
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > parity_num)
cpu_num1 = parity_num;
//cpu_num1 = 0; // 2nd encode の実験用に 1st encode を停止する
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num <= 2)) if ((src_max < CACHE_MIN_NUM) || (src_max > CACHE_MAX_NUM))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = CACHE_MAX_NUM; // 不明または極端な場合は、規定値にする
//cpu_num1 = 0; // 2nd encode の実験用に 1st encode を停止する
#ifdef TIMER #ifdef TIMER
printf("\n read all source blocks, and keep all parity blocks (GPU)\n"); printf("\n read all source blocks, and keep all parity blocks (GPU)\n");
printf("buffer size = %I64d MB, io_size = %d, split = %d\n", file_off >> 20, io_size, (block_size + io_size - 1) / io_size); printf("buffer size = %I64d MB, io_size = %d, split = %d\n", file_off >> 20, io_size, (block_size + io_size - 1) / io_size);
printf("cache: limit size = %d, chunk_size = %d, chunk_num = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num); printf("cache: limit size = %d, chunk_size = %d, chunk_num = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num);
printf("unit_size = %d, cpu_num1 = %d, src_max = %d\n", unit_size, cpu_num1, src_max); printf("unit_size = %d, cpu_num1 = %d, cpu_num2 = %d\n", unit_size, cpu_num1, cpu_num2);
#endif #endif
if (io_size < block_size){ // スライスが分割される場合だけ、途中までのハッシュ値を保持する if (io_size < block_size){ // スライスが分割される場合だけ、途中までのハッシュ値を保持する
@@ -1534,15 +1504,14 @@ int encode_method4( // 全てのブロックを断片的に保持する場合 (G
// OpenCL の初期化 // OpenCL の初期化
vram_max = source_num; vram_max = source_num;
i = init_OpenCL(unit_size, len, &vram_max); i = init_OpenCL(unit_size, &vram_max);
if (i != 0){ if (i != 0){
if (i != 3) // GPU が見つからなかった場合はエラー表示しない if (i != 3) // GPU が見つからなかった場合はエラー表示しない
printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8); printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
i = free_OpenCL(); i = free_OpenCL();
if (i != 0) if (i != 0)
printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8); printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
OpenCL_method = 0; // GPU を使わない設定にする OpenCL_method = 0; // GPU を使えなかった印
// GPU を使わずに計算を続行する場合は以下をコメントアウト
err = -2; // CPU だけの方式に切り替える err = -2; // CPU だけの方式に切り替える
goto error_end; goto error_end;
} }
@@ -1552,10 +1521,14 @@ int encode_method4( // 全てのブロックを断片的に保持する場合 (G
// マルチ・スレッドの準備をする // マルチ・スレッドの準備をする
th->mat = constant; th->mat = constant;
th2->mat = constant;
th->buf = p_buf; th->buf = p_buf;
th2->buf = g_buf;
th->size = unit_size; th->size = unit_size;
th2->size = unit_size;
th->len = len; // chunk size th->len = len; // chunk size
for (j = 0; j < cpu_num; j++){ // サブ・スレッドごとに th2->len = 0; // GPUのエラー通知用にする
for (j = 0; j < cpu_num2; j++){ // サブ・スレッドごとに
hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする
if (hRun[j] == NULL){ if (hRun[j] == NULL){
print_win32_err(); print_win32_err();
@@ -1572,12 +1545,13 @@ int encode_method4( // 全てのブロックを断片的に保持する場合 (G
goto error_end; goto error_end;
} }
// サブ・スレッドを起動する // サブ・スレッドを起動する
if (j == cpu_num2 - 1){ // 最後のスレッドを GPU 管理用にする
th2->run = hRun[j];
th2->end = hEnd[j];
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode_gpu, (LPVOID)th2, 0, NULL);
} else {
th->run = hRun[j]; th->run = hRun[j];
th->end = hEnd[j]; th->end = hEnd[j];
//_mm_sfence(); // メモリーへの書き込みを完了してからスレッドを起動する
if ((j == cpu_num - 1) && (OpenCL_method != 0)){ // 最後のスレッドを GPU 管理用にする
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode_gpu, (LPVOID)th, 0, NULL);
} else {
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode3, (LPVOID)th, 0, NULL); hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode3, (LPVOID)th, 0, NULL);
} }
if (hSub[j] == NULL){ if (hSub[j] == NULL){
@@ -1590,7 +1564,6 @@ int encode_method4( // 全てのブロックを断片的に保持する場合 (G
} }
WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない) WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない)
} }
th->len = 0; // GPUのエラー通知用にする
// ソース・ブロック断片を読み込んで、パリティ・ブロック断片を作成する // ソース・ブロック断片を読み込んで、パリティ・ブロック断片を作成する
time_last = GetTickCount(); time_last = GetTickCount();
@@ -1708,6 +1681,7 @@ skip_count++;
time_read += GetTickCount() - time_start; time_read += GetTickCount() - time_start;
#endif #endif
memset(g_buf, 0, (size_t)unit_size * parity_num); // 待機中に GPU用の領域をゼロ埋めしておく
WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ
src_off += 1; // 計算を開始するソース・ブロックの番号 src_off += 1; // 計算を開始するソース・ブロックの番号
if (src_off > 0){ if (src_off > 0){
@@ -1733,72 +1707,166 @@ skip_count++;
len = io_size; len = io_size;
} }
// GPU と CPU のどちらに最適化するかが難しい th2->size = 0; // 計算前の状態にしておく (th->size は既に 0 になってる)
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する cpu_end = gpu_end = 0;
if (src_num > vram_max){ // VRAM に収まらない場合は、VRAM のサイズに応じて分割する
src_num = vram_max & ~1; // 減らして偶数にする(元が奇数なら分割数が増えるかも)
i = (source_num - src_off + src_num - 1) / src_num; // 何回に分けて処理するか
src_num = (source_num - src_off + i - 1) / i; // 一度に処理する量を平均化する
src_num = (src_num + 1) & ~1; // 増やして偶数にする
}
#ifdef TIMER #ifdef TIMER
printf("remain = %d, src_off = %d, src_num = %d\n", source_num - src_off, src_off, src_num); printf("remain = %d, src_off = %d, src_max = %d\n", source_num - src_off, src_off, src_max);
#endif #endif
while (src_off < source_num){ while (src_off < source_num){
// ソース・ブロックを何個ずつ処理するか // GPUスレッドと CPUスレッドのどちらかが待機中になるまで待つ
if (src_off + src_num > source_num){ do {
src_num = source_num - src_off; th_act = 0;
#ifdef TIMER // CPUスレッドの動作状況を調べる
printf("last1: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForMultipleObjects(cpu_num2 - 1, hEnd, TRUE, 0) == WAIT_TIMEOUT){
#endif th_act |= 1; // CPUスレッドが動作中
} else if (src_off + src_num * 2 - 1 >= source_num){ } else if (th->size > 0){ // CPUスレッドの計算量を加算する
src_num = source_num - src_off; prog_num += th->size * parity_num;
if (src_num > vram_max){ // VRAM のサイズまでにする th->size = 0;
src_num = (src_num + 1) / 2; // 半分にする
src_num = (src_num + 1) & ~1; // 偶数にする
} }
#ifdef TIMER // GPUスレッドの動作状況を調べる
printf("last2: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForSingleObject(hEnd[cpu_num2 - 1], 0) == WAIT_TIMEOUT){
#endif th_act |= 2; // GPUスレッドが動作中
} else if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
} }
prog_num += th2->size * parity_num;
// GPU と CPU がスレッドごとにパリティ・ブロックを計算する th2->size = 0;
th->buf = buf + (size_t)unit_size * src_off;
th->off = src_off; // ソース・ブロックの番号にする
th->size = src_num;
th->now = -1; // 初期値 - 1
//_mm_sfence();
//for (j = cpu_num - 1; j >= 0; j--){ // GPU から先に計算を開始する?
for (j = 0; j < cpu_num; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
} }
if (th_act == 3){ // 両方が動作中なら
//if (th_act == 1){ // CPUスレッドだけが動作中か調べる実験
//if (th_act == 2){ // GPUスレッドだけが動作中か調べる実験
// サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する // サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){ while (WaitForMultipleObjects(cpu_num2, hEnd, FALSE, UPDATE_TIME) == WAIT_TIMEOUT){
// th-now が最高値なので、計算が終わってるのは th-now + 1 - cpu_num 個となる // th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
j = th->now + 1 - cpu_num; i = th2->now;
if (j < 0) if (i < 0){
i = 0;
} else {
i *= th2->size;
}
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0; j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず) // 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + src_num * j) * 1000) / prog_base))){ if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
time_last = GetTickCount(); time_last = GetTickCount();
} }
if (th->len != 0){ // エラー発生 }
i = th->len; } while (th_act == 3);
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8); //} while (th_act == 1); // CPUスレッドの終了だけを待つ実験
err = 1; //} while (th_act == 2); // GPUスレッドの終了だけを待つ実験
goto error_end;
// どちらかのスレッドでパリティ・ブロックを計算する
if ((th_act & 1) == 0){ // CPUスレッドを優先的に開始する
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_off + src_num * 2 - 1 >= source_num){
src_num = source_num - src_off;
#ifdef TIMER
printf("CPU last: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
cpu_end += src_num;
th->buf = buf + (size_t)unit_size * src_off;
th->off = src_off; // ソース・ブロックの番号にする
th->size = src_num;
th->now = -1; // CPUスレッドの初期値 - 1
//_mm_sfence();
for (j = 0; j < cpu_num2 - 1; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
}
} else { // CPUスレッドが動作中なら、GPUスレッドを開始する
src_num = (source_num - src_off) * gpu_end / (cpu_end + gpu_end); // 残りブロック数に対する割合
if (src_num < src_max){
if (gpu_end == 0){ // 最初に負担するブロック数は CPUスレッドの 2倍まで
src_num = (source_num - src_off) / (cpu_num2 + 2);
if (src_num < src_max){
src_num = src_max;
} else if (src_num > src_max * 2){
src_num = src_max * 2;
}
} else if (gpu_end * 2 < cpu_end){ // GPU が遅い場合は最低負担量も減らす
if (gpu_end * 4 < cpu_end){
if (src_num < src_max / 4)
src_num = src_max / 4;
} else if (src_num < src_max / 2){
src_num = src_max / 2;
}
} else {
src_num = src_max; // 最低でも CPUスレッドと同じ量を担当する
}
}
if (src_num > vram_max)
src_num = vram_max;
if (src_off + src_num >= source_num){
src_num = source_num - src_off;
#ifdef TIMER
printf("GPU last 1: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
} else if (src_off + src_num + src_max > source_num){
src_num = source_num - src_off - src_max;
// src_num が 0にならないように、src_num == src_max なら上の last1 にする
if ((src_num < src_max) && (src_num + src_max <= vram_max) && (gpu_end * 2 > cpu_end)){
src_num += src_max; // GPU担当量が少なくて、余裕がある場合は、残りも全て任せる
#ifdef TIMER
printf("GPU last +: src_off = %d, src_num = %d + %d\n", src_off, src_num - src_max, src_max);
} else {
printf("GPU last 2: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
#ifdef TIMER
} else {
printf("GPU: remain = %d, src_off = %d, src_num = %d\n", source_num - src_off, src_off, src_num);
#endif
}
gpu_end += src_num;
th2->buf = buf + (size_t)unit_size * src_off;
th2->off = src_off; // ソース・ブロックの番号にする
th2->size = src_num;
th2->now = -1; // GPUスレッドの初期値 - 1
//_mm_sfence();
ResetEvent(hEnd[cpu_num2 - 1]); // リセットしておく
SetEvent(hRun[cpu_num2 - 1]); // サブ・スレッドに計算を開始させる
} }
// 経過表示 // 経過表示
prog_num += src_num * parity_num;
if (GetTickCount() - time_last >= UPDATE_TIME){ if (GetTickCount() - time_last >= UPDATE_TIME){
if (print_progress((int)((prog_num * 1000) / prog_base))){ if (th2->size == 0){
i = 0;
} else {
// th2-now がGPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
@@ -1808,12 +1876,58 @@ skip_count++;
src_off += src_num; src_off += src_num;
} }
// 全スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num2, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){
if (th2->size == 0){
i = 0;
} else {
// th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2;
goto error_end;
}
time_last = GetTickCount();
}
if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
}
prog_num += th2->size * parity_num;
}
if (th->size > 0) // CPUスレッドの計算量を加算する
prog_num += th->size * parity_num;
#ifdef TIMER #ifdef TIMER
time_start = GetTickCount(); time_start = GetTickCount();
#endif #endif
// パリティ・ブロックを書き込む // パリティ・ブロックを書き込む
work_buf = p_buf; work_buf = p_buf;
for (i = 0; i < parity_num; i++){ for (i = 0; i < parity_num; i++){
// CPUスレッドと GPUスレッドの計算結果を合わせる
galois_align_xor(g_buf + (size_t)unit_size * i, work_buf, unit_size);
// パリティ・ブロックのチェックサムを検証する // パリティ・ブロックのチェックサムを検証する
checksum16_return(work_buf, hash, io_size); checksum16_return(work_buf, hash, io_size);
if (memcmp(work_buf + io_size, hash, HASH_SIZE) != 0){ if (memcmp(work_buf + io_size, hash, HASH_SIZE) != 0){
@@ -1931,7 +2045,8 @@ if (prog_num != prog_base)
error_end: error_end:
InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する
for (j = 0; j < cpu_num; j++){ InterlockedExchange(&(th2->now), INT_MAX / 2);
for (j = 0; j < cpu_num2; j++){
if (hSub[j]){ // サブ・スレッドを終了させる if (hSub[j]){ // サブ・スレッドを終了させる
SetEvent(hRun[j]); SetEvent(hRun[j]);
WaitForSingleObject(hSub[j], INFINITE); WaitForSingleObject(hSub[j], INFINITE);
@@ -1965,24 +2080,26 @@ int encode_method5( // ソース・ブロックの一部とパリティ・ブロ
source_ctx_c *s_blk, // ソース・ブロックの情報 source_ctx_c *s_blk, // ソース・ブロックの情報
unsigned short *constant) unsigned short *constant)
{ {
unsigned char *buf = NULL, *p_buf; unsigned char *buf = NULL, *p_buf, *g_buf;
int err = 0, i, j, last_file, chunk_num; int err = 0, i, j, last_file, chunk_num;
int source_off, read_num, packet_off; int source_off, read_num, packet_off;
int cpu_num1, src_off, src_num, src_max, vram_max; int cpu_num1, src_off, src_num, src_max;
int cpu_num2, vram_max, cpu_end, gpu_end, th_act;
unsigned int unit_size, len; unsigned int unit_size, len;
unsigned int time_last, prog_read, prog_write; unsigned int time_last, prog_read, prog_write;
__int64 prog_num = 0, prog_base; __int64 prog_num = 0, prog_base;
size_t mem_size; size_t mem_size;
HANDLE hFile = NULL; HANDLE hFile = NULL;
HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU]; HANDLE hSub[MAX_CPU], hRun[MAX_CPU], hEnd[MAX_CPU];
RS_TH th[1]; RS_TH th[1], th2[1];
PHMD5 file_md_ctx, blk_md_ctx; PHMD5 file_md_ctx, blk_md_ctx;
memset(hSub, 0, sizeof(HANDLE) * MAX_CPU); memset(hSub, 0, sizeof(HANDLE) * MAX_CPU);
unit_size = (block_size + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1); // MEM_UNIT の倍数にする unit_size = (block_size + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1); // MEM_UNIT の倍数にする
// 作業バッファーを確保する // 作業バッファーを確保する
read_num = read_block_num(parity_num, 1, MEM_UNIT); // ソース・ブロックを何個読み込むか // CPU計算スレッドと GPU計算スレッドで保存先を別けるので、パリティ・ブロック分を倍確保する
read_num = read_block_num(parity_num * 2, 1, MEM_UNIT); // ソース・ブロックを何個読み込むか
if (read_num == 0){ if (read_num == 0){
#ifdef TIMER #ifdef TIMER
printf("cannot keep enough blocks, use another method\n"); printf("cannot keep enough blocks, use another method\n");
@@ -1991,7 +2108,7 @@ int encode_method5( // ソース・ブロックの一部とパリティ・ブロ
} }
//read_num = (read_num + 1) / 2 + 1; // 2分割の実験用 //read_num = (read_num + 1) / 2 + 1; // 2分割の実験用
//read_num = (read_num + 2) / 3 + 1; // 3分割の実験用 //read_num = (read_num + 2) / 3 + 1; // 3分割の実験用
mem_size = (size_t)(read_num + parity_num) * unit_size; mem_size = (size_t)(read_num + parity_num * 2) * unit_size;
buf = _aligned_malloc(mem_size, MEM_UNIT); // GPU 用の境界 buf = _aligned_malloc(mem_size, MEM_UNIT); // GPU 用の境界
if (buf == NULL){ if (buf == NULL){
printf("malloc, %Id\n", mem_size); printf("malloc, %Id\n", mem_size);
@@ -1999,40 +2116,34 @@ int encode_method5( // ソース・ブロックの一部とパリティ・ブロ
goto error_end; goto error_end;
} }
p_buf = buf + (size_t)unit_size * read_num; // パリティ・ブロックを記録する領域 p_buf = buf + (size_t)unit_size * read_num; // パリティ・ブロックを記録する領域
g_buf = p_buf + (size_t)unit_size * parity_num; // GPUスレッド用
prog_read = (parity_num + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする prog_read = (parity_num + 31) / 32; // 読み書きの経過をそれぞれ 3% ぐらいにする
prog_write = (source_num + 31) / 32; prog_write = (source_num + 31) / 32;
prog_base = (__int64)(source_num + prog_write) * parity_num + prog_read * source_num; // ブロックの合計掛け算個数 + 書き込み回数 prog_base = (__int64)(source_num + prog_write) * parity_num + prog_read * source_num; // ブロックの合計掛け算個数 + 書き込み回数
len = try_cache_blocking(unit_size); len = try_cache_blocking(unit_size);
chunk_num = (unit_size + len - 1) / len; chunk_num = (unit_size + len - 1) / len;
cpu_num1 = 0; // 読み込み中はスレッド数を減らす(シングル・スレッドの時は 0にする cpu_num1 = calc_thread_num2(parity_num, &cpu_num2); // 使用するスレッド数を調節する
i = 1;
while (i * 2 <= cpu_num){ // 1=0, 2~3=1, 4~7=2, 8~15=3, 16~31=4, 32=5
cpu_num1++;
i *= 2;
}
if (cpu_num1 > parity_num)
cpu_num1 = parity_num;
src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する src_max = cpu_cache & 0xFFFE; // CPU cache 最適化のため、同時に処理するブロック数を制限する
if ((src_max < 8) || (cpu_num <= 2)) if ((src_max < CACHE_MIN_NUM) || (src_max > CACHE_MAX_NUM))
src_max = 0x8000; // 不明または少な過ぎる場合は、制限しない src_max = CACHE_MAX_NUM; // 不明または極端な場合は、規定値にする
//cpu_num1 = 0; // 2nd encode の実験用に 1st encode を停止する
#ifdef TIMER #ifdef TIMER
printf("\n read some source blocks, and keep all parity blocks (GPU)\n"); printf("\n read some source blocks, and keep all parity blocks (GPU)\n");
printf("buffer size = %Id MB, read_num = %d, round = %d\n", mem_size >> 20, read_num, (source_num + read_num - 1) / read_num); printf("buffer size = %Id MB, read_num = %d, round = %d\n", mem_size >> 20, read_num, (source_num + read_num - 1) / read_num);
printf("cache: limit size = %d, chunk_size = %d, chunk_num = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num); printf("cache: limit size = %d, chunk_size = %d, chunk_num = %d\n", cpu_flag & 0x7FFF0000, len, chunk_num);
printf("unit_size = %d, cpu_num1 = %d, src_max = %d\n", unit_size, cpu_num1, src_max); printf("unit_size = %d, cpu_num1 = %d, cpu_num2 = %d\n", unit_size, cpu_num1, cpu_num2);
#endif #endif
// OpenCL の初期化 // OpenCL の初期化
vram_max = read_num; // 読み込める分だけにする vram_max = read_num; // 読み込める分だけにする
i = init_OpenCL(unit_size, len, &vram_max); i = init_OpenCL(unit_size, &vram_max);
if (i != 0){ if (i != 0){
if (i != 3) // GPU が見つからなかった場合はエラー表示しない if (i != 3) // GPU が見つからなかった場合はエラー表示しない
printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8); printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
i = free_OpenCL(); i = free_OpenCL();
if (i != 0) if (i != 0)
printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8); printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
OpenCL_method = 0; // GPU を使わない設定にする OpenCL_method = 0; // GPU を使えなかった印
// GPU を使わずに計算を続行する場合は以下をコメントアウト
err = -3; // CPU だけの方式に切り替える err = -3; // CPU だけの方式に切り替える
goto error_end; goto error_end;
} }
@@ -2043,10 +2154,14 @@ int encode_method5( // ソース・ブロックの一部とパリティ・ブロ
// マルチ・スレッドの準備をする // マルチ・スレッドの準備をする
th->mat = constant; th->mat = constant;
th2->mat = constant;
th->buf = p_buf; th->buf = p_buf;
th2->buf = g_buf;
th->size = unit_size; th->size = unit_size;
th2->size = unit_size;
th->len = len; // chunk size th->len = len; // chunk size
for (j = 0; j < cpu_num; j++){ // サブ・スレッドごとに th2->len = 0; // GPUのエラー通知用にする
for (j = 0; j < cpu_num2; j++){ // サブ・スレッドごとに
hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする hRun[j] = CreateEvent(NULL, FALSE, FALSE, NULL); // Auto Reset にする
if (hRun[j] == NULL){ if (hRun[j] == NULL){
print_win32_err(); print_win32_err();
@@ -2063,12 +2178,13 @@ int encode_method5( // ソース・ブロックの一部とパリティ・ブロ
goto error_end; goto error_end;
} }
// サブ・スレッドを起動する // サブ・スレッドを起動する
if (j == cpu_num2 - 1){ // 最後のスレッドを GPU 管理用にする
th2->run = hRun[j];
th2->end = hEnd[j];
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode_gpu, (LPVOID)th2, 0, NULL);
} else {
th->run = hRun[j]; th->run = hRun[j];
th->end = hEnd[j]; th->end = hEnd[j];
//_mm_sfence(); // メモリーへの書き込みを完了してからスレッドを起動する
if ((j == cpu_num - 1) && (OpenCL_method != 0)){ // 最後のスレッドを GPU 管理用にする
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode_gpu, (LPVOID)th, 0, NULL);
} else {
hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode3, (LPVOID)th, 0, NULL); hSub[j] = (HANDLE)_beginthreadex(NULL, STACK_SIZE, thread_encode3, (LPVOID)th, 0, NULL);
} }
if (hSub[j] == NULL){ if (hSub[j] == NULL){
@@ -2081,7 +2197,6 @@ int encode_method5( // ソース・ブロックの一部とパリティ・ブロ
} }
WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない) WaitForSingleObject(hEnd[j], INFINITE); // 設定終了の合図を待つ (リセットしない)
} }
th->len = 0; // GPUのエラー通知用にする
// 何回かに別けてソース・ブロックを読み込んで、パリティ・ブロックを少しずつ作成する // 何回かに別けてソース・ブロックを読み込んで、パリティ・ブロックを少しずつ作成する
time_last = GetTickCount(); time_last = GetTickCount();
@@ -2229,6 +2344,8 @@ time_start = GetTickCount();
time_read += GetTickCount() - time_start; time_read += GetTickCount() - time_start;
#endif #endif
if (source_off == 0)
memset(g_buf, 0, (size_t)unit_size * parity_num); // 待機中に GPU用の領域をゼロ埋めしておく
WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE); // サブ・スレッドの計算終了の合図を待つ
src_off += 1; // 計算を開始するソース・ブロックの番号 src_off += 1; // 計算を開始するソース・ブロックの番号
if (src_off == 0) // 1st encode しなかった場合src_off = 0は、生成ブロックをゼロ埋めする if (src_off == 0) // 1st encode しなかった場合src_off = 0は、生成ブロックをゼロ埋めする
@@ -2238,72 +2355,162 @@ time_read += GetTickCount() - time_start;
printf("partial encode = %d / %d (%d.%d%%), source_off = %d\n", src_off - source_off, read_num, j / 10, j % 10, source_off); printf("partial encode = %d / %d (%d.%d%%), source_off = %d\n", src_off - source_off, read_num, j / 10, j % 10, source_off);
#endif #endif
// GPU と CPU のどちらに最適化するかが難しい th2->size = 0; // 計算前の状態にしておく (th->size は既に 0 になってる)
cpu_end = gpu_end = 0;
src_off -= source_off; // バッファー内でのソース・ブロックの位置にする src_off -= source_off; // バッファー内でのソース・ブロックの位置にする
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_num > vram_max){ // VRAM に収まらない場合は、VRAM のサイズに応じて分割する
src_num = vram_max & ~1; // 減らして偶数にする(元が奇数なら分割数が増えるかも)
i = (read_num - src_off + src_num - 1) / src_num; // 何回に分けて処理するか
src_num = (read_num - src_off + i - 1) / i; // 一度に処理する量を平均化する
src_num = (src_num + 1) & ~1; // 増やして偶数にする
}
#ifdef TIMER #ifdef TIMER
printf("remain = %d, src_off = %d, src_num = %d\n", read_num - src_off, src_off, src_num); printf("remain = %d, src_off = %d, src_max = %d\n", read_num - src_off, src_off, src_max);
#endif #endif
while (src_off < read_num){ while (src_off < read_num){
// ソース・ブロックを何個ずつ処理するか // GPUスレッドと CPUスレッドのどちらかが待機中になるまで待つ
if (src_off + src_num > read_num){ do {
src_num = read_num - src_off; th_act = 0;
#ifdef TIMER // CPUスレッドの動作状況を調べる
printf("last1: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForMultipleObjects(cpu_num2 - 1, hEnd, TRUE, 0) == WAIT_TIMEOUT){
#endif th_act |= 1; // CPUスレッドが動作中
} else if (src_off + src_num * 2 - 1 >= read_num){ } else if (th->size > 0){ // CPUスレッドの計算量を加算する
src_num = read_num - src_off; prog_num += th->size * parity_num;
if (src_num > vram_max){ // VRAM のサイズまでにする th->size = 0;
src_num = (src_num + 1) / 2; // 半分にする
src_num = (src_num + 1) & ~1; // 偶数にする
} }
#ifdef TIMER // GPUスレッドの動作状況を調べる
printf("last2: src_off = %d, src_num = %d\n", src_off, src_num); if (WaitForSingleObject(hEnd[cpu_num2 - 1], 0) == WAIT_TIMEOUT){
#endif th_act |= 2; // GPUスレッドが動作中
} else if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
} }
prog_num += th2->size * parity_num;
// GPU と CPU がスレッドごとにパリティ・ブロックを計算する th2->size = 0;
th->buf = buf + (size_t)unit_size * src_off;
th->off = source_off + src_off; // ソース・ブロックの番号にする
th->size = src_num;
th->now = -1; // 初期値 - 1
//_mm_sfence();
for (j = 0; j < cpu_num; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
} }
if (th_act == 3){ // 両方が動作中なら
// サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する // サブ・スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){ while (WaitForMultipleObjects(cpu_num2, hEnd, FALSE, UPDATE_TIME) == WAIT_TIMEOUT){
// th-now が最高値なので、計算が終わってるのは th-now + 1 - cpu_num 個となる // th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
j = th->now + 1 - cpu_num; i = th2->now;
if (j < 0) if (i < 0){
i = 0;
} else {
i *= th2->size;
}
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0; j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず) // 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + src_num * j) * 1000) / prog_base))){ if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
time_last = GetTickCount(); time_last = GetTickCount();
} }
if (th->len != 0){ // エラー発生 }
i = th->len; } while (th_act == 3);
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1; // どちらかのスレッドでパリティ・ブロックを計算する
goto error_end; if ((th_act & 1) == 0){ // CPUスレッドを優先的に開始する
src_num = src_max; // 一度に処理するソース・ブロックの数を制限する
if (src_off + src_num * 2 - 1 >= read_num){
src_num = read_num - src_off;
#ifdef TIMER
printf("CPU last: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
cpu_end += src_num;
th->buf = buf + (size_t)unit_size * src_off;
th->off = source_off + src_off; // ソース・ブロックの番号にする
th->size = src_num;
th->now = -1; // CPUスレッドの初期値 - 1
//_mm_sfence();
for (j = 0; j < cpu_num2 - 1; j++){
ResetEvent(hEnd[j]); // リセットしておく
SetEvent(hRun[j]); // サブ・スレッドに計算を開始させる
}
} else { // CPUスレッドが動作中なら、GPUスレッドを開始する
src_num = (read_num - src_off) * gpu_end / (cpu_end + gpu_end); // 残りブロック数に対する割合
if (src_num < src_max){
if (gpu_end == 0){ // 最初に負担するブロック数は CPUスレッドの 2倍まで
src_num = (read_num - src_off) / (cpu_num2 + 2);
if (src_num < src_max){
src_num = src_max;
} else if (src_num > src_max * 2){
src_num = src_max * 2;
}
} else if (gpu_end * 2 < cpu_end){ // GPU が遅い場合は最低負担量も減らす
if (gpu_end * 4 < cpu_end){
if (src_num < src_max / 4)
src_num = src_max / 4;
} else if (src_num < src_max / 2){
src_num = src_max / 2;
}
} else {
src_num = src_max; // 最低でも CPUスレッドと同じ量を担当する
}
}
if (src_num > vram_max)
src_num = vram_max;
if (src_off + src_num >= read_num){
src_num = read_num - src_off;
#ifdef TIMER
printf("GPU last 1: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
} else if (src_off + src_num + src_max > read_num){
src_num = read_num - src_off - src_max;
if ((src_num < src_max) && (src_num + src_max <= vram_max) && (gpu_end * 2 > cpu_end)){
src_num += src_max; // GPU担当量が少なくて、余裕がある場合は、残りも全て任せる
#ifdef TIMER
printf("GPU last +: src_off = %d, src_num = %d + %d\n", src_off, src_num - src_max, src_max);
} else {
printf("GPU last 2: src_off = %d, src_num = %d\n", src_off, src_num);
#endif
}
#ifdef TIMER
} else {
printf("GPU: remain = %d, src_off = %d, src_num = %d\n", read_num - src_off, src_off, src_num);
#endif
}
gpu_end += src_num;
th2->buf = buf + (size_t)unit_size * src_off;
th2->off = source_off + src_off; // ソース・ブロックの番号にする
th2->size = src_num;
th2->now = -1; // GPUスレッドの初期値 - 1
//_mm_sfence();
ResetEvent(hEnd[cpu_num2 - 1]); // リセットしておく
SetEvent(hRun[cpu_num2 - 1]); // サブ・スレッドに計算を開始させる
} }
// 経過表示 // 経過表示
prog_num += src_num * parity_num;
if (GetTickCount() - time_last >= UPDATE_TIME){ if (GetTickCount() - time_last >= UPDATE_TIME){
if (print_progress((int)((prog_num * 1000) / prog_base))){ if (th2->size == 0){
i = 0;
} else {
// th2-now がGPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2; err = 2;
goto error_end; goto error_end;
} }
@@ -2313,6 +2520,50 @@ time_read += GetTickCount() - time_start;
src_off += src_num; src_off += src_num;
} }
// 全スレッドの計算終了の合図を UPDATE_TIME だけ待ちながら、経過表示する
while (WaitForMultipleObjects(cpu_num2, hEnd, TRUE, UPDATE_TIME) == WAIT_TIMEOUT){
if (th2->size == 0){
i = 0;
} else {
// th2-now が GPUスレッドの最高値なので、計算が終わってるのは th2-now 個となる
i = th2->now;
if (i < 0){
i = 0;
} else {
i *= th2->size;
}
}
if (th->size == 0){
j = 0;
} else {
// th-now が CPUスレッドの最高値なので、計算が終わってるのは th-now + 2 - cpu_num2 個となる
j = th->now + 2 - cpu_num2;
if (j < 0){
j = 0;
} else {
j /= chunk_num; // chunk数で割ってブロック数にする
j *= th->size;
}
}
// 経過表示UPDATE_TIME 時間待った場合なので、必ず経過してるはず)
if (print_progress((int)(((prog_num + i + j) * 1000) / prog_base))){
err = 2;
goto error_end;
}
time_last = GetTickCount();
}
if (th2->size > 0){ // GPUスレッドの計算量を加算する
if (th2->len != 0){ // エラー発生
i = th2->len;
printf("error, gpu-thread, %d, %d\n", i & 0xFF, i >> 8);
err = 1;
goto error_end;
}
prog_num += th2->size * parity_num;
}
if (th->size > 0) // CPUスレッドの計算量を加算する
prog_num += th->size * parity_num;
source_off += read_num; source_off += read_num;
} }
@@ -2322,7 +2573,7 @@ time_start = GetTickCount();
memcpy(common_buf + common_size, common_buf, common_size); // 後の半分に前半のをコピーする memcpy(common_buf + common_size, common_buf, common_size); // 後の半分に前半のをコピーする
// 最後にパリティ・ブロックのチェックサムを検証して、リカバリ・ファイルに書き込む // 最後にパリティ・ブロックのチェックサムを検証して、リカバリ・ファイルに書き込む
err = create_recovery_file_1pass(file_path, recovery_path, packet_limit, block_distri, err = create_recovery_file_1pass(file_path, recovery_path, packet_limit, block_distri,
packet_num, common_buf, common_size, footer_buf, footer_size, rcv_hFile, p_buf, unit_size); packet_num, common_buf, common_size, footer_buf, footer_size, rcv_hFile, p_buf, g_buf, unit_size);
#ifdef TIMER #ifdef TIMER
time_write = GetTickCount() - time_start; time_write = GetTickCount() - time_start;
#endif #endif
@@ -2337,7 +2588,8 @@ if (prog_num != prog_base - prog_write * parity_num)
error_end: error_end:
InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する InterlockedExchange(&(th->now), INT_MAX / 2); // サブ・スレッドの計算を中断する
for (j = 0; j < cpu_num; j++){ InterlockedExchange(&(th2->now), INT_MAX / 2);
for (j = 0; j < cpu_num2; j++){
if (hSub[j]){ // サブ・スレッドを終了させる if (hSub[j]){ // サブ・スレッドを終了させる
SetEvent(hRun[j]); SetEvent(hRun[j]);
WaitForSingleObject(hSub[j], INFINITE); WaitForSingleObject(hSub[j], INFINITE);

View File

@@ -18,9 +18,7 @@ __kernel void method1(
__global uint *src, __global uint *src,
__global uint *dst, __global uint *dst,
__global ushort *factors, __global ushort *factors,
int blk_num, int blk_num)
int offset,
int length)
{ {
__local uint mtab[512]; __local uint mtab[512];
int i, blk; int i, blk;
@@ -29,15 +27,14 @@ __kernel void method1(
const int work_size = get_global_size(0); const int work_size = get_global_size(0);
const int table_id = get_local_id(0); const int table_id = get_local_id(0);
src += offset; for (i = work_id; i < BLK_SIZE; i += work_size)
for (i = work_id; i < length; i += work_size)
dst[i] = 0; dst[i] = 0;
for (blk = 0; blk < blk_num; blk++){ for (blk = 0; blk < blk_num; blk++){
calc_table(mtab, table_id, factors[blk]); calc_table(mtab, table_id, factors[blk]);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for (i = work_id; i < length; i += work_size){ for (i = work_id; i < BLK_SIZE; i += work_size){
v = src[i]; v = src[i];
sum = mtab[(uchar)(v >> 16)] ^ mtab[256 + (v >> 24)]; sum = mtab[(uchar)(v >> 16)] ^ mtab[256 + (v >> 24)];
sum <<= 16; sum <<= 16;
@@ -53,9 +50,7 @@ __kernel void method2(
__global uint *src, __global uint *src,
__global uint *dst, __global uint *dst,
__global ushort *factors, __global ushort *factors,
int blk_num, int blk_num)
int offset,
int length)
{ {
__local uint mtab[512]; __local uint mtab[512];
int i, blk, pos; int i, blk, pos;
@@ -64,8 +59,7 @@ __kernel void method2(
const int work_size = get_global_size(0) * 2; const int work_size = get_global_size(0) * 2;
const int table_id = get_local_id(0); const int table_id = get_local_id(0);
src += offset; for (i = work_id; i < BLK_SIZE; i += work_size){
for (i = work_id; i < length; i += work_size){
dst[i ] = 0; dst[i ] = 0;
dst[i + 1] = 0; dst[i + 1] = 0;
} }
@@ -74,7 +68,7 @@ __kernel void method2(
calc_table(mtab, table_id, factors[blk]); calc_table(mtab, table_id, factors[blk]);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for (i = work_id; i < length; i += work_size){ for (i = work_id; i < BLK_SIZE; i += work_size){
pos = (i & ~7) + ((i & 7) >> 1); pos = (i & ~7) + ((i & 7) >> 1);
lo = src[pos ]; lo = src[pos ];
hi = src[pos + 4]; hi = src[pos + 4];
@@ -96,9 +90,7 @@ __kernel void method4(
__global uint *src, __global uint *src,
__global uint *dst, __global uint *dst,
__global ushort *factors, __global ushort *factors,
int blk_num, int blk_num)
int offset,
int length)
{ {
__local int table[16]; __local int table[16];
__local uint cache[256]; __local uint cache[256];
@@ -107,8 +99,7 @@ __kernel void method4(
const int work_id = get_global_id(0); const int work_id = get_global_id(0);
const int work_size = get_global_size(0); const int work_size = get_global_size(0);
src += offset; for (i = work_id; i < BLK_SIZE; i += work_size)
for (i = work_id; i < length; i += work_size)
dst[i] = 0; dst[i] = 0;
for (blk = 0; blk < blk_num; blk++){ for (blk = 0; blk < blk_num; blk++){
@@ -122,7 +113,7 @@ __kernel void method4(
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for (i = work_id; i < length; i += work_size){ for (i = work_id; i < BLK_SIZE; i += work_size){
pos = i & 255; pos = i & 255;
cache[pos] = src[i]; cache[pos] = src[i];
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);

View File

@@ -1,2 +1,2 @@
#define FILE_VERSION "1.3.3.0" // ファイルのバージョン番号 #define FILE_VERSION "1.3.3.1" // ファイルのバージョン番号
#define PRODUCT_VERSION "1.3.3" // 製品のバージョン番号 #define PRODUCT_VERSION "1.3.3" // 製品のバージョン番号