間欠モードだが電池を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 を使った方が
    遅くなる現象が見られました。発生条件は調査中です。

nbody を cupy で書いてみた

CUDA サンプルにもある nbody (多体問題) を cupy で解いてみた。
うまくはまれば namba を使うよりも少ない行数で記述できる。
実行速度は namba とほぼ変わらなかった。

初速は銀河面上を円運動するように与えているため、シミュレーション
時間を長くしてもあまり面白い絵になっていませんので、
適宜工夫してみて下さい。

ソースコードは以下の通り。

# -*- coding: utf-8 -*-

import datetime
import cupy as cp
import numpy as np
import matplotlib.pyplot as plt

# 質点数
N = 1024
REP = 5000
AREAX = 1e3
AREAY = 1e3
FLIP = 0

# 物理定数
cp.cuda.set_allocator(cp.cuda.MemoryPool().malloc)
MASS = cp.float32(1.0)
M = cp.ones(N, dtype=cp.float32) * MASS
G = cp.float32(1.0)
EPSILON = cp.float32(1e-6)
DT = cp.float32(0.01)

# 分布、初速を決定する
DAX = cp.zeros(N, dtype=cp.float32)
DAY = cp.zeros(N, dtype=cp.float32)

def advance(x0, y0, xx1, xx2, yy1, yy2, M, EPSILON):
        diffx = x0[xx2] - x0[xx1]
        diffy = y0[yy2] - y0[yy1]
        r = cp.sqrt(diffx * diffx + diffy * diffy)
        r3 = r * r * r + EPSILON
        dax = cp.sum(M * diffx / r3, axis = 1)
        day = cp.sum(M * diffy / r3, axis = 1)
        return dax, day

def update(x0, y0, vx0, vy0, DAX, DAY, G, DT):
        vx1 = vx0 + DAX * G * DT
        vy1 = vy0 + DAY * G * DT
        x1 = x0 + vx1 * DT
        y1 = y0 + vy1 * DT
        return x1, y1, vx1, vy1

def eccentric(x):
        return (cp.exp(2 - 8 * x) - cp.exp(2)) / (cp.exp(2 - 8) - cp.exp(2))

def init_dist():
        global AREAX, AREAY, MASS, G
        RR = eccentric(cp.random.rand(N)) + 1e-7
        THETA = cp.random.rand(N)
        XS = RR * cp.cos(THETA * 2 * cp.pi)
        YS = RR * cp.sin(THETA * 2 * cp.pi)
        VAREA = cp.ones(N, dtype=cp.float32) * cp.sqrt(G/N) / RR
        RR = cp.random.rand(N) / 4
        THETA = cp.random.rand(N)
        VSX = RR * cp.cos(THETA * 2 * cp.pi)
        VSY = RR * cp.sin(THETA * 2 * cp.pi)

        r = cp.sqrt(XS * XS + YS * YS)
        VSX = VAREA * (-YS + VSX * (1 - r))
        VSY = VAREA * (XS + VSY * (1 - r))
        XS *= AREAX
        YS *= AREAY

        # 重心を求める
        x0 = cp.sum(XS) / N
        y0 = cp.sum(YS) / N
        vx0 = cp.sum(VSX) / N
        vy0 = cp.sum(VSY) / N

        # 重心座標系に移行する
        XS -= x0
        YS -= y0
        VSX -= vx0
        VSY -= vy0
        return XS, YS, VSX, VSY

# 分布、初速を決定する
XS, YS, VSX, VSY = init_dist()
X = [XS, cp.zeros(N, dtype=cp.float32)]
Y = [YS, cp.zeros(N, dtype=cp.float32)]
VX = [VSX, cp.zeros(N, dtype=cp.float32)]
VY = [VSY, cp.zeros(N, dtype=cp.float32)]
print("intialized")

# REP 時間繰り返す
x0s = range(N)
y0s = range(N)
xx1, xx2 = np.meshgrid(x0s, x0s)
yy1, yy2 = np.meshgrid(y0s, y0s)
st = datetime.datetime.now()
for t in range(REP):
        DAX, DAY = advance(X[FLIP], Y[FLIP], xx1, xx2, yy1, yy2, M, EPSILON)
        X[1-FLIP], Y[1-FLIP], VX[1-FLIP], VY[1-FLIP] = \
                update(X[FLIP], Y[FLIP], VX[FLIP], VY[FLIP], \
                DAX, DAY, G, DT)
        if t % (REP // 50) == 0:
                print('*', end = '', flush = True)
        # 結果をグラフで表示
        if t == 0:
                plt.plot(cp.asnumpy(X[FLIP]), cp.asnumpy(Y[FLIP]), 'x', color='red')
        FLIP = 1 - FLIP

ed = datetime.datetime.now()
print("\nelapsed time:", ed - st)
x = cp.asnumpy(X[FLIP])
y = cp.asnumpy(Y[FLIP])
plt.plot(x, y, '.', color='green')
plt.show()

 

cupy のメモリアロケータ(2)

前項からの続き。

set_allocator() を実行すると、どうやらメモリキャッシュが解放されるようだ。
メモリプールを使い回す設定で2回同じテストを実行したところ、
実行速度が一番速くなった。

なお、今回は初期値を一様乱数にしてある。

numpy: 0:00:00.012221 sec
cupy: 0:00:00.000184 sec
cupy(memory pool): 0:00:00.000442 sec
cupy(memory pool2): 0:00:00.000096 sec

 

cupy のメモリアロケータ

cupy の中の人の投稿によると、cupy.cuda.set_allocator() で
メモリプールを使い回すようにすると高速になると書いてある。

実際そのとおりなのだが、投稿にある test() を2回連続で
回すと、アロケータをセットした方が圧倒的に遅くなった。
ソースコードを見ていないが、どこかに実行結果をキャッシュしている?

numpy: 0:00:00.004890 sec
cupy: 0:00:00.000179 sec
cupy(memory pool): 0:00:00.000456 sec

なお、記事では触れていないが、いわいる「バーンイン」を
行わないで GPU 回りのメソッドを呼ぶと初期化設定のため、
最初の1回だけは非常に時間がかかることに注意。

使用したテストコードは以下の通り。

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

def test(xp):
    x = xp.arange(1000000).reshape(1000, -1)
    return x.T * 2

def test_with_time(title, xp):
    t1 = datetime.datetime.now()
    test(xp)
    t2 = datetime.datetime.now()
    print(title, t2 - t1, "sec")

# バーンイン
#cp.arange(1)
test(cp)
 
test_with_time("numpy:", np)
test_with_time("cupy:", cp)
cp.cuda.set_allocator(cp.cuda.MemoryPool().malloc)
test_with_time("cupy(memory pool):", cp)