Release note of version 1.3.3.1

README.md

@@ -1,6 +1,6 @@
 # MultiPar
 
-### v1.3.3.0 is public
+### v1.3.3.1 is public
 
   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. 
@@ -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. 
 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
 - 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
-  - Matrix inversion may use more threads.
+  - GPU acceleration would become faster.
-  - L3 cache optimization was improved for recent CPUs.
 
 
 [ 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, 
-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 
 [GitHub](https://github.com/Yutaka-Sawada/MultiPar/releases) or 

alpha/ReadMe_English.txt

@@ -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 ]
 
- 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.
 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".
 You may need to re-start OS after overwrite install or uninstall rarely.
 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,
 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 ]
 
- 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.
 
  You can create short-cut icon or send-to link at Option window later.

alpha/ReadMe_Japanese.txt

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

alpha/Update_English.txt

@@ -1,5 +1,19 @@
 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)
 
 GUI update

alpha/Update_Japanese.txt

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

alpha/help/0409/index.htm

@@ -51,7 +51,7 @@ There are command-line manuals in "<tt>help</tt>" folder.
 </table>
 
 <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>
 </html>

alpha/help/0409/readme3.htm

@@ -8,14 +8,13 @@
 
 <h3>Install or uninstall with installer package</h3>
 <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. 
 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". 
 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, 
-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>&nbsp
 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>
 <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.
 </p>
 <p>&nbsp

alpha/help/0411/index.htm

@@ -51,7 +51,7 @@
 </table>
 
 <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>
 </html>

alpha/help/0411/readme3.htm

@@ -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>
 <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>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>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>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>&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
@@ -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>
 <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><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

alpha/help/0804/index.htm

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

alpha/help/0804/readme3.htm

@@ -8,7 +8,7 @@
 
 <h3>使用安装包安装或卸载</h3>
 <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>&nbsp
 您可以通过Windows操作系统的控制面板卸载程序，或双击MultiPar安装文件夹中的<tt>unins000.exe</tt>。由于卸载程序不会删除设置文件或安装后新添加的文件，因此您可以自行删除它们。
@@ -26,7 +26,7 @@
 
 <h3>使用压缩包安装</h3>
 <p>&nbsp
-在文件夹中解压压缩文件（<tt>MultiPar131.zip</tt>或类似名称文件）。
+在文件夹中解压压缩文件（<tt>MultiPar133.zip</tt>或类似名称文件）。
 <tt>MultiPar.exe</tt>是MultiPar的启动程序。
 </p>
 <p>&nbsp

source/par2j/Command_par2j.txt

@@ -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.
 
@@ -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.
 254: It uses one less threads than 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;
-+1024 to disable CLMUL (and use old SSSE3 code), 
++1024 to disable CLMUL (and use slower SSSE3 code), 
 +2048 to disable JIT (for SSE2), 
 +4096 to disable SSSE3,
 +8192 to disable AVX2,

source/par2j/common2.c

@@ -1,5 +1,5 @@
 // common2.c
-// Copyright : 2023-09-23 Yutaka Sawada
+// Copyright : 2023-10-13 Yutaka Sawada
 // License : GPL
 
 #ifndef _UNICODE
@@ -1848,7 +1848,7 @@ int sqrt32(int num)
 /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
 
 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 サイズから計算した制限サイズ
 unsigned int cpu_flag = 0;
 unsigned int cpu_cache = 0;	// 上位 16-bit = L3 cache の制限サイズ, 下位 16-bit = 同時処理数

source/par2j/create.c

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

source/par2j/create.h

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

source/par2j/gf16.c

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

source/par2j/lib_opencl.c

@@ -1,5 +1,5 @@
 // lib_opencl.c
-// Copyright : 2023-09-23 Yutaka Sawada
+// Copyright : 2023-10-22 Yutaka Sawada
 // License : GPL
 
 #ifndef _WIN32_WINNT
@@ -102,7 +102,6 @@ API_clEnqueueNDRangeKernel gfn_clEnqueueNDRangeKernel;
 入力
 OpenCL_method : どのデバイスを選ぶか
 unit_size : ブロックの単位サイズ
-chunk_size: 分割された断片サイズ
 src_max : ソース・ブロック個数
 
 出力
@@ -112,7 +111,7 @@ OpenCL_method : 動作フラグいろいろ
 */
 
 // 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;
 	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);
 #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);
 			if (ret != CL_SUCCESS)
 				continue;
 #ifdef DEBUG_OUTPUT
-			printf("ADDRESS_BITS = %d\n", param_value);
 			printf("MAX_MEM_ALLOC_SIZE = %I64d MB\n", param_value8 >> 20);
 #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);
 			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 数にする
 				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);
 				if (ret == CL_SUCCESS){
 #ifdef DEBUG_OUTPUT
 					printf("GLOBAL_MEM_SIZE = %I64d MB\n", param_value8 >> 20);
 #endif
-					// 領域一個あたりのサイズは全体の 1/4 までにする
+					// 領域一個あたりのサイズは全体の 1/4 までにする（VRAMを使いすぎると不安定になる）
 					param_value8 /= 4;
 					if ((cl_ulong)alloc_max > 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;
 
 	// 計算方式を選択する
-	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 ならデータの並び替え対応版を使う
 	} else if (((cpu_flag & 128) != 0) && (sse_unit == 256)){
 		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 数が必要以上に多い場合は減らす
 	if (OpenCL_method == 2){
 		// work item 一個が 8バイトずつ計算する、256個なら work group ごとに 2KB 担当する
-		data_size = chunk_size / 2048;
+		data_size = unit_size / 2048;
 	} else {
 		// work item 一個が 4バイトずつ計算する、256個なら work group ごとに 1KB 担当する
-		data_size = chunk_size / 1024;
+		data_size = unit_size / 1024;
 	}
 	if (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);
 #endif
 
-	// 出力先はchunk 1個分だけあればいい
+	// 出力先は1ブロック分だけあればいい
 	// 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);
 	if (ret != CL_SUCCESS)
 		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
 	unsigned short *mat,	// Matrix of numbers to multiply by
 	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;
 	size_t global_size, local_size;
@@ -720,14 +707,6 @@ int gpu_multiply_chunks(
 	ret = gfn_clSetKernelArg(OpenCL_kernel, 3, sizeof(int), &src_num);
 	if (ret != CL_SUCCESS)
 		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に固定する
@@ -738,18 +717,18 @@ int gpu_multiply_chunks(
 		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)
 		return (ret << 8) | 12;
 
 	// 8バイトごとに XOR する (SSE2 で XOR しても速くならず)
 	src = vram;
 	dst = (unsigned __int64 *)buf;
-	while (length > 0){
+	while (len > 0){
 		*dst ^= *src;
 		dst++;
 		src++;
-		length -= 2;
+		len -= 8;
 	}
 
 	// ホスト側でデータを変更しなくても、clEnqueueMapBufferと対で呼び出さないといけない

source/par2j/lib_opencl.h

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

source/par2j/list.c

@@ -1,5 +1,5 @@
 // list.c
-// Copyright : 2022-10-14 Yutaka Sawada
+// Copyright : 2023-10-15 Yutaka Sawada
 // License : GPL
 
 #ifndef _UNICODE
@@ -348,7 +348,7 @@ if (time_start > 0){
 // SSD 上で複数ファイルを同時に検査する
 
 // MAX_MULTI_READ の２倍ぐらいにする？
-#define MAX_READ_NUM 10
+#define MAX_READ_NUM 12
 
 int check_file_complete_multi(
 	char *ascii_buf,
@@ -370,11 +370,9 @@ unsigned int time_start = GetTickCount();
 	memset(hSub, 0, sizeof(HANDLE) * MAX_READ_NUM);
 	// Core数に応じてスレッド数を増やす
 	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
 		multi_read = 2;
 	}

source/par2j/md5_crc.c

@@ -1,5 +1,5 @@
 // md5_crc.c
-// Copyright : 2023-08-28 Yutaka Sawada
+// Copyright : 2023-10-29 Yutaka Sawada
 // License : GPL
 
 #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;
 #endif
 
+#define MAX_BUF_SIZE	2097152	// ヒープ領域を使う場合の最大サイズ
+
 // ファイルのハッシュ値と各スライスのチェックサムを同時に計算する
 int file_hash_crc(
 	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))
 			break;
 	}
@@ -866,7 +868,7 @@ DWORD WINAPI file_hash_crc2(LPVOID lpParameter)
 
 	// バッファー・サイズが大きいのでヒープ領域を使う
 	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))
 			break;
 		prog_tick++;
@@ -1303,7 +1305,7 @@ DWORD WINAPI file_hash_background(LPVOID lpParameter)
 	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))
 			break;
 	}

source/par2j/par2.c

@@ -1,5 +1,5 @@
 // par2.c
-// Copyright : 2023-09-21 Yutaka Sawada
+// Copyright : 2023-10-15 Yutaka Sawada
 // License : GPL
 
 #ifndef _UNICODE
@@ -181,7 +181,7 @@ int par2_create(
 			}
 		} 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);
 			} else {
 				common_size = set_common_packet(common_buf, &packet_num, (switch_p & 2) >> 1, files);
@@ -529,7 +529,7 @@ int par2_verify(
 
 	// ソース・ファイルが完全かどうかを調べる
 	// ファイルの状態は 完全、消失、追加、破損(完全なブロックの数) の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);
 	} else {
 		err = check_file_complete(ascii_buf, uni_buf, files, s_blk);
@@ -741,7 +741,7 @@ int par2_repair(
 
 	// ソース・ファイルが完全かどうかを一覧表示する
 	// ファイルの状態は 完全、消失、追加、破損(完全なブロックの数) の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);
 	} else {
 		err = check_file_complete(ascii_buf, uni_buf, files, s_blk);

source/par2j/par2_cmd.c

@@ -1,5 +1,5 @@
 // par2_cmd.c
-// Copyright : 2023-09-28 Yutaka Sawada
+// Copyright : 2023-10-15 Yutaka Sawada
 // License : GPL
 
 #ifndef _UNICODE
@@ -87,35 +87,37 @@ static void print_environment(void)
 	printf("CPU thread\t: %d / %d\n", cpu_num & 0xFFFF, cpu_num >> 24);
 	cpu_num &= 0xFFFF;	// 利用するコア数だけにしておく
 	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:");
-	if (cpu_flag & 1){
+	if (((cpu_flag & 16) != 0) && ((cpu_flag & 256) == 0)){
+		printf(" AVX2");
+	} else if (cpu_flag & 1){
 		if (cpu_flag & 256){
-			printf(" SSSE3(old)");
+			printf(" SSSE3(slow)");
 		} else {
 			printf(" SSSE3");
 		}
-	} else if (cpu_flag & 128){
+	} else if (((cpu_flag & 128) != 0) && ((cpu_flag & 256) == 0)){
 		printf(" SSE2");
 	} else {
 		printf(" MMX");
 	}
-#else	// 64-bit 版は SSE2, SSSE3 を表示する
+#else	// 64-bit 版は SSE2, SSSE3, AVX2 を表示する
 	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){
-			printf(" SSSE3(old)");
+			printf(" SSSE3(slow)");
 		} else {
 			printf(" SSSE3");
 		}
-	} else if (cpu_flag & 128){
+	} else if (((cpu_flag & 128) != 0) && ((cpu_flag & 256) == 0)){
 		printf(" SSE2");
 	}
 #endif
 	if (cpu_flag & 8)
 		printf(" CLMUL");
-	if (cpu_flag & 16)
-		printf(" AVX2");
 	printf("\nMemory usage\t: ");
 	if (memory_use & 7){
 		printf("%d/8", memory_use & 7);
@@ -1486,8 +1488,8 @@ ri= switch_set & 0x00040000
 				} else if (k & 512){
 					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) を使わない
 					cpu_flag &= 0xFFFFFF7F;
 				if (k & 4096)	// SSSE3 を使わない
@@ -1506,10 +1508,10 @@ ri= switch_set & 0x00040000
 					} else if (k == 254){	// 物理コア数より減らす
 						k = ((cpu_num & 0x00FF0000) >> 16) - 1;
 					} 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%になるけど、速くはならない・・・
+						// k = (k & 0xFF) + ((k >> 8) - (k & 0xFF)) / 4;	// 物理コア数の 5/4 にする？
 					}
 					if (k > MAX_CPU){
 						k = MAX_CPU;

source/par2j/reedsolomon.c

@@ -1,5 +1,5 @@
 // reedsolomon.c
-// Copyright : 2023-09-28 Yutaka Sawada
+// Copyright : 2023-10-26 Yutaka Sawada
 // License : GPL
 
 #ifndef _UNICODE
@@ -204,6 +204,48 @@ int read_block_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 で必ず２スレッド使う
+	num2 = cpu_num;
+	if (num2 < 2)
+		num2 = 2;
+	*cpu_num2 = num2;
+
+	return num1;
+}
+
 /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
 // 戸川 隼人 の「演習と応用FORTRAN77」の逆行列の計算方法を参考にして
 // Gaussian Elimination を少し修正して行列の数を一つにしてみた
@@ -960,8 +1002,7 @@ time_matrix = GetTickCount() - time_matrix;
 		// ブロック数が多いなら、ブロックごとにスレッドを割り当てる (GPU を使う)
 		if (memory_use & 16){
 			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 方式を使う
 		} else {
 			err = -4;	// メモリー不足なら Read all 方式でブロックを断片化させる
@@ -970,8 +1011,7 @@ time_matrix = GetTickCount() - time_matrix;
 		// ソース・ブロックを全て断片的に読み込むか、いくつかを丸ごと読み込むかを決める
 		if (memory_use & 16){
 			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 方式を使う
 		} else {
 			err = -2;	// メモリー不足なら Read all 方式でブロックを断片化させる

source/par2j/reedsolomon.h

@@ -17,6 +17,10 @@ extern "C" {
 #define READ_MIN_RATE	1	// 保持するブロック数の 1/2 = 50%
 #define READ_MIN_NUM	16
 
+// CPU cache 最適化のため、同時に処理するブロック数を制限する
+#define CACHE_MIN_NUM	8
+#define CACHE_MAX_NUM	128
+
 /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
 
 // Cache Blocking を試みる
@@ -35,6 +39,12 @@ int read_block_num(
 	size_t trial_alloc,		// 確保できるか確認するのか
 	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);
+
 /* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */
 
 // リード・ソロモン符号を使ってエンコードする

source/par2j/res_par2j.rc

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

source/par2j/rs_decode.c

@@ -1,5 +1,5 @@
 // rs_decode.c
-// Copyright : 2023-09-21 Yutaka Sawada
+// Copyright : 2023-10-29 Yutaka Sawada
 // License : GPL
 
 #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 では呼ばれない)
 static DWORD WINAPI thread_decode_gpu(LPVOID lpParameter)
 {
-	unsigned char *s_buf, *p_buf;
+	unsigned char *s_buf, *g_buf;
 	unsigned short *factor;
-	int i, j, block_lost, max_num, chunk_num;
+	int i, j, block_lost;
 	int src_num;
-	unsigned int unit_size, len, off, chunk_size;
+	unsigned int unit_size;
 	HANDLE hRun, hEnd;
 	RS_TH *th;
 #ifdef TIMER
@@ -296,18 +296,14 @@ unsigned int time_start2, time_encode2 = 0, loop_count2 = 0;
 #endif
 
 	th = (RS_TH *)lpParameter;
-	p_buf = th->buf;
+	g_buf = th->buf;
 	unit_size = th->size;
-	chunk_size = th->len;
 	block_lost = th->count;
 	hRun = th->run;
 	hEnd = th->end;
 	//_mm_sfence();
 	SetEvent(hEnd);	// 設定完了を通知する
 
-	chunk_num = (unit_size + chunk_size - 1) / chunk_size;
-	max_num = chunk_num * block_lost;
-
 	WaitForSingleObject(hRun, INFINITE);	// 計算開始の合図を待つ
 	while (th->now < INT_MAX / 2){
 #ifdef TIMER
@@ -325,17 +321,10 @@ time_start2 = GetTickCount();
 			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){
 				th->len = i;
 				InterlockedExchange(&(th->now), INT_MAX / 3);	// サブ・スレッドの計算を中断する
@@ -359,7 +348,6 @@ time_encode2 += GetTickCount() - time_start2;
 		WaitForSingleObject(hRun, INFINITE);	// 計算開始の合図を待つ
 	}
 #ifdef TIMER
-loop_count2 /= chunk_num;	// chunk数で割ってブロック数にする
 printf("gpu-thread :\n");
 if (time_encode2 > 0){
 	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 = ((len + 2) / 3 + (sse_unit - 1)) & ~(sse_unit - 1);	// 1/3の実験用
 	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 最適化のため、同時に処理するブロック数を制限する
-	if ((src_max < 8) || (cpu_num == 1))
+	if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
 		src_max = 0x8000;	// 不明または少な過ぎる場合は、制限しない
 #ifdef TIMER
 	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;	// ブロックの合計掛け算個数 + 読み書き回数
 	len = try_cache_blocking(unit_size);
 	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 最適化のため、同時に処理するブロック数を制限する
-	if ((src_max < 8) || (cpu_num == 1))
+	if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
 		src_max = 0x8000;	// 不明または少な過ぎる場合は、制限しない
 #ifdef TIMER
 	printf("\n read some blocks, and keep all recovering blocks\n");
@@ -1364,27 +1338,29 @@ int decode_method4(	// 全てのブロックを断片的に保持する場合 (G
 	parity_ctx_r *p_blk,		// パリティ・ブロックの情報
 	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;
 	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 time_last, prog_read, prog_write;
 	__int64 file_off, prog_num = 0, prog_base;
 	HANDLE hFile = NULL;
 	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);
 	id = mat + (block_lost * source_num);	// 何番目の消失ソース・ブロックがどのパリティで代替されるか
 
 	// 作業バッファーを確保する
 	// 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 + 2) / 3 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE;	// 3分割の実験用
 	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 用の境界
 	if (buf == NULL){
 		printf("malloc, %I64d\n", file_off);
@@ -1392,42 +1368,36 @@ int decode_method4(	// 全てのブロックを断片的に保持する場合 (G
 		goto error_end;
 	}
 	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_read = (block_lost + 31) / 32;	// 読み書きの経過をそれぞれ 3% ぐらいにする
 	prog_write = (source_num + 31) / 32;
 	prog_base *= (__int64)(source_num + prog_write) * block_lost + prog_read * source_num;	// 全体の断片の個数
 	len = try_cache_blocking(unit_size);
 	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 最適化のため、同時に処理するブロック数を制限する
-	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
 	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("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
 
 	// OpenCL の初期化
 	vram_max = source_num;
-	i = init_OpenCL(unit_size, len, &vram_max);
+	i = init_OpenCL(unit_size, &vram_max);
 	if (i != 0){
 		if (i != 3)	// GPU が見つからなかった場合はエラー表示しない
 			printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
 		i = free_OpenCL();
 		if (i != 0)
 			printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
-		OpenCL_method = 0;	// GPU を使わない設定にする
+		OpenCL_method = 0;	// GPU を使えなかった印
-		// GPU を使わずに計算を続行する場合は以下をコメントアウト
 		err = -2;	// CPU だけの方式に切り替える
 		goto error_end;
 	}
@@ -1437,10 +1407,14 @@ int decode_method4(	// 全てのブロックを断片的に保持する場合 (G
 
 	// マルチ・スレッドの準備をする
 	th->buf = p_buf;
+	th2->buf = g_buf;
 	th->size = unit_size;
+	th2->size = unit_size;
 	th->count = block_lost;
+	th2->count = block_lost;
 	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 にする
 		if (hRun[j] == NULL){
 			print_win32_err();
@@ -1457,12 +1431,13 @@ int decode_method4(	// 全てのブロックを断片的に保持する場合 (G
 			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->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);
 		}
 		if (hSub[j] == NULL){
@@ -1475,7 +1450,6 @@ int decode_method4(	// 全てのブロックを断片的に保持する場合 (G
 		}
 		WaitForSingleObject(hEnd[j], INFINITE);	// 設定終了の合図を待つ (リセットしない)
 	}
-	th->len = 0;	// GPUのエラー通知用にする
 
 	// ブロック断片を読み込んで、消失ブロック断片を復元する
 	print_progress_text(0, "Recovering slice");
@@ -1629,6 +1603,7 @@ skip_count++;
 time_read += GetTickCount() - time_start;
 #endif
 
+		memset(g_buf, 0, (size_t)unit_size * block_lost);	// 待機中に GPU用の領域をゼロ埋めしておく
 		WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE);	// サブ・スレッドの計算終了の合図を待つ
 		src_off += 1;	// 計算を開始するソース・ブロックの番号
 		if (src_off > 0){	// 計算不要なソース・ブロックはとばす
@@ -1647,74 +1622,164 @@ skip_count++;
 		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);
 #endif
+
 		recv_now = -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
-		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
 		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;
+						} else {
 							j /= chunk_num;	// chunk数で割ってブロック数にする
+							j *= th->size;
+						}
 						// 経過表示（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;
 							goto error_end;
 						}
 						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 (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;
 					goto error_end;
 				}
@@ -1724,6 +1789,50 @@ skip_count++;
 			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
 time_start = GetTickCount();
 #endif
@@ -1738,6 +1847,8 @@ time_start = GetTickCount();
 			}
 			//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);
 			if (memcmp(work_buf + io_size, hash, HASH_SIZE) != 0){
@@ -1817,7 +1928,8 @@ if (prog_num != prog_base)
 
 error_end:
 	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]){	// サブ・スレッドを終了させる
 			SetEvent(hRun[j]);
 			WaitForSingleObject(hSub[j], INFINITE);
@@ -1843,31 +1955,33 @@ int decode_method5(	// 復元するブロックだけ保持する場合 (GPU対
 	parity_ctx_r *p_blk,	// パリティ・ブロックの情報
 	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;
 	int err = 0, i, j, last_file, chunk_num, recv_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 time_last, prog_read, prog_write;
 	__int64 file_off, prog_num = 0, prog_base;
 	HANDLE hFile = NULL;
 	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);
 	id = mat + (block_lost * source_num);	// 何番目の消失ソース・ブロックがどのパリティで代替されるか
 	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){
 		//printf("cannot keep enough blocks, use another method\n");
 		return -4;	// スライスを分割して処理しないと無理
 	}
 	//read_num = (read_num + 1) / 2 + 1;	// 2分割の実験用
 	//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 用の境界
 	if (buf == NULL){
 		printf("malloc, %I64d\n", file_off);
@@ -1875,41 +1989,35 @@ int decode_method5(	// 復元するブロックだけ保持する場合 (GPU対
 		goto error_end;
 	}
 	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_write = (source_num + 31) / 32;
 	prog_base = (__int64)(source_num + prog_write) * block_lost + prog_read * source_num;	// ブロックの合計掛け算個数 + 書き込み回数
 	len = try_cache_blocking(unit_size);
 	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 最適化のため、同時に処理するブロック数を制限する
-	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
 	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("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
 
 	// OpenCL の初期化
 	vram_max = read_num;	// 読み込める分だけにする
-	i = init_OpenCL(unit_size, len, &vram_max);
+	i = init_OpenCL(unit_size, &vram_max);
 	if (i != 0){
 		if (i != 3)	// GPU が見つからなかった場合はエラー表示しない
 			printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
 		i = free_OpenCL();
 		if (i != 0)
 			printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
-		OpenCL_method = 0;	// GPU を使わない設定にする
+		OpenCL_method = 0;	// GPU を使えなかった印
-		// GPU を使わずに計算を続行する場合は以下をコメントアウト
 		err = -3;	// CPU だけの方式に切り替える
 		goto error_end;
 	}
@@ -1919,10 +2027,14 @@ int decode_method5(	// 復元するブロックだけ保持する場合 (GPU対
 
 	// マルチ・スレッドの準備をする
 	th->buf = p_buf;
+	th2->buf = g_buf;
 	th->size = unit_size;
+	th2->size = unit_size;
 	th->count = block_lost;
+	th2->count = block_lost;
 	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 にする
 		if (hRun[j] == NULL){
 			print_win32_err();
@@ -1939,12 +2051,13 @@ int decode_method5(	// 復元するブロックだけ保持する場合 (GPU対
 			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->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);
 		}
 		if (hSub[j] == NULL){
@@ -1957,7 +2070,6 @@ int decode_method5(	// 復元するブロックだけ保持する場合 (GPU対
 		}
 		WaitForSingleObject(hEnd[j], INFINITE);	// 設定終了の合図を待つ (リセットしない)
 	}
-	th->len = 0;	// GPUのエラー通知用にする
 
 	// 何回かに別けてブロックを読み込んで、消失ブロックを少しずつ復元する
 	print_progress_text(0, "Recovering slice");
@@ -2086,6 +2198,8 @@ read_count++;
 time_read += GetTickCount() - time_start;
 #endif
 
+		if (source_off == 0)
+			memset(g_buf, 0, (size_t)unit_size * block_lost);	// 待機中に GPU用の領域をゼロ埋めしておく
 		WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE);	// サブ・スレッドの計算終了の合図を待つ
 		src_off += 1;	// 計算を開始するソース・ブロックの番号
 		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;
 		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
+
 		recv_now = -1;	// 消失ブロックの本来のソース番号
 		last_file = -1;
-
+		th2->size = 0;	// 計算前の状態にしておく (th->size は既に 0 になってる)
-		// GPU と CPU のどちらに最適化するかが難しい
+		cpu_end = gpu_end = 0;
 		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
-		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
 		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;
+						} else {
 							j /= chunk_num;	// chunk数で割ってブロック数にする
+							j *= th->size;
+						}
 						// 経過表示（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;
 							goto error_end;
 						}
 						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 (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;
 					goto error_end;
 				}
@@ -2172,6 +2376,50 @@ time_read += GetTickCount() - time_start;
 			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;
 	}
 
@@ -2189,6 +2437,8 @@ time_start = GetTickCount();
 		}
 		//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);
 		if (memcmp(work_buf + unit_size - HASH_SIZE, hash, HASH_SIZE) != 0){
@@ -2252,7 +2502,8 @@ if (prog_num != prog_base)
 
 error_end:
 	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]){	// サブ・スレッドを終了させる
 			SetEvent(hRun[j]);
 			WaitForSingleObject(hSub[j], INFINITE);

source/par2j/rs_encode.c

@@ -1,5 +1,5 @@
 // rs_encode.c
-// Copyright : 2023-09-21 Yutaka Sawada
+// Copyright : 2023-10-29 Yutaka Sawada
 // License : GPL
 
 #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 では呼ばれない)
 static DWORD WINAPI thread_encode_gpu(LPVOID lpParameter)
 {
-	unsigned char *s_buf, *p_buf;
+	unsigned char *s_buf, *g_buf;
 	unsigned short *constant, *factor;
-	int i, j, max_num, chunk_num;
+	int i, j;
 	int src_off, src_num;
-	unsigned int unit_size, len, off, chunk_size;
+	unsigned int unit_size;
 	HANDLE hRun, hEnd;
 	RS_TH *th;
 #ifdef TIMER
@@ -308,17 +308,14 @@ unsigned int time_start2, time_encode2 = 0, loop_count2 = 0;
 
 	th = (RS_TH *)lpParameter;
 	constant = th->mat;
-	p_buf = th->buf;
+	g_buf = th->buf;
 	unit_size = th->size;
-	chunk_size = th->len;
 	hRun = th->run;
 	hEnd = th->end;
 	//_mm_sfence();
 	SetEvent(hEnd);	// 設定完了を通知する
 
 	factor = constant + source_num;
-	chunk_num = (unit_size + chunk_size - 1) / chunk_size;
-	max_num = chunk_num * parity_num;
 
 	WaitForSingleObject(hRun, INFINITE);	// 計算開始の合図を待つ
 	while (th->now < INT_MAX / 2){
@@ -337,27 +334,19 @@ time_start2 = GetTickCount();
 			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 は定数行列の乗数になる
 			for (i = 0; i < src_num; i++)
 				factor[i] = galois_power(constant[src_off + i], first_num + j);
 
 			// 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){
 				th->len = i;
 				InterlockedExchange(&(th->now), INT_MAX / 3);	// サブ・スレッドの計算を中断する
 				break;
 			}
-
 #ifdef TIMER
 loop_count2 += src_num;
 #endif
@@ -375,7 +364,6 @@ time_encode2 += GetTickCount() - time_start2;
 		WaitForSingleObject(hRun, INFINITE);	// 計算開始の合図を待つ
 	}
 #ifdef TIMER
-loop_count2 /= chunk_num;	// chunk数で割ってブロック数にする
 printf("gpu-thread :\n");
 if (time_encode2 > 0){
 	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 = ((len + 2) / 3 + (sse_unit - 1)) & ~(sse_unit - 1);	// 1/3の実験用
 	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 最適化のため、同時に処理するブロック数を制限する
-	if ((src_max < 8) || (cpu_num == 1))
+	if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
 		src_max = 0x8000;	// 不明または少な過ぎる場合は、制限しない
 #ifdef TIMER
 	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;	// ブロックの合計掛け算個数 + 読み書き回数
 	len = try_cache_blocking(unit_size);
 	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 最適化のため、同時に処理するブロック数を制限する
-	if ((src_max < 8) || (cpu_num == 1))
+	if ((src_max < CACHE_MIN_NUM) || (cpu_num == 1))
 		src_max = 0x8000;	// 不明または少な過ぎる場合は、制限しない
 #ifdef TIMER
 	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);	// 後の半分に前半のをコピーする
 	// 最後にパリティ・ブロックのチェックサムを検証して、リカバリ・ファイルに書き込む
 	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
 time_write = GetTickCount() - time_start;
 #endif
@@ -1463,26 +1437,28 @@ int encode_method4(	// 全てのブロックを断片的に保持する場合 (G
 	parity_ctx_c *p_blk,		// パリティ・ブロックの情報
 	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 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 time_last, prog_read, prog_write;
 	__int64 file_off, prog_num = 0, prog_base;
 	HANDLE hFile = NULL;
 	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;
 
 	memset(hSub, 0, sizeof(HANDLE) * MAX_CPU);
 
 	// 作業バッファーを確保する
 	// 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 + 2) / 3 + HASH_SIZE + (MEM_UNIT - 1)) & ~(MEM_UNIT - 1)) - HASH_SIZE;	// 3分割の実験用
 	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 用の境界
 	if (buf == NULL){
 		printf("malloc, %I64d\n", file_off);
@@ -1490,30 +1466,24 @@ int encode_method4(	// 全てのブロックを断片的に保持する場合 (G
 		goto error_end;
 	}
 	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_read = (parity_num + 31) / 32;	// 読み書きの経過をそれぞれ 3% ぐらいにする
 	prog_write = (source_num + 31) / 32;
 	prog_base *= (__int64)(source_num + prog_write) * parity_num + prog_read * source_num;	// 全体の断片の個数
 	len = try_cache_blocking(unit_size);
 	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 最適化のため、同時に処理するブロック数を制限する
-	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
 	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("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
 
 	if (io_size < block_size){	// スライスが分割される場合だけ、途中までのハッシュ値を保持する
@@ -1534,15 +1504,14 @@ int encode_method4(	// 全てのブロックを断片的に保持する場合 (G
 
 	// OpenCL の初期化
 	vram_max = source_num;
-	i = init_OpenCL(unit_size, len, &vram_max);
+	i = init_OpenCL(unit_size, &vram_max);
 	if (i != 0){
 		if (i != 3)	// GPU が見つからなかった場合はエラー表示しない
 			printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
 		i = free_OpenCL();
 		if (i != 0)
 			printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
-		OpenCL_method = 0;	// GPU を使わない設定にする
+		OpenCL_method = 0;	// GPU を使えなかった印
-		// GPU を使わずに計算を続行する場合は以下をコメントアウト
 		err = -2;	// CPU だけの方式に切り替える
 		goto error_end;
 	}
@@ -1552,10 +1521,14 @@ int encode_method4(	// 全てのブロックを断片的に保持する場合 (G
 
 	// マルチ・スレッドの準備をする
 	th->mat = constant;
+	th2->mat = constant;
 	th->buf = p_buf;
+	th2->buf = g_buf;
 	th->size = unit_size;
+	th2->size = unit_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 にする
 		if (hRun[j] == NULL){
 			print_win32_err();
@@ -1572,12 +1545,13 @@ int encode_method4(	// 全てのブロックを断片的に保持する場合 (G
 			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->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);
 		}
 		if (hSub[j] == NULL){
@@ -1590,7 +1564,6 @@ int encode_method4(	// 全てのブロックを断片的に保持する場合 (G
 		}
 		WaitForSingleObject(hEnd[j], INFINITE);	// 設定終了の合図を待つ (リセットしない)
 	}
-	th->len = 0;	// GPUのエラー通知用にする
 
 	// ソース・ブロック断片を読み込んで、パリティ・ブロック断片を作成する
 	time_last = GetTickCount();
@@ -1708,6 +1681,7 @@ skip_count++;
 time_read += GetTickCount() - time_start;
 #endif
 
+		memset(g_buf, 0, (size_t)unit_size * parity_num);	// 待機中に GPU用の領域をゼロ埋めしておく
 		WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE);	// サブ・スレッドの計算終了の合図を待つ
 		src_off += 1;	// 計算を開始するソース・ブロックの番号
 		if (src_off > 0){
@@ -1733,72 +1707,166 @@ skip_count++;
 			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
-		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
 		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 だけ待ちながら、経過表示する
-			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;
+						} else {
 							j /= chunk_num;	// chunk数で割ってブロック数にする
+							j *= th->size;
+						}
 						// 経過表示（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;
 							goto error_end;
 						}
 						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 (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;
 					goto error_end;
 				}
@@ -1808,12 +1876,58 @@ skip_count++;
 			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
 time_start = GetTickCount();
 #endif
 		// パリティ・ブロックを書き込む
 		work_buf = p_buf;
 		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);
 			if (memcmp(work_buf + io_size, hash, HASH_SIZE) != 0){
@@ -1931,7 +2045,8 @@ if (prog_num != prog_base)
 
 error_end:
 	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]){	// サブ・スレッドを終了させる
 			SetEvent(hRun[j]);
 			WaitForSingleObject(hSub[j], INFINITE);
@@ -1965,24 +2080,26 @@ int encode_method5(	// ソース・ブロックの一部とパリティ・ブロ
 	source_ctx_c *s_blk,		// ソース・ブロックの情報
 	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 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 time_last, prog_read, prog_write;
 	__int64 prog_num = 0, prog_base;
 	size_t mem_size;
 	HANDLE hFile = NULL;
 	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;
 
 	memset(hSub, 0, sizeof(HANDLE) * MAX_CPU);
 	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){
 #ifdef TIMER
 		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 + 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 用の境界
 	if (buf == NULL){
 		printf("malloc, %Id\n", mem_size);
@@ -1999,40 +2116,34 @@ int encode_method5(	// ソース・ブロックの一部とパリティ・ブロ
 		goto error_end;
 	}
 	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_write = (source_num + 31) / 32;
 	prog_base = (__int64)(source_num + prog_write) * parity_num + prog_read * source_num;	// ブロックの合計掛け算個数 + 書き込み回数
 	len = try_cache_blocking(unit_size);
 	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 最適化のため、同時に処理するブロック数を制限する
-	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
 	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("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
 
 	// OpenCL の初期化
 	vram_max = read_num;	// 読み込める分だけにする
-	i = init_OpenCL(unit_size, len, &vram_max);
+	i = init_OpenCL(unit_size, &vram_max);
 	if (i != 0){
 		if (i != 3)	// GPU が見つからなかった場合はエラー表示しない
 			printf("init_OpenCL, %d, %d\n", i & 0xFF, i >> 8);
 		i = free_OpenCL();
 		if (i != 0)
 			printf("free_OpenCL, %d, %d", i & 0xFF, i >> 8);
-		OpenCL_method = 0;	// GPU を使わない設定にする
+		OpenCL_method = 0;	// GPU を使えなかった印
-		// GPU を使わずに計算を続行する場合は以下をコメントアウト
 		err = -3;	// CPU だけの方式に切り替える
 		goto error_end;
 	}
@@ -2043,10 +2154,14 @@ int encode_method5(	// ソース・ブロックの一部とパリティ・ブロ
 
 	// マルチ・スレッドの準備をする
 	th->mat = constant;
+	th2->mat = constant;
 	th->buf = p_buf;
+	th2->buf = g_buf;
 	th->size = unit_size;
+	th2->size = unit_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 にする
 		if (hRun[j] == NULL){
 			print_win32_err();
@@ -2063,12 +2178,13 @@ int encode_method5(	// ソース・ブロックの一部とパリティ・ブロ
 			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->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);
 		}
 		if (hSub[j] == NULL){
@@ -2081,7 +2197,6 @@ int encode_method5(	// ソース・ブロックの一部とパリティ・ブロ
 		}
 		WaitForSingleObject(hEnd[j], INFINITE);	// 設定終了の合図を待つ (リセットしない)
 	}
-	th->len = 0;	// GPUのエラー通知用にする
 
 	// 何回かに別けてソース・ブロックを読み込んで、パリティ・ブロックを少しずつ作成する
 	time_last = GetTickCount();
@@ -2229,6 +2344,8 @@ time_start = GetTickCount();
 time_read += GetTickCount() - time_start;
 #endif
 
+		if (source_off == 0)
+			memset(g_buf, 0, (size_t)unit_size * parity_num);	// 待機中に GPU用の領域をゼロ埋めしておく
 		WaitForMultipleObjects(cpu_num1, hEnd, TRUE, INFINITE);	// サブ・スレッドの計算終了の合図を待つ
 		src_off += 1;	// 計算を開始するソース・ブロックの番号
 		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);
 #endif
 
-		// GPU と CPU のどちらに最適化するかが難しい
+		th2->size = 0;	// 計算前の状態にしておく (th->size は既に 0 になってる)
+		cpu_end = gpu_end = 0;
 		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
-		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
 		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 だけ待ちながら、経過表示する
-			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;
+						} else {
 							j /= chunk_num;	// chunk数で割ってブロック数にする
+							j *= th->size;
+						}
 						// 経過表示（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;
 							goto error_end;
 						}
 						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 (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;
 					goto error_end;
 				}
@@ -2313,6 +2520,50 @@ time_read += GetTickCount() - time_start;
 			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;
 	}
 
@@ -2322,7 +2573,7 @@ time_start = GetTickCount();
 	memcpy(common_buf + common_size, common_buf, common_size);	// 後の半分に前半のをコピーする
 	// 最後にパリティ・ブロックのチェックサムを検証して、リカバリ・ファイルに書き込む
 	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
 time_write = GetTickCount() - time_start;
 #endif
@@ -2337,7 +2588,8 @@ if (prog_num != prog_base - prog_write * parity_num)
 
 error_end:
 	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]){	// サブ・スレッドを終了させる
 			SetEvent(hRun[j]);
 			WaitForSingleObject(hSub[j], INFINITE);

source/par2j/source.cl

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

source/par2j/version.h

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