minosysScript 更新中

minosysScript 実装中です。

今日までに代入、++、– –、比較演算子、+, * を実装しました。
今週は地道に演算子の実装を行う予定です。

今週中に演算子全ての実装が終わると思いましたが、
演算子の数が思ったよりも多く、ちょっと無理そうです。

minosysScript を1年ぶりに更新中

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

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

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

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

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

 

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

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

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

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

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

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

マルチコア vs SIMD vs CUDA

前回あまりにも Numba が遅いのでおかしいと思い、マルチコア、SIMD、CUDA を比較するためのテストプログラムを作ってみた。

結果は、以下の通り。シングル CPU の非SIMD演算が最も遅く、CUDA が最も速いという予想通りの結果になりました。

version: 201307
#20 processors, 20 threads
start parallel calculation.
<<single-nosimd>>
elapsed 1.115483 seconds

<<single-simd>>
elapsed 0.535289 seconds

<<multicore-nosimd>>
elapsed 0.189586 seconds

<<multicore-simd>>
elapsed 0.128129 seconds

<<cuda>>
elapsed 0.021330 seconds

1.234571 1.581539 1.109199 1.103452 0.831745 1.106268 0.878185 1.868425 1.353009 0.748571 
1.234571 1.581539 1.109199 1.103452 0.831745 1.106268 0.878185 1.868425 1.353009 0.748571 

ただし、cudaMallocManaged() による Unified Memory を使うと速度が極端に低下する現象が見られました。(20倍くらい遅くなる)

やはり PCIe バスの遅さが実行時間に影響していると考えてよいようです。Numba のように自動的にメモリ配置してしまうライブラリは却って速度低下の原因となることがわかりました。

というわけで、CUDA を使うときはなるべく Device Memory に閉じるように設計しましょう、という当たり前の結論となりました。

この下は今回使ったソースコードです。マルチコア化には OpenMP を、SIMDには AVX2 を使用しました。

[Nelem.h]

#define NELEM (400 * 1000 * 1000) # 400M elements

[sample_cuda.h]

#if !defined(SAMPLE_CUDA_H_)
#define SAMPLE_CUDA_H_

#if defined(__cplusplus)
extern "C"
#endif
void prepare(float *a, float *b, float *c);

#if defined(__cplusplus)
extern "C"
#endif
float *getResult(float *);

#if defined(__cplusplus)
extern "C"
#endif
void release();

#if defined(__cplusplus)
extern "C"
#endif
void wrapper_for_cuda(void);

#endif

[sample.c]

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <unistd.h>
#include <omp.h>
#include <time.h>
#include <x86intrin.h>
#include "Nelem.h"
#include "sample_cuda.h"

#define alignof(x) __alignof__(x)

float *a, *b, *c;

double elapsed(struct timespec start, struct timespec end) {
  time_t second = end.tv_sec - start.tv_sec;
  long nsecond = end.tv_nsec - start.tv_nsec;
  return (double)second + (double)nsecond / (1000.0 * 1000.0 * 1000.0);
}

typedef void (*CBFUNC)(void);

static inline void simd_add(int i) {
    __m256 ma = _mm256_load_ps(&a[i]);
    __m256 mb = _mm256_load_ps(&b[i]);
    __m256 mc = _mm256_add_ps(ma, mb);
    _mm256_store_ps(&c[i], mc);
}

static void single_nosimd_add(void) {
  int i;
  for (i = 0; i < NELEM; ++i) {
    c[i] = a[i] + b[i];
  }
}

static void single_simd_add(void) {
  int i;
  for (i = 0; i < NELEM; i += 8) {
    simd_add(i);
  }
}

static void multicore_nosimd_add(void) {
  int i;

  #pragma omp parallel for private(i)
  for (i = 0; i < NELEM; ++i) {
    c[i] = a[i] + b[i];
  }
}

static void multicore_simd_add(void) {
  int i;
  #pragma omp parallel for private(i)
  for (i = 0; i < NELEM; i += 8) {
    simd_add(i);
  }
}

char *cbtitle[] = { "single-nosimd", "single-simd",
                    "multicore-nosimd", "multicore-simd",
                    "cuda" };
CBFUNC cbfunc[] = {
  single_nosimd_add, single_simd_add,
  multicore_nosimd_add, multicore_simd_add,
  wrapper_for_cuda
};
#define NCBFUNC (sizeof(cbfunc)/sizeof(cbfunc[0]))

int main() {
  int i;
  float *f;
  struct timespec start, end;

#if defined(_OPENMP)
  {
    int np = omp_get_num_procs();
    int nth = omp_get_max_threads();
    printf("version: %d\n", _OPENMP);
    printf("#%d processors, %d threads\n", np, nth);
  }
#endif
 
  a = (float *)_mm_malloc(sizeof(float) * NELEM, 32);
  b = (float *)_mm_malloc(sizeof(float) * NELEM, 32);
  c = (float *)_mm_malloc(sizeof(float) * NELEM, 32);

  for (i = 0; i < NELEM; ++i) {
    a[i] = ((float)random() / (float)RAND_MAX);
    b[i] = ((float)random() / (float)RAND_MAX);
  }
  prepare(a, b, c);

  printf("start parallel calculation.\n");

  for (i = 0; i < NCBFUNC; ++i) {
    clock_gettime(CLOCK_REALTIME, &start);
    printf("<<%s>>\n", cbtitle[i]);
    (*cbfunc[i])();
    clock_gettime(CLOCK_REALTIME, &end);
    printf("elapsed %f seconds\n\n", elapsed(start, end));
  }
  

  for (i = 0; i < 10 && i < NELEM; ++i) {
    printf("%f ", c[i]);
  }
  fputs("\n", stdout);

  f = getResult(c);
  for (i = 0; i < 10 && i < NELEM; ++i) {
    printf("%f ", f[i]);
  }
  fputs("\n", stdout);
 
  _mm_free(a);
  _mm_free(b);
  _mm_free(c);
  release();
  return 0;
}

[sample_cuda.cu]

#include 
#include 
#include "Nelem.h"

#define NTHREAD (256)

float *pa, *pb, *pc;

__global__
void kernel(float *xc, float *a, float *b) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < NELEM) {
    xc[i] = a[i] + b[i];
  }
  __syncthreads();
}

__host__
void invoke_kernel(float *c, float *a, float *b) {
  int block = NELEM / NTHREAD;
  kernel<<<block, NTHREAD>>>(c, a, b);
}

extern "C" void prepare(float *a, float *b, float *c) {
  cudaMalloc((void**)&pa, sizeof(float) * NELEM);
  cudaMalloc((void**)&pb, sizeof(float) * NELEM);
  cudaMalloc((void**)&pc, sizeof(float) * NELEM);
  cudaMemcpy(pa, a, sizeof(float) * NELEM, cudaMemcpyHostToDevice);
  cudaMemcpy(pb, b, sizeof(float) * NELEM, cudaMemcpyHostToDevice);
}

extern "C" float *getResult(float *c) {
  cudaMemcpy(c, pc, sizeof(float) * NELEM, cudaMemcpyDeviceToHost);
  return c;
}

extern "C" void release(void) {
  cudaFree(pa);
  cudaFree(pb);
  cudaFree(pc);
}

extern "C" void wrapper_for_cuda(void) {
  invoke_kernel(pc, pa, pb);
  cudaDeviceSynchronize();
}

[Makefile]

SRC=sample.c
OBJ=$(SRC:.c=.o)
CUDASRC=sample_cuda.cu
CUDAOBJ=$(CUDASRC:.cu=.o)
TARGET_OPENMP=sample

all: $(TARGET_OPENMP)

$(TARGET_OPENMP): $(OBJ) $(CUDAOBJ)
        nvcc -g -O0 -o $(TARGET_OPENMP) --gpu-architecture=compute_61 --gpu-code=sm_61 --compiler-options="-fopenmp -DAVX2 -march=native" $(OBJ) $(CUDAOBJ)

$(CUDAOBJ): $(CUDASRC)
        nvcc --gpu-architecture=compute_61 --gpu-code=sm_61 -c $(CUDASRC)

.c.o:
        gcc -c -g -O0 -fopenmp -DAVX2 -march=native $<

clean:
        -rm $(TARGET_OPENMP) $(OBJ) $(CUDAOBJ)

 

unordered_map から unordered_set は作れないのか?

連想配列を使っていると、キーの集合だけ、値の集合だけを作成したいことがある。そんな時はどうするのがスマートなのだろうか。

1.for 文で地道に拾う

C++ っぽくないので却下。

2.for_each 式+ラムダ式で拾う

using namespace std;

void separate(unordered_set<MapKey> key_set, unordered_key<MapValue> value_set, const unordered_map<MapKey, MapValue> map) {
  for_each(begin(map), end(map), [&key_set, &key_value] (const pair<MapKey, MapValue> &i) { key_set.insert(i->first), value_set.insert(i->second) });
}

Java と違って標準の抜き出しメソッドを用意していないのは set とか他のクラスに実装する可能性を考えているからだろうか。

どっちが正解??

スクリプト処理系で配列変数を定義するにあたり、どちらの実装にするべきか悩んでいる。

  1. 演算子 “[” を定義し、第1引数にベース、第2引数にキーを設定する。
  2. 変数に添え字配列を用意する。

1.の利点は固定数の引数に対して演算子を定義すればよいことで、実装が簡単になる。その分、実行エンジンは遅くなりますが。

2.はイメージとしては

class Content {
  Content *c;
  vector <Content *> pc;
};

として、pc が empty でない場合は連想配列の引数とみなすようにする。

2.の利点は実装が多少煩雑になるが、多重配列を定義する場合の実行速度を上げることができる。(ツリーをたどらず、O(1) の配列要素が利用できるので。)

最初1、の実装で考えていたが、途中で考え直して2.の実装に切り替えた。(コード上はその残骸が残っていますが)ただし、まだ変数を左辺値に指定すると NullException が発生するというバグがある。