minosysScript を1年ぶりに更新中

実行エンジンを中途半端に実装して1年間
ほったらかしていた minosys-scrypt ですが、
1年ぶりに実装を追加することになりました。

minosys-script.git を取得するすると最近の
ソースコードが入手できます。

とりあえず、今週末までには class, package
関連以外の内容を追加したいと考えております。
毎日コミット予定なので、言語が出来ていく様子が
リアルタイムで見られるかも?!

minosys-script は最終的には www サーバ
として稼働して www に特化した処理(html の
ダウンロードや画像ファイルの range 転送)
を実装する予定ですので、乞うご期待。

何年先になるか分かりませんが。

 

フリーランス

開業届を税務署に出してきました。

知人には今年の年賀状でフリーランス宣言していましたが、
これで、名実ともにフリーランスということになりました。

当面は請負で食べていくつもりですが、技術コンサルや
製品開発のお手伝いまで間口を広げられればいいなと思っています。

皆様、よろしくお願いします。

自作PCを無償譲渡

お世話になった支援施設に自作PCを進呈することになりました。
スペックは Core i7 860 と初代 Core i7 ですが、Windows10 を
載せてみると意外に快適に動きました。

このPC、今は懐かしい東芝の SpursEngine が乗っていて、
ビデオ圧縮の速さは今なおダントツです。
(ビデオカードは nVidia GT640なのでCUDAも動きますが、
非力です)

その他、HDDが2TB×2発乗っていたりして、完全にビデオ編集向け
仕様なのですが、きっとオーバースペックだろうなあ。

TWELite が電池を消費する問題(3)

TWELite を変更して2週間経過したが、電源電圧は3.08Vと
3V台をキープしている。このまま行けば、最低1か月は
持つのではないだろうか。

変更したのは TWELite デバイスと繰り返し精度を「低」に
にしたことだが、繰り返し精度が消費電力に影響するとは
思えないので、やはり購入したデバイスが外れだったと
考えるのが妥当だろう。

TWELite が2日で電池を消費する問題(2)

PWM, ADC, OUTポート、I2C, SIO と次々に切っていっても
電池を消費してしまうため、試しに SHT31 だけを電池に
つないで放置してみた。

2時間経過しても電源電圧に変化なしだったので、
これは TWELite 側の問題だなと確信し、デバイスを別の
ものに変えてみた。

3時間経過後、電源電圧は降下していていなかったので、
おそらく予想が当たっていたのではないだろうか。

間欠モードだが電池を2日で消費

SHT31 ベースの温湿度計を動かして2日しか経たないのに
単4電池の電圧が2Vに低下して停止してしまった。

どうやらどこかで暗電流が流れているらしい。

仕方がないので、ポート, ADC, PWM のプルアップを全て停止し、
電池を単3電池に変更して実地検証している。(オプション
ビットは0x800836)

これでも電池が持たないようなら… 回路を検討し直さないと
いけないかも。

TWELite のモード7でwatchdog timeout? 発生

久しぶりにモノワイアレス社の TWELite を購入し、秋月電子通商の
SHT31キットを使って無線温湿度計を作ってみた。

モード7を60秒間欠で動かして消費電力を抑える。
単4電池2本でどのくらい持つだろうか?

ところで、ウォッチドッグと思われるタイムアウトに悩まされたのでメモ。

  1. ボタン入力、ADC、PWM といったI/O系をすべて外してはまる。
    シリアル入力から何も入力できず、5秒後にリセットされる現象が発生。
    何が原因か判然としないが、ボタン入力を復活させたら起きなくなった。
  2. モード7しか使わなくてもモード0は残すべし。
    ライターはモード0で起動するため、そうしないとEEPROMの
    書き換えが出来なくなる。
  3. PWMを使っていなければPWMタイマーは触るな。
    初期化されていないタイマーを止めると動作が停止し、5秒後にリセット
    されるようだ。vAHI_StopWatchdog() を呼び出してもこの現象は
    変わらなかった。

以前、SHT11ベースの温湿度計を作成したが、それとの差が
かなりある状態。簡易マニュアルには室温になれるのに数日要する
ようなことが書いてあるので、気長に待ちます。

numpy, cupy, numba 速度比較まとめ

1000要素のベクトルのドット積を求めるプログラムで、
numpy, cupy, numba の実行速度を比較してみました。

numba 64bit: 0:00:00.001167
numba 32bit: 0:00:00.000692
cupy 64bit: 0:00:00.000074
cupy 32bit: 0:00:00.000060
numpy 64bit: 0:00:00.000006
numpy 32bit: 0:00:00.000003

となり、numpy 圧勝。cupy が若干遅くなっているのはメイン
メモリとグラフィックメモリの相互移動が発生しているためと
推察します。
numba が著しく遅いのは和の計算に atomic add を使っている
ためでしょう。
(バイナリツリーを用いると劇的に速くなりますが、
面倒なので実装せず。)

最後にテストに用いたプログラムを載せておきます。

# -*- coding: utf-8 -*-
from numba import cuda
import numpy as np
import cupy as cp
import datetime

@cuda.jit
def dotProduct(c, a, b):
        i = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
        cuda.atomic.add(c, 0, a[i] * b[i])

N = 1000
a = np.random.rand(N).astype(np.float64)
b = np.random.rand(N).astype(np.float64)
d = np.zeros(1, dtype=np.float64)

a32 = np.random.rand(N).astype(np.float32)
b32 = np.random.rand(N).astype(np.float32)
d32 = np.zeros(1, dtype=np.float32)

ac = cp.asarray(a)
bc = cp.asarray(b)
dc = cp.asarray(d)

ac32 = cp.asarray(a32)
bc32 = cp.asarray(b32)
dc32 = cp.asarray(d32)

threadN = 256
blockN = (N + threadN - 1) // threadN
# burn-in
dotProduct[blockN, threadN](d, a, b)
dotProduct[blockN, threadN](d32, a32, b32)
dc = cp.dot(ac, bc)
dc32 = cp.dot(ac32, bc32)
d_gold = np.dot(a, b)
d32 = np.dot(a32, b32)

d = np.zeros(1, dtype=np.float64)
st = datetime.datetime.now()
dotProduct[blockN, threadN](d, a, b)
ed = datetime.datetime.now()
print("numba 64bit:", ed - st)

d32 = np.zeros(1, dtype=np.float32)
st = datetime.datetime.now()
dotProduct[blockN, threadN](d32, a32, b32)
ed = datetime.datetime.now()
print("numba 32bit:", ed - st)

st = datetime.datetime.now()
dc = cp.dot(ac, bc)
ed = datetime.datetime.now()
print("cupy 64bit:", ed - st)

st = datetime.datetime.now()
dc32 = cp.dot(ac32, bc32)
ed = datetime.datetime.now()
print("cupy 32bit:", ed - st)

st = datetime.datetime.now()
d_gold = np.dot(a, b)
ed = datetime.datetime.now()
print("numpy 64bit:", ed - st)

st = datetime.datetime.now()
d_gold32 = np.dot(a32, b32)
ed = datetime.datetime.now()
print("numpy 32bit:", ed - st)

 

CUDA の Tensor Core を使ってみた

折角 TITAN V を購入したので、SM 7.0 以上で利用できる
Tensor Core を使った行列計算テストプログラム
書いてみました。

CUDA Toolkit のサンプルとして Tensor Core を
使ったプログラムがありますが、余計な最適化処理を
していてちょっと分かりづらい。

さらに2のn乗(n >= 4)でないと動作しないコードだった
ので、任意の行列に対応するようにしました。

で、その効果ですが、やや微妙です。。。

start wmma version...
calc_fm elapsed time without TensorCore:7.1383ms
calc_fm elapsed time without TensorCore:4.08064ms
calc_fm elapsed time with TensorCore:3.6567ms
start CPU gold...
calc_fm_gold elapsed time:43.6952m

MNIST を VGG とかで解析する際の第1段めを想定した実験
ですが、効果は10%程度となりました。っていうか、Tensor Core
を使わなくても TITAN V の CUDA Core が速すぎるというか。

WMMA を使うにあたっていくつか TIPS がありました。

  1. 行列の内積を求める際は row_major を使う。
    サンプルでは matrix_a に対して col_major を指定しています
    が、row_major にしないと正しく積が計算されません。
  2. __syncthreads() を呼ばなくてもいい。
    wmma ロジック内部でスレッド同期してから返される
    ようなので、__syncthreads() 等のスレッド同期を外部で
    行う必要はありませんでした。
  3. タイル数をむやみに増やしても効果がない。
    nVidia のサンプルでは全ての部分行列(タイル)を読み込んで
    高速化を図っていますが、行列をキャッシュして使い回す以上の
    意味はなく、それ以外の積和演算のパフォーマンスへの
    インパクトがありません。
    これは wmma のアクセス関数が *_sync と同期タイプしか
    提供されていないためだと思われます。
    (正直、どうやったら 640 コア使えるのか分かりません。)
  4. 部分行列の総和は matrix_c を読み書きする。
    accumulator はload/storeが可能です。
    というか、この機能がないとwmmaの外で総和を計算
    することになり、若干不便です。(私が複数タイル用に
    書いたようなコードを埋め込む必要が出てきます。)
  5. 行列の要素数によっては効果がないことがある。
    一般に、要素数が多いほど効果は上がるのですが、
    n, m, k の関係によっては Tensor Core を使った方が
    遅くなる現象が見られました。発生条件は調査中です。