仮眠プログラマーのつぶやき

自分がプログラムやっていて、思いついたことをつぶやいていきます。

自作ゲームやツール、ソースなどを公開しております。
①ポンコツ自動車シュライシュラー
DOWNLOAD
②流体力学ソース付き
汚いほうDOWNLOAD
綺麗なほうDOWNLOAD
③ミニスタヲズ
DOWNLOAD
④地下鉄でGO
DOWNLOAD
⑤ババドン
DOWNLOAD
⑥圧縮拳(ツール)
DOWNLOAD
⑦複写拳
DOWNLOAD
⑧布シミュレーション
DOWNLOAD
⑨minecraft巨大電卓地形データ
DOWNLOAD
⑩フリュードランダー
デジゲー博頒布α版
DOWNLOAD
⑪パズドラルート解析GPGPU版
DOWNLOAD
⑫ゲーム「流体de月面着陸」
DOWNLOAD

【Stable Diffusion】sd-webui-train-toolsをlocalのweb UIに入れるときハマったポイント

LoRAでオリジナルの画像を学習させたい準備編

完成形
kanseikei


Stable Diffusion web UI上でLoRAでオリジナル画像を学習させる方法について紹介と補足をしたいと思います。Google Colabではなくローカル環境(自分のPC)で実行することを前提としています。

【Stable Diffusion】超簡単にLoRAを自作できる拡張機能「sd-webui-train-tools」の使い方!
すでにわかりやすいサイトがあり、この通りにやるのが簡単そうだとわかりました。しかし「sd-webui-train-tools」をインストールするまでいろいろハマりポイントがありました。
この手の操作では、自分が操作を正しく行なえているかの確認がとても重要で、気づかぬうちに間違ったことをして変なエラーを起こしてしまうものです。
この確認するという手順も含めて解説できたらとおもいます。

初期時点の環境

OS:Windows10
GPU 1番目のPCIスロット:GTX 1650 (VRAM 4GB)
GPU 2番目のPCIスロット:RTX 2080Ti (VRAM 11GB)
Pythonバージョン:python 3.9.12(Anaconda)
使用Webブラウザ:Google Chrome (ver 116)

まずはサイト通りにやってみる

ひとまずこの3点が重要なようです
  • Python 3.10.6
  • web UIのバージョンはコミット「a9fed7c364061ae6efb37f797b6b522cb3cf7aa2」
  • xformers0.0.16以外のバージョン
次点で以下のことをやらないように気をつける必要があります。
  • 「Stable Diffusion web UI」のデータが入ったフォルダを日本語(全角)や半角スペースが含まれるフォルダに保存する
  • 「Stable Diffusion web UI」のデータが入ったフォルダを別の場所に移動して使う
下記でzipファイルを解凍していろいろやりますが、このzipファイルを解凍する場所(パス)に日本語名や半角スペースが含まれてないことが望ましいです。zipファイルを解凍した直後に移動するのは問題ないとおもわれますが、いろいろ進めたあとにフォルダごと移動するのはダメのようですね。


やること1点目のPython 3.10.6については、すでに自分のパソコンにはAnacondaがインストールされていて、バージョンは3.9.12とやや古いものでした。

現在のPythonのバージョンの確認方法

コマンドプロンプトを開きます。
画面一番左下のWindowsマークを左クリックしcmdと入力、Enter

cmdpy


その後でてきた画面で「Python」と打ち込みEnter

py3912

これでバージョンがわかりました。

3.10.6がほしいので
まずはAnaconda経由でPythonをアップデートしようと
conda install python=3.10.6

をコマンドプロンプトから実行しましたがこれが遅い遅い。solving environmentが永遠におわらないのですがタスクマネージャーをみるとプロセスは稼働しているようで、結局おわるのに一晩かかりました。
実際は待ちきれずに
https://www.python.org/downloads/release/python-3106/
ここからAnacondaじゃなく普通のPythonインストーラーをダウンロードしてインストールしました。

2点目コミット「a9fed7c364061ae6efb37f797b6b522cb3cf7aa2」は上記サイトにもあるよう

https://github.com/AUTOMATIC1111/stable-diffusion-webui/tree/a9fed7c364061ae6efb37f797b6b522cb3cf7aa2

からダウンロードすれば良いです。

ちなみに3点目のxformersのバージョンに関しては、特になにかする必要はありませんでした。

web UIのインストール

さて
stable-diffusion-webui-a9fed7c364061ae6efb37f797b6b522cb3cf7aa2.zip
が手に入ったとおもうのでこれを解凍して中の
webui-user.bat
をメモ帳などで編集します。

このときwebui-user.batの中身は最初こうなっています。
memo

Pythonのパスを記入

先程の「現在のPythonのバージョンの確認方法」ですでに3.10.6だった人や、Python3.10.6インストール後パスを通している人はここの記入は不要です。


まずは
PYTHON=
のところに先程インストールしたPython 3.10.6のパスを入力する必要があります。

set PYTHON=C:\Users\9920x\AppData\Local\Programs\Python\Python310\python.exe

私の名前(9920x)は人によってかわるところです。
Python 3.10.6の場所(python.exe)がわからい人はEverythingなどのファイル検索ソフトで「python.exe」の名前を検索し、それっぽいところを見つけてください。
それらしいpython.exeをダブルクリック実行してPython 3.10.6と表示されればそれが3.10.6のpython.exeです。

--xformersを記入

おまじない--xformersをCOMMANDLINE_ARGS=のあとに記入します。
set COMMANDLINE_ARGS=--xformers

この時点でNVIDIAのGPUじゃないとダメなので、AMDのGPUしかもってない方はここは記入しないで進むか、あきらめてください。AMDのGPUでは検証してません。

webui-user.batを実行

webui-user.batをダブルクリックして実行するとインストールが始まります。

私の環境では
torch==1.13.1+cu117
のバージョンが勝手に選択されてインストールされてました。
しばらくするとエラーがでます。

cannot import name '_compare_version' from 'torchmetrics.utilities.imports'

このエラーがでたらこの記事のように対処します。
ローカルでstable diffusion web uiを起動するまでの手順メモ

pip install torchmetrics==0.11.4

このコマンドをコマンドプロンプトに貼り付けて実行します。
ここで注意ですが、これを実行するときのカレントディレクトリが
(解凍したフォルダ)\venv\Scripts
でないと多分ダメです。

torchm

画像のように、解凍したフォルダを開いてvenv→Scriptsとフォルダを開いて、画面上部の矢印のところにcmdと入力してエンターを押すと、その開いてる場所がカレントディレクトリになってコマンドプロンプトが起動します。
この「開いてる場所」が間違ってるとダメなので画像を見ながらよく確認してください。
このフォルダ内にpip.exeが入っており、これをコマンドのpip installで起動してるんですね。

web UIのPythonが3.10.6か確認する方法

ついでに、web UIがちゃんとPython3.10.6でインストールされているかも確認できます。
このフォルダ内に「pip3.10.exe」というファイルがある時点でなんとなく察しはつきますが、フォルダ内の「python.exe」を実行してください。

p3106

このようになっているはずです。ここが3.10.6以外だとweb UIがPython 3.10.6で動くことはないのでさきほどの手順をやり直してください。
(そもそもwebui-user.batを実行したときの2行目に書いてありますが)

webui-user.batを再度実行

torchmetricsのバージョンが解決したので再度webui-user.batを実行します。
in2

今度はなにやら3.97GBのデータのダウンロードが始まりました。
こういった重いデータは、そのサイズやファイル名はなんとなく覚えておいたほうがいいです。ダウンロードが途切れて後にエラーに遭遇した場合、ファイルサイズからダウンロードが100%できてなかったことに気づけることがあります。

web UIをインストールできた状態(起動した状態)

なんかごちゃごちゃ書いてありますが、「error」と書かれてなければエラーはでていません
end1

これが正しくインストールし起動できた状態です。
127.0.0.1:7860 はブラウザに貼り付けるURLです。
私の場合はChromeを使いました。

ざっくりいうと127.0.0.1は自分自身へアクセスするという意味なので、これでweb UIの画面がでなければ少なくとも「起動」に失敗してるということがわかります。

起動できてるかはGPU-Zなどでも確認できます。
gpuz

今回の場合、なぜか2番目のGPUのRTX2080Tiでちゃんと起動できてるようで、メモリが2.8GB消費されているのがわかります。起動できてなければ300-700MBくらいになっているはずです。
GPU-Zやタスクマネージャーでメモリ消費量を見れば、どのGPUで起動できているか、できてないかが一目でわかるのでおすすめです。
外付けGPUを使う場合でも、そもそもGPUを認識しているかの確認ができるのでよく使うソフトです。

sd-webui-train-toolsのインストール

【Stable Diffusion】超簡単にLoRAを自作できる拡張機能「sd-webui-train-tools」の使い方!
引き続きこの記事の
「ローカル版の場合」
を参考に進めます。


「Extensions」→「Install from URL」の「URL for extension’s git repository」に下記URLを入力しInstallをクリックします。
https://github.com/liasece/sd-webui-train-tools

さて何が起こりましたか?
Loadingというちっちゃい文字がしばらく出たあと私はこうなりました。
stderr: ERROR: Could not install packages due to an OSError: [WinError 5] ANZX\u06c2\u0702B: 'z:\\sd\\venv\\lib\\site-packages\\google\\~rotobuf\\internal\\_api_implementation.cp310-win_amd64.pyd'
Check the permissions.


[notice] A new release of pip available: 22.2.1 -> 23.2.1
[notice] To update, run: Z:\sd\venv\Scripts\python.exe -m pip install --upgrade pip

err0



err1

コマンドプロンプトではERRORがでて、ブラウザ側はインストールに成功したような文字がでています。
とりあえずいわれた通りInstalledタブからrestartします。
しかし「Train Tools」タブはでてきません。

詰んだ・・・

webui-user.batから再起動

一旦コマンドプロンプトを消して、再度webui-user.batを起動したところ

tool

なぜか現れていました。
とりあえずこれで先に進めそうです。

ファイルをドロップする

学習させたい画像ファイルを読み込ませます。
さきほどのサイトの通りに進めていきますが、私がつまづいたので一応メモしておきます。
ファイルをドロップするのところにフォルダごとドロップしてはダメです。
ファイルを複数選択しドラッグドロップです。(web UIにも書いてあるのに間違えた)

もし間違えるとこのようなエラーがでます。
  File "Z:\sd\venv\lib\site-packages\gradio\routes.py", line 337, in run_predict
    output = await app.get_blocks().process_api(
  File "Z:\sd\venv\lib\site-packages\gradio\blocks.py", line 1013, in process_api
    inputs = self.preprocess_data(fn_index, inputs, state)
  File "Z:\sd\venv\lib\site-packages\gradio\blocks.py", line 911, in preprocess_data
    processed_input.append(block.preprocess(inputs[i]))
  File "Z:\sd\venv\lib\site-packages\gradio\components.py", line 2360, in preprocess
    return [process_single_file(f) for f in x]
  File "Z:\sd\venv\lib\site-packages\gradio\components.py", line 2360, in <listcomp>
    return [process_single_file(f) for f in x]
  File "Z:\sd\venv\lib\site-packages\gradio\components.py", line 2334, in process_single_file
    file = processing_utils.decode_base64_to_file(
  File "Z:\sd\venv\lib\site-packages\gradio\processing_utils.py", line 276, in decode_base64_to_file
    data, extension = decode_base64_to_binary(encoding)
  File "Z:\sd\venv\lib\site-packages\gradio\processing_utils.py", line 263, in decode_base64_to_binary
    extension = get_extension(encoding)
  File "Z:\sd\venv\lib\site-packages\gradio\processing_utils.py", line 73, in get_extension
    encoding = encoding.replace("audio/wav", "audio/x-wav")
AttributeError: 'NoneType' object has no attribute 'replace'

また、日本語ファイル名のものでも意外とエラーは起きなかったのですが念のため注意しておく必要があります。

deepbooruはなくBLIPを使う

アニメ絵や二次絵ならdeepbooruがいいようですが、今回リアルよりな風景を学習させたかったのでBLIPにチェックを入れました。
blip

ちなみに今回学習させる画像は1280x720と横長ですが、Preprocess imagesのWidthとHeightは512,512でやって特に問題なく動きました。(学習もまともにできていた印象でした)

次にUpdate Datasetを押します。


dl1

また855MBの重いデータがダウンロードされはじめました。
ここらへんも念のため目を光らせておきます。
ご丁寧にファイルの保存先が書いてあります。

TypeError: 'NoneType' object is not subscriptable

もしダウンロード中になんらかのトラブルがあって855MB全部ダウンロードできてないとこうなります。

err2

Train Tools: dataset update error 'NoneType' object is not subscriptable Traceback (most recent call last): File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\dateset_ui.py", line 107, in on_ui_update_dataset_click preprocess.preprocess(None, origin_preload_data_path, processed_output_path, File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\tools\preprocess.py", line 17, in preprocess shared.interrogator.load() File "Z:\sd\modules\interrogate.py", line 123, in load self.blip_model = self.load_blip_model() File "Z:\sd\modules\interrogate.py", line 103, in load_blip_model blip_model = models.blip.blip_decoder(pretrained=files[0], image_size=blip_image_eval_size, vit='base', med_config=os.path.join(paths.paths["BLIP"], "configs", "med_config.json")) File "Z:\sd\repositories\BLIP\models\blip.py", line 175, in blip_decoder model,msg = load_checkpoint(model,pretrained) File "Z:\sd\repositories\BLIP\models\blip.py", line 224, in load_checkpoint state_dict = checkpoint['model'] TypeError: 'NoneType' object is not subscriptable


この場合283MBダウンロードしたところで止まってしまったようです。

dl2

再度ダウンロードすれば良いので一旦このファイルは消して、web UIを再起動しUpdate Datasetボタンを押すと今度はうまくいくはずです。

下の方の設定

Update Datasetが問題なく終了したら下のほうの設定にうつります。
基本的にサイトと同じ設定で、アニメ絵を学習させるわけではないので

・Train base model
v1-5-pruned-emaonly.safetensors(最初からついてる)

・Clip skip
2

・Save every n epochs
5

・Batch size
1

・Number of epochs
5

・Learning rate
0.0001

・Net dim
128

・Alpha
64

・Optimizer type
Lion

・Mixed precision
fp16

としました。

なおUse xformersにチェックがデフォルトで入っていますが、これはwebui-user.batで--xformersを記入してないとこの先でエラーがおきます。
Traceback (most recent call last):
  File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\train_ui.py", line 127, in on_train_begin_click
    train.train(cfg)
  File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\train.py", line 37, in train
    train_network.train(args)
  File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\sd_scripts\train_network.py", line 171, in train
    train_util.replace_unet_modules(unet, args.mem_eff_attn, args.xformers)
  File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\sd_scripts\library\train_util.py", line 1865, in replace_unet_modules
    replace_unet_cross_attn_to_xformers()
  File "Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\sd_scripts\library\train_util.py", line 1913, in replace_unet_cross_attn_to_xformers
    raise ImportError("No xformers / xformersがインストールされていないようです")
ImportError: No xformers / xformersがインストールされていないようです

設定が終わったら
「Begin Train」を押します。

コマンドプロンプトに大量の白い文字が流れていきますがerrorの文字が出てなければうまくいってます。

A matching Triton is not available, some optimizations will not be enabled.
Error caught was: No module named 'triton'

のエラーは出てしまいましたが、なぜか学習はうまく進んでいるようです。
調べると無視しても問題ないようなエラーらしく(ほんとか)、ここではGPU-Zでもながめておとなしくまちましょう。

3


Vram消費量10081MB、結構ギリです・・。GPU Loadをみると47%であり、そこそこの負荷で稼働しているのがわかります。

学習が終了するとsafetensorsファイルが生成されます。

(解凍フォルダ)\outputs\train_tools\projects\test0\versions\v1\trains\v1-5-pruned-emaonly-bs-1-ep-5-op-Lion-lr-0_0001-net-128-ap-64

このディレクトリに

test0-v1.safetensors

こんな名前のsafetensorsファイルができているのでこれを

(解凍フォルダ)\models\Lora

のフォルダ内にいれると、img2imgやtex2imgでsafetensorsファイルが読み込めるようになります。(使うにはゴミ箱ボタンの右のボタンを1回押す必要があります)
kansei2


これですきな絵を学習させることができるようになりました!

おわりの雑談

今回私はゲームのステージの作成にAIを使おうとしており、今後うまく行ったらまた記事にしようと考えています。
(参考)画像生成AIで絵柄が一貫した2Dゲームのステージを制作する

ここにたどり着くまでつまづきポイントが多くあり、エラー内容でググる人もいるとおもい、少しでも解決の糸口になればと今回記事にしました。
もとのサイトでもコメント欄がすごいことになってますしね

エラーが長い場合、エラーの先頭付近と最後付近の行に注目するとヒントが得られる場合があります。ググるときもそこらへんにアタリをつけて検索するといいと思います。
また検索時
Z:\sd\extensions\sd-webui-train-tools\liasece_sd_webui_train_tools\train_ui.py
みたいにパスを含めて検索することもあると思いますが、"Z:\sd\"までは個々の環境次第でかわってくる部分なので、そこは検索に含めないなど工夫をするとより精度のよい情報にたどり着けるかもしれません。
GoogleだけじゃなくてTwitter(X)でも有益な情報をつぶやいてる人がいるので、検索で使うのもいいとおもいます。

エラー内容から個人情報が抜かれることはあまりないとおもいますが、私の例では「9920x」の箇所のように、Windowsをセットアップするときにユーザー名をいれるところに本名をいれてしまった人は、パスの中にその本名がでてくるかもしれないので注意が必要です。
(ちなみに9920xはCPUの型番からとりました)


2023/9現在、まだこのような難しい手順をとっていく必要がありますが、この界隈は発展が速いので、そのうちもっと簡単にLora学習ができるようになると思います。


またAI学習に関してはNVIDIA GPUが他のベンダーより2歩3歩さきをいっている印象です。CUDAソフトウェアの資産は歴史が長くユーザーも多いことから一朝一夕にAMDやIntelが逆転できる状況ではないと考えています。
--xformersの指定も、Vram消費量を抑えたり高速化に寄与している点でほぼ必須です。これはNVIDIAのGPUでしか使えない機能のようです。
もしGPUをAI画像生成目的に購入しようと考えているなら間違いなくNVIDIA製GPU一択であり、求められるOSはWindows、次点でLinuxと考えます。

WindowsのノートPCしか持ってない方なら、外付けGPU Boxの選択肢もありです。実際私はAKiTiOの外付けGPU Box(約3万円)にRTX3080やRTX2080Tiを搭載してStable Diffusion web UIを動かすことに成功しています。ただノートPC側にThunderBolt対応USB-C端子がないとこれも無理な話ですが。

UnityのComputeShader関連のエラー Failed to present D3D11 swapchain due to device reset/removed.について

Failed to present D3D11 swapchain due to device reset/removed.

というエラーメッセージについて、個人的にハマったこと、解決したことをまとめようと思います。
結論から言うとこういうことです。

環境:Windows 11、Unity 2021.2.7f1、GPUはGeForce MX150(mem 2GB)
・Windowsで描画担当のGPUが高負荷で数秒応答しなくなった時、ドライバが強制終了する仕様がある
・Unityのプログラムで初期地形生成にGPUのComputeShaderを使っていた。低スペックGPUだと3-5秒くらいかかるので、タイムアウト時間によりOSがドライバをリセット、Unityで上記エラーが発生

さて、これはバグでもなくUnity側の不具合でもなく、WindowsOSの仕様が引き起こしたことなのでバグとは呼ばず「エラー」と言うこととします。
では時系列順に起こったことを書いていきます。

Unity Editorでエラー&強制終了

ある日からUnity Editorで再生ボタンを押すとこのようにエラーマークが出てEditorごと落ちるようになった。
100%落ちるわけではなく20%くらいの確率で普通にうまくプログラムが起動する。
また、別の高スペックなほうのPCでやると一切エラーマークは発生せず、ちゃんと起動する。
e1

Waiting for Unity's code to finish executing

続いてエラーが出るときにEditor Windowの後ろのほうに出ているメッセージになにかヒントがないか調べてみた。
"EnterPlayMode"や"Waiting for Unity's code to finish executing"などで検索するも有益な情報なし。
そもそも今回のエラーと関係あるのか確証もなかった。

e2

Buildしたらどうなるか

Editor上で実行するときだけ起こる問題なのかと思い、今度はBuildでexeを生成してから実行してみた。
すると100%の確率で強制終了。画像のようなエラーマークがでる。
別のPCでやるとちゃんと起動できる。なのでPCのスペックに依存したエラーかなとあたりがつく
e3

クラッシュログを見る

Editorでエラー終了するときにどうもクラッシュログが生成されていることに気づく。
クラッシュログをみるとGPUのメモリ確保(TextureやComputeBuffer)に全て失敗しているようだった。
これでGPU関連のプログラムでなにかまずいことをしているかもと絞り込んでいろいろ試行錯誤開始。
crashlog

エラーが発生しない状況を突き止める

その結果、地形生成部分のGPUの計算負荷(ComputeShaderを使っている)を軽くするとエラーが発生しないことがわかった。具体的には1616×1616マスの地形生成を208×208に減らした。
なので今後は1616×1616のうち208×208の地形生成を最初に行い、数秒後に残りの部分の地形生成を行いエラーが発生するか確かめてみることにした。

GPU Timeoutの原因がexpensive workloads

そうしたところエラー発生。ついに有益なメッセージが得られた。
Failed to present D3D11 swapchain due to device reset/removed.
This error can happen if you draw or dispatch very expensive workloads to the GPU,
which can cause Windows to detect a GPU Timeout and reset the device.
~~~~~~~~~

e4

要約するとGPU負荷がすごすぎてGPUタイムアウト、Windowsがデバイスをリセットした、ということのよう。
これはWindows上でCUDAで遊んだり機械学習してる人なら見慣れている光景かもしれないが、レジストリのTdrDelayを書き換えれば解決するやつ。

TdrDelayを書き換え解決

リンクのやつを参考にTdrDelayを60秒にして再起動をかける。

振り返り

振り返ってみると、ある日からエラーが発生するようになったのはWindows 11にアップデートしたからでした。もともとWindows10で作業していたときは別件ですでに上記のレジストリをいじっていたのでエラーは発生しなかったのだと思われます。これがOSアップデートで戻され、タイムアウト時間が数秒程度になっていたのでしょう。
Unity側からしてみると、プログラム実行中にドライバが突然リセットされるので何が起こったかわかるはずもありません。最初まともなエラーメッセージがでなかったのも無理はありません。

多分、WindowsOSがGPUドライバを強制リセットすることがあるということを知らないとなかなか解決には至らなかったかと思うので今回記事にさせていただきました。

パディングなしでバンクコンフリクトを解消する方法 CUDA 行列転置

シェアードメモリとバンク

GPUを使った行列転置でほぼ頻出話題のバンクコンフリクト(Bank Conflict)とパディング(Padding)。前提知識としてCUDA等のGPUプログラミングでシェアードメモリ(共有メモリ)を明示的に扱う際に、Bankを意識しないと遅くなっちゃいますよという話があります。
bank

シェアードメモリは4byte刻みで32個のBankに分かれています。
例えばthread 0がSMEM[0]、thread 1がSMEM[32]にアクセスするようなプログラムだと(SMEMは4byte型とします)、どちらも同じBank0にアクセスすることになるのでリプレイ(replay)が発生して遅くなります。

行列転置バンクコンフリクトありver

つまり普通にシェアードメモリを使って行列転置をしようとすると
jurai

こんな感じにShared Memoryから値を読み込む際にバンクコンフリクトが起こってしまっています。
ちなみにここではわかりやすくBankは0~3までとしています。

行列転置パディングありver

そこでパディングをいれることで参照Bankがずれるのでこれを解消できます。そこそこ有名なテクニックです。
Kepler GPUアーキテクチャとプログラム最適化
プログラムを高速化する話Ⅱ 〜GPGPU編〜

sample0

概念図
juraip

これでShared Memoryに書き込むときも読み込むときも、Bank競合が起きてないことが分かります。
でも貴重なShared Memoryに未使用領域があるのはちょっと気持ち悪くないですか。

行列転置パディングなしver

上の概念図がすでに大きなヒントでした。
sinki
これでも読み書きともバンクコンフリクトは起こってないですね。
メモリアクセスの仕方としては
code0

こんな感じにx方向のズレ幅をyごとにずらしてやればいいです。

一応コードの全貌を載せておきます。

#include<stdio.h>
#include<cuda.h>
#include<cuda_runtime.h>

// --matrix size --
const int nimax=8192,njmax=8192;
// -- the number of threads per block --
const int BS=32;
__device__ float idata[nimax][njmax],odata[nimax][njmax];

__global__ void transposeNaive(){
  int ni,nj;
  ni=blockIdx.y*blockDim.y+threadIdx.y;
  nj=blockIdx.x*blockDim.x+threadIdx.x;
  odata[nj][ni]=idata[ni][nj];
  return;
}

__global__ void transposeShared(){
  int ni,nj;
  __shared__ float tile[BS][BS];
  ni=blockIdx.y*blockDim.y+threadIdx.y;
  nj=blockIdx.x*blockDim.x+threadIdx.x;
  tile[threadIdx.y][threadIdx.x]=idata[ni][nj];
  __syncthreads();
  ni=blockIdx.x*blockDim.x+threadIdx.y;
  nj=blockIdx.y*blockDim.y+threadIdx.x;
  odata[ni][nj]=tile[threadIdx.x][threadIdx.y];
  return;
}

__global__ void transposeSharedNobankconf(){
  int ni,nj;
  __shared__ float tile[BS][BS+1];
  ni=blockIdx.y*blockDim.y+threadIdx.y;
  nj=blockIdx.x*blockDim.x+threadIdx.x;
  tile[threadIdx.y][threadIdx.x]=idata[ni][nj];
  __syncthreads();
  ni=blockIdx.x*blockDim.x+threadIdx.y;
  nj=blockIdx.y*blockDim.y+threadIdx.x;
  odata[ni][nj]=tile[threadIdx.x][threadIdx.y];
  return;
}


__global__ void transposeSharedNobankconf_nopading(){
  int ni,nj;
  __shared__ float tile[BS][BS];
  ni=blockIdx.y*blockDim.y+threadIdx.y;
  nj=blockIdx.x*blockDim.x+threadIdx.x;
  tile[threadIdx.y][(threadIdx.x+threadIdx.y)%BS]=idata[ni][nj];
  __syncthreads();
  ni=blockIdx.x*blockDim.x+threadIdx.y;
  nj=blockIdx.y*blockDim.y+threadIdx.x;
  odata[ni][nj]=tile[threadIdx.x][(threadIdx.x+threadIdx.y)%BS];
  return;
}


void transpose(float hidata[][njmax],float hodata[][njmax]){
  int ni,nj;
  for(ni=0;ni<nimax;ni++){
    for(nj=0;nj<njmax;nj++){
      hodata[nj][ni]=hidata[ni][nj];
    }
  }
  return;
}

void check_mat(float C[][njmax],float CC[][njmax]){
  int ni,nj,flg;
  flg=0;
  for(ni=0;ni<nimax;ni++){
    for(nj=0;nj<njmax;nj++){
      if(fabs(C[ni][nj]-CC[ni][nj])>1.0e-8){
        flg=1;
        printf("%d %d %lf %lf\n",ni,nj,C[ni][nj],CC[ni][nj]);
      }
    }
  }
  if(flg==1){
    printf("Calculation error.\n");
  }else{
    printf("OK.\n");
  }
  return;
}




float hidata[nimax][njmax],hodata[nimax][njmax],htdata[nimax][njmax];

int main(){
  int ni,nj;
  dim3 grid,blck;
  grid.x=nimax/BS; grid.y=njmax/BS;
  blck.x=BS; blck.y=BS;
  // -- set initial value --
  srand(248309);
  
  
  for(ni=0;ni<nimax;ni++)
  {
    for(nj=0;nj<njmax;nj++)
    {
      hidata[ni][nj]=(float)rand()*1.1;
    }
  }
  
  
  cudaMemcpyToSymbol(idata,hidata,nimax*njmax*sizeof(float));
  transpose(hidata,htdata);
  // -- 1. transpose by simple kernel --
  transposeNaive<<<grid,blck>>>();
  transposeNaive<<<grid,blck>>>();
  cudaMemcpyFromSymbol(hodata,odata,nimax*njmax*sizeof(float));
  check_mat(htdata,hodata);
  
  // -- 2. transpose with shared memory --
  transposeShared<<<grid,blck>>>();
  transposeShared<<<grid,blck>>>();
  cudaMemcpyFromSymbol(hodata,odata,nimax*njmax*sizeof(float));
  check_mat(htdata,hodata);
  
  // -- 3. transpose with shared memory and without bank conflict --
  transposeSharedNobankconf<<<grid,blck>>>();
  transposeSharedNobankconf<<<grid,blck>>>();
  cudaMemcpyFromSymbol(hodata,odata,nimax*njmax*sizeof(float));
  check_mat(htdata,hodata);

  // -- 4. transpose with shared memory and without bank conflict and nopading --
  transposeSharedNobankconf_nopading<<<grid,blck>>>();
  transposeSharedNobankconf_nopading<<<grid,blck>>>();
  cudaMemcpyFromSymbol(hodata,odata,nimax*njmax*sizeof(float));
  check_mat(htdata,hodata);
  

  cudaDeviceReset();
  
  return 0;
}


8192×8192要素の行列の転置をするコードとなっています。

速度比較


normal

1.transposeNaive:シェアードメモリすら使ってない。グローバルメモリへコアレスアクセスできてないので一番遅い
2.transposeShared:シェアードメモリ内でバンクコンフリクト発生
3.transposeNobankconf:パディング入れてるやつ
4.transposeNobankconf_nopading:パディング入れてないやつ


まずグローバルメモリにコアレスアクセスできてない1番、これは論外です。次に2,3のバンクコンフリクトのあるなし、これだけで2倍違うことがわかります。
今回考案した4番は、速度的に3番に劣ってないことがわかります。

で、これは何の役に立つんですか?

さきほど書いたよう、パディング入りは無駄なメモリ領域を確保しています。なのでOccupancyが100%になる前にシェアードメモリ容量が律速になるような問題では以下のようなことが発生します。

1Blockあたり16.25KBのシェアードメモリを使うプログラムの場合
nvvp2

1Blockあたり16.0KBのシェアードメモリを使うプログラムの場合
nvvp

このようにOccupancyが上がることが期待できます。

なお、上記の行列転置のコードのカーネルはすべてOccupancy100%です。(Shared Memoryを使い切る前にOccupancy100%になる)

CUDA Unified Memoryを扱う方法 (CとPyCUDAのコード)

最小限のコードであまり落ちてなかったので自分のメモ用としても

まずはCUDA版(nvccでコンパイルするやつ)
#include <stdio.h>
__global__ void VecAdd(float* A, float* B, float* C) {
	int id = 256 * blockIdx.x + threadIdx.x;
	C[id]=A[id]+B[id];
}

int main() {
	int N = 1024;
	// Allocate 3 arrays on GPU
	float *d_A, * d_B, * d_C;
	cudaMallocManaged(&d_A, N * sizeof(float));
	cudaMallocManaged(&d_B, N * sizeof(float));
	cudaMallocManaged(&d_C, N * sizeof(float));
	// CPU init
	for(int i=0;i<N;i++){
		d_A[i]=d_B[i]=1.0*i;
	}

	//gpu kernel
	VecAdd <<<N/256, 256>>> (d_A, d_B, d_C);
	cudaDeviceSynchronize();//wait

	for(int i = 0;i<N;i++){
		printf("%f ",d_C[i]);
	}
	
	cudaFree(d_A);
	cudaFree(d_B);
	cudaFree(d_C);
	return 0;
}

つぎにPyCUDA版(2020.1)
from pycuda.autoinit import context
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
import numpy as np

mod = SourceModule("""
__global__ void doublify(float *a)
{
    a[threadIdx.x] *= 2;
}
""")
doublify = mod.get_function("doublify")

a = cuda.managed_empty(shape=12, dtype=np.float32, mem_flags=cuda.mem_attach_flags.GLOBAL)
a[:] = np.linspace(0, 11, len(a)) # Fill array on host
doublify(a, grid=(1,1), block=(len(a),1,1))
context.synchronize() # Wait for kernel completion before host access
print(a)
del a #これでガベージコレクションに削除してもらうことで、GPUメモリも解放される

ほんの19行でかけた aにCPUからもGPUからもアクセスできている。

Xcode+Unity+NCMB+古すぎるPCでハマった。Cannot find protocol declaration for 'ASAuthorizationControllerDelegate'

なかなか厄介なエラーにぶち当たったのでメモ

Mac OS : 10.13.6(High Sierra)
Xcode : 10.1
Unity 2020.1.2f1
NCMB v4.2.0


Unityで自作アプリを作成して、Xcodeでビルドして自前のiPhoneでデバッグしようとしたが
Cannot find protocol declaration for 'ASAuthorizationControllerDelegate': did you mean 'UINavigationControllNCMBAppleAuth.m'

とかいうエラーが発生。
e1

e2

e3


これは16個あるうちの最初のエラーで他にもいろいろあるようだけど、何が根本的な原因か全然わからなかった。
ちなみにNCMBを使っているのでデバッグにもApple Developer登録(1万円払うやつ)が必要だと思って、それはすでに取得してある。


でこれの解決方法だが、「PCを新しくする」をして解決した。

ここで使用していたPCはMacBook Air 2011 Midとかなり古かった。調べるとこのモデルはOSがどれだけ新しくしてもMac OS  10.13.6(High Sierra)まで。
(2020/8時点で最新の)Mac OS 10.15 Catalinaじゃないと使えないXcodeバージョンというのがあり、それでBuildを試したら謎のエラーはなくなった。

開発環境は古すぎるとダメという教訓であった。
プロフィール

toropippi

記事検索
アクセスカウンター

    QRコード
    QRコード
    • ライブドアブログ