DCGANでテキスト生成を試してみた2

こんにちは、Link-Uの町屋敷です。

前回は、Word2Vec(以下W2V)とDCGANで英語のテキストを生成しているサイトがあったので、日本語でも試しよう!と言う事で実装してみたところなかなか残念な結果だったので、今回はもう少しなんとかならないか続きをします。

生成結果を視覚化する

結果を良くするにはパラメーター調整をしていかなければならないんですが、画像と違って文章生成の場合生成された結果が直感でわかりにくくなっています。

そこで、わかりやすくするためにジェネレーターが生成した結果をWord2Vecで文章に戻すのではなく、それを次元圧縮してプロットします。

    def show_tsne(self, epoch, real, fake, save_pic):
        from sklearn.manifold import TSNE
        col = np.zeros(len(real))
        col = np.append(col, np.ones(len(fake)))
        
        comp = TSNE(n_components=2, random_state=0).fit_transform(np.reshape(np.vstack((real,fake)), (-1, self.img_cols * self.img_rows * self.channels)))
        plt.scatter(comp[:, 0], comp[:, 1], c=col)
        plt.xlim(-30,30)
        plt.ylim(-30,30)
        if save_pic:
            plt.savefig("images/sent_gan_%d.png" % epoch)
        else:
            plt.show()
        plt.close()

本物の文章と偽物の文章を受け取ってTSNEで圧縮してプロット。今回は本物を紫、偽物を黄色でプロットする

いい偽物の文章が生成されていれば、2色のプロットは混ざり合うが、そうでなければ分離される。

ためしに前回の最終結果で試してみるとこんな感じ。

Epoch 100

Epoch 1000

完全に分離されているのがわかります。

MeCabからの品詞の情報を使う

上記の確認方法を使ってDやGのモデル、W2Vのベクトルの長さ、学習率などのパラメーターをいろいろいじくり回しましたが良い結果にはなりませんでした。

出力された結果を見てみると、品詞がガバガバになっていたので、MeCabから品詞の情報を抜き出して、W2Vの後ろに追加したら多少マシになるんじゃないたと思って追加しました。

    def AnalyzeWord(self, word):
        gen = self.mecab.parse(word, as_nodes=True)
        for w in gen:
            return w.posid, w.surface

品詞はposidから取得できます。posidさえあれば単語の活用は単語の原型から変形させればいいので、簡単のために入力する単語をすべて原型にしました。

結果はこんな感じ。

Epoch 100

Epoch 1000

多少マシなはなったがまだひどい。

ちなみに生成された文章はこんな感じ。

掘り起こしぎゅれんずつ欄オアフまた。。

掘り起こしぎゅれんずつプロバスケットボールリーグオアフまた。。

掘り起こしぎゅれんずつシティーハンタースペシャルオアフまた。。

最後には。がつくっていうのくらいしか学習できてない感がある。

品詞の並びを生成してそれを単語に変換する

品詞の並びを生成する

流石にこのままでは終われないので作戦変更(妥協)。

直接雑音から文章を生成するのを諦め、DCGANと2つ使った文章の生成を目指す。

1つめのDCGANで雑音から品詞の列を生成するジェネレーターを生成し、

品詞を単語に変換するニューラルネットを2つめのDCGANで生成することを目指す。

まず雑音から品詞の列を生成する。これが出来ないなら話にならない。

今までW2Vを入力していたところの先程の品詞の列のみを入力してDCGAN。

Epoch 100

Epoch 300

Epoch 500

Epoch 1000

Epoch 5000

どうやら品詞の列ぐらいは生成できるようだ。

学習した重みを保存する。転移学習に使えるよう各レイヤーに名前わつけて、(name = ‘PC1’など)combinedではなけgenerateorの重みを保存する。

品詞の列を生成するニューラルネットをpos_makerと名付ける

    def build_pos_maker(self):
        model = Sequential()
    
        model.add(Dense(256 * int(self.img_rows * self.img_cols), activation="relu",
                          input_dim=self.latent_dim, name = 'PD1'))
        model.add(Reshape((self.img_cols , self.img_rows, 256)))
        model.add(UpSampling2D())
        model.add(Conv2D(128, kernel_size=3, strides=2, padding="same", name = 'PC1'))
        model.add(BatchNormalization(momentum=0.8, name = 'PB1'))
        model.add(Activation("relu"))
        model.add(UpSampling2D())
        model.add(Conv2D(128, kernel_size=3, strides=2, padding="same", name = 'PC2'))
        model.add(BatchNormalization(momentum=0.8, name = 'PB2'))
        model.add(Activation("relu"))
        model.add(UpSampling2D())
        model.add(Conv2D(64, kernel_size=3, strides=2, padding="same", name = 'PC3'))
        model.add(BatchNormalization(momentum=0.8, name = 'PB3'))
        model.add(Activation("relu"))
        model.add(UpSampling2D())
        model.add(Conv2D(32, kernel_size=3, strides=2, padding="same", name = 'PC4'))
        model.add(BatchNormalization(momentum=0.8, name = 'PB4'))
        model.add(Activation("relu"))
        model.add(Conv2D(1, kernel_size=3, strides=1, padding="same", name = 'PC5'))
        model.add(Activation("tanh"))
    
        model.summary()
        
        noise = Input(shape=(self.latent_dim,))
        print(noise)
        pos = model(noise)
    
        return Model(noise, Container(noise, pos)(noise), name='pos_gen')

//略

        self.generator.save_weights('pos_gan_weight.h5')

2つの学習機をつなげる

        # The generator takes noise as input and generates imgs
        z = Input(shape=(self.latent_dim,))
        pos = self.pos_maker(z)
        img = Input(shape=self.txt_shape)
        
        fake = self.generator(pos)
        # For the combined model we will only train the generator
        #Fix pos_maker weights
        self.pos_maker.load_weights('pos_gan_weight.h5', True)
        self.pos_maker.trainable = False

z(雑音)をpos_makerで品詞の列に変換して、generatorで単語にする。pos_makerの重みは更新しない。

2つめの学習機を初期化する

入力に品詞出力に単語のW2V列をセットして学習するこれを初期値とする。

    def train_pos_to_word(self):
        pos_to_w2v = self.generator
        pos_to_w2v.compile(loss='mse',
        optimizer='adam')
        X_train = joblib.load('{0}/pos_labels.pkl'.format(WRITE_JOBLIB_DIR))
        X_train = np.reshape(X_train, (-1,self.img_cols, self.img_rows, 1))
        X_test = joblib.load('{0}/vectorized_words.pkl'.format(WRITE_JOBLIB_DIR))
        X_test = np.reshape(X_test, (-1, self.img_cols, self.img_rows , self.channels))
        pos_to_w2v.fit(X_train, X_test, batch_size=32,
            epochs=12,
            verbose=1)

結果

かかった時間は手元のパソコンとGPUサーバーで2倍ほどの差があった。

Epoch5

Epoch40

Epoch400

長く学習するとモードがなくなって悪化する。このほうほうではすく偏るらしく元のサイトでも課題になっていた。

Epoch40のときの生成分がこちら。

収録オニヅカ原典する通りは単体。

がけっぷち相当し初めてリリース。。。

詳細は記述実際は彼奴。。

すべてリリース必ずた。。。。

巻が第発売全リリース。。

は本来の年リリース。。。

はエピローグ相当する両性それぞれ取消。

が」全巻。。。。

巻から日発売。。。。

スケリグカン・シヌ。。。。。。

第初めてリリースは取消。。。

ミッションイー。。。。。。。

元データが会話文ではないので名詞が多い。

またこの手のやつは大体そうだが、別に意味を理解しているわけではないので、意味不明な文も結構生成されている。

まとめ

DCGANで文章生成は品詞の列は作れることはわかった。

ただ文章生成になると微妙なので、別の手法(別のGANとかRNN系)も試してみて比較したほうが良さそう。

DCGANでテキスト生成を試してみた

こんにちは、Link-Uの町屋敷です。

今回は生成モデルで有名なGANさんで、Word2Vecを使って英語の文章を生成するものを見つけたので、日本語でもできるのかやってみました。

DCGANのサンプルを動かしてみる

今回はMITライセンスのこのプロジェクトのDCGANをフォークして使っていきます。

まずはそのままのコードを確認していきましょう。

DCGANはGANの一種です。
GAN(Generative Adversarial Networks)は、2つの学習器を互いに争わせることで、学習を進めていきます。2つの学習器とは、
入力された乱数からほしいもの(画像や文章など)を作るジェネレーター(以降G:生成器)と
画像や文章などを入力とし、それがGから生成されたなのかデータセットとして収集してきた本物なのかを見分けるディスクリミネーター(D:識別器)です。

サンプルでは、build_generator()でGの、build_discrimanator内でDの構造が定義されています。

今回はDCGANなので、両方の学習器で畳み込み層を使用しますがプーリング層はありません。

このサンプルでは、コンストラクタで使うネットワークの形と学習方法、損失関数などを定義しています。

Dでは、生成された画像と本物の画像を見分けることを目ざします。生成された画像に0のラベルを、本物の画像に1のラベルをつければ、通常の画像分類と同様なので実装は簡単です。

    # Build and compile the discriminator
        self.discriminator = self.build_discriminator()
        self.discriminator.compile(loss='binary_crossentropy',
            optimizer=optimizer2,
            metrics=['accuracy'])

Gもやっていることは簡単で、ノイズを入力して生成された画像をDに入れて、より高い確率で本物の画像たど判別されるように学習します。

プログラムにするとこうです。

何も考えずに作ってしまうと、Gのモデルの中にDが含まれているので、Gを学習するときにDも学習されてしまいます。

そこでGを学習するときにはDの値を変えないようself.discriminator.trainable = Falseで制限をかけています。

        # The generator takes noise as input and generates imgs
        z = Input(shape=(self.latent_dim,)) #雑音の形の定義
        img = self.generator(z) #生成された画像がimg

         # For the combined model we will only train the generator
        self.discriminator.trainable = False

        # The discriminator takes generated images as input and determines validity
        valid = self.discriminator(img) #Validが判定結果

        # The combined model  (stacked generator and discriminator)
        # Trains the generator to fool the discriminator
        self.combined = Model(z, valid) #入力が雑音で、出力が判定結果のモデル
        self.combined.compile(loss='binary_crossentropy', optimizer=optimizer)

あとは、train関数内でデータを入れたら完成です。

実際に走らせてみましょう。

0Epoch

ただの雑音が

200Epoch

4000Epoch

読める数字になっていっています。

文章生成用に改造する

さてここからが本題です。

先程のサンプルでは28*28の画像を入力にしていましたが、代わりに文字を変換して入力します。

そのままの文字を入力してもどうにもならないので、何らかの変換をしないといけないんですが、それが今回使うWord2Vecです。

正解の文章は前にWikipediaから取ってきた漫画の記事(約5000記事)のデータを使いまわします。

Word2Vecはその名の通り単語をベクトルに変換する技術で、これを使うことによって単語と単語の計算ができるようになります。

        all_sentences = joblib.load('{0}/all_sentences.pkl'.format(WRITE_JOBLIB_DIR))
        sentences = [[word for word in document.lower().split()] for document in all_sentences]
        
        print("Building Word2Vec")
        word_model = Word2Vec(sentences, size=63, min_count=1, window=5) 
        joblib.dump(word_model, '{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))

変換はgensimを使って行います。Word2Vec関数に文章を渡せば辞書が出来ます。。簡単!

先にMecabでの分かち書きを忘れないよう注意。引数のsizeは変換後のベクトルの長さ、windowは。文章中の単語と単語の関係性を計算する距離を表す。この例だと5単語まで。

def VectorizeWord(self):
        all_sentences = joblib.load('{0}/all_sentences.pkl'.format(WRITE_JOBLIB_DIR))
        sentences = [[word for word in document.lower().split()] for document in all_sentences]
        word_model = joblib.load('{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
               
        n_words = 8
        # converted_sentences = [] 
        # converted_sentence = np.zeros(word_model.syn0.shape[0])
        input_data = []
        for s in sentences:
            vectorized_sentence = []
            word_count = 0
            for w in s:
                vector = word_model.wv[w]
                vectorized_sentence.append(vector)
                word_count += 1
                if w == '。':
                    if word_count < n_words:
                        for i in range(n_words - word_count):
                            vector = word_model.wv['。']
                            vectorized_sentence.append(np.append(vector, -1))
                            word_count += 1
                    if word_count == n_words:
                        input_data.append(vectorized_sentence)
                        vectorized_sentence = []
                        word_count = 0
                    elif word_count > n_words:
                        vectorized_sentence = []
                        word_count = 0
        print(np.shape(input_data))
        print(n_words * (word_model.layer1_size + 1))
        input_data = np.reshape(input_data, (-1, n_words * (word_model.layer1_size + 1)))  # 品詞IDの分がたされる
        print(np.shape(input_data))
        print(n_words * len(sentences))
        word_model = joblib.dump(input_data, '{0}/vectorized_words.pkl'.format(WRITE_JOBLIB_DIR))
        print(1)
            

作った辞書を使って単語をベクトルに変換していきます。(vector = word_model.wv[w])の部分。

ついでに1文章に含まれる単語数はまちまちなので、ある一定数に揃えます。一定数に満たない文章は空白で埋めたかったんですが、辞書に登録されてない単語は変換するときにエラーになるので「。」で埋めました。「行き,まし,た,。」 => 「行き,まし,た,。,。,。,。,。」

これで準備完了。先程の画像の代わりに変換した単語を入れていきます。

さてここで問題になるのは、どの形で、単語を渡すかです。

一つは、一単語を行ベクトルを列として(単語数*ベクトル数*1)の画像みたいに渡す方法。

ベクトルを画像で言うRGBみたいなものだと考えて(1*単語数*ベクトル数)もあります。

結論から言うと前者はあまりよろしくなかったです。

生成例がこれなんですけど

epoch0

・ メルチェリーダぷれりゅうどぷれりゅうどガンムマーズクロニクル歌野ゃがんはがちりんにとぶアプトアプトばろそぐいアダルトグッズショップばろケルビム・ゲート・キーパーゃがんはがちりんにとぶキザキ売り込もゃがんはがちりんにとぶゃがんはがちりんにとぶオゲ売り込もキザキ売り込も受け流そ歌野歌野ザ・ビーンズ歌野臨もザ・ビーンズ歌野ザ・ビーンズメルチェリーダ弖虎弖虎弖虎弖虎弖虎メルチェリーダアオノキセキウキウキペディアカードウキウキペディアカード

epoch400

・ 使用二鷹野大遠藤広隆小暮昌広程度桁異なり本当に他紙足取り後ろ他付属ボーナストラックトラック柔弱。。。

epoch0でも兆候が見えてるですが、epoch400ではなんか前の方に名詞固まってるし、トラックトラックみたいな繰り返しがよく起こってます。

とくにこの繰り返しが至るところに出現します。多分ゲネレーターの畳み込み層のせいだと思ったので、入力を後者の方に変更。単語数64も多すぎたので8まで減らした。

結果はこちら

epoch 0

・ 若作り繋っシャンツァカッシーギャグイニチアシブノレ混交デイブレイク

・ 年代デイブレイクビジネスインフラ・リミテッドオフィシャルパンフレット商業ドルドルーヴォドラゴン・サーガジャケットイラストギャラリー

・ 交ぜるスパークオンウェイヴオフィシャルパンフレット描くソガシイナ京田辺イリヤルートドラゴンボールメニュー

・ 福見オールドマン・パーサキュライナーツノートぬおイニチアシブゲームエッセイマンガ木城

・ あかしあ台ノスタルジーゅせんきょうジミナニードルアイフレスポ鬩ホラー

epoch 100

・ 選評再現カレイドスター・ウィンディステージ数。。。。

・ ホワイトナックル・クリムゾンオーブハンター両替キャストバトルオブフェアリーテイル選り抜いゲームパッケージイラスト全て。

・ カードダスステーション終盤ぶった斬る以後後半フルタッチ・アクション。。

・ 温州ニコラ・ケフェウス一般に約前後。。。

・ 時点キディ・ガーランドとおり前後予定同権。。

epoch 200

・ キャラウムカフェウルトラジャンプエッグサイト差し置き。日本一が上回り驚愕

・ ワイディーネ学期エロティック・コメディ。メジャーが上回りから

・ キャラウムカフェ学期エロティック・コメディガールズ・デート年代は上回り共に

・ ドルビーステレオコンサンタルト学期エロティック・コメディガールズ・デートアニヴァーサリーが上回り共に

・ キャラウムカフェ学期エロティック・コメディガールズ・デート以後が上回り奪え

さっきよりは多少マシになったが、まだまだ不自然。

epoch200ではモード崩壊気味でなぜか末尾の「。」が消えた。

ここから改善を始めていくが続きは来月に。

まとめ

一応生成は出来たが精度はまだまだ。

改善点は色々あるので来月はそれをやっていこうと思う。

具体的にはパラメーター調整とかMeCabの情報を使ってConditional GAN化してみるとかそもそもDCGANをやめるとかかなあ。

プログラム全文

'''
MIT License

Copyright (c) 2017 Erik Linder-Norén

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.
'''

import re
import json

from timeit import default_timer as timer

import keras.backend as K

import numpy as np
import matplotlib
from matplotlib import pyplot as plt
import joblib

from natto import MeCab
import gensim
from gensim.models.word2vec import Word2Vec

try:
    from keras.engine.topology import Container
except:
    from keras.engine.network import Network as Container
from keras.datasets import mnist
from keras.layers.recurrent import LSTM
from keras.layers.embeddings import Embedding
from keras.models import Model, Sequential
from keras.layers import Dense, Activation, Reshape, UpSampling2D, Conv2D, BatchNormalization, Input, Dropout, ZeroPadding2D, Flatten
from keras.optimizers import Adam
from keras.layers.advanced_activations import LeakyReLU
from keras.losses import binary_crossentropy

TEXT_JSON_DIR = '../WikipediaComic/whole_data'
INFOBOX_JSON_DIR = '.'
INFOBOX_FILE_NAME = 'wiki_infobox_Infobox_animanga_Manga.json'
WRITE_JSON_DIR = '.'
WRITE_JSON_FILE_NAME = 'joined.json' 
WRITE_TEXT_RESULT_DIR = '.'
AUTO_ENCODER_DIR = '.'

WRITE_JOBLIB_DIR = '.'

# Matplotlibの日本語設定
font_path = '/usr/share/fonts/truetype/takao-gothic/TakaoPGothic.ttf'
font_prop = matplotlib.font_manager.FontProperties(fname=font_path)
matplotlib.rcParams['font.family'] = font_prop.get_name()

def cross_entropy_plus_simmilarity(y_true, y_pred, X_pred):
    simi = 0
    for i , v in enumerate(X_pred):
        for j, _ in enumerate(X_pred):
            if i > j:
                simi += K.square(X_pred[i] - X_pred[j])
    n = len(X_pred)
    return simi/(n*(n-1)/2) + binary_crossentropy(y_true, y_pred)

class DCGAN():

    def __init__(self, row, col, additional_featire_count = 0):
        self.additional_feature_count = additional_featire_count
        # Input shape
        self.img_rows = 1
        self.img_cols = col
        self.channels = row + self.additional_feature_count
        
        self.txt_shape = (self.img_cols, 1, self.channels)
        self.latent_dim = 200

        optimizer = Adam(0.0002, 0.5)

        # Build the generator
        self.generator = self.build_generator()
#        self.generator = self.WordGenerator()

        # Build and compile the discriminator
        self.discriminator = self.build_discriminator()
        self.discriminator.compile(loss='binary_crossentropy',
            optimizer=optimizer,
            metrics=['accuracy'])

        # The generator takes noise as input and generates imgs
        z = Input(shape=(self.latent_dim,))
        img = self.generator(z)

        # For the combined model we will only train the generator
        self.discriminator.trainable = False

        # The discriminator takes generated images as input and determines validity
        valid = self.discriminator(img)

        # The combined model  (stacked generator and discriminator)
        # Trains the generator to fool the discriminator
        self.combined = Model(z, valid)
        self.combined.compile(loss='binary_crossentropy', optimizer=optimizer)
        
    def build_generator(self):
    
        model = Sequential()
    
        model.add(Dense(128 * int((self.img_cols) * self.img_rows), activation="relu", input_dim=self.latent_dim))
        model.add(Reshape((self.img_cols , self.img_rows, 128)))
        model.add(Conv2D(128, kernel_size=3, padding="same"))
        model.add(BatchNormalization(momentum=0.8))
        model.add(Activation("relu"))
        model.add(Conv2D(64, kernel_size=3, strides=1, padding="same"))
        model.add(BatchNormalization(momentum=0.8))
        model.add(Activation("relu"))
        model.add(Conv2D(32, kernel_size=3, strides=1, padding="same"))
        model.add(BatchNormalization(momentum=0.8))
        model.add(Activation("relu"))
        model.add(Conv2D(self.channels, kernel_size=3, padding="same"))
        model.add(Activation("tanh"))
    
        model.summary()
    
        noise = Input(shape=(self.latent_dim,))
        img = model(noise)
    
        return Model(noise, Container(noise, img)(noise)) 

    def build_discriminator(self):

        model = Sequential()
        print(self.txt_shape)
        model.add(Conv2D(32, kernel_size=3, strides=2, input_shape=self.txt_shape, padding="same"))
        model.add(LeakyReLU(alpha=0.2))
        model.add(Conv2D(64, kernel_size=3, strides=2, padding="same"))
        model.add(ZeroPadding2D(padding=((0, 1), (0, 1))))
        model.add(BatchNormalization(momentum=0.8))
        model.add(LeakyReLU(alpha=0.2))
        model.add(Dropout(0.5))
        model.add(Conv2D(128, kernel_size=3, strides=2, padding="same"))
        model.add(BatchNormalization(momentum=0.8))
        model.add(LeakyReLU(alpha=0.2))
        model.add(Flatten())
        model.add(Dense(1, activation='sigmoid'))

        model.summary()

        img = Input(shape=self.txt_shape)
        validity = model(img)

        return Model(img, Container(img, validity)(img))

    def train(self, epochs, batch_size=128, save_interval=50):
        word_model = joblib.load('{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
        X_train = joblib.load('{0}/vectorized_words.pkl'.format(WRITE_JOBLIB_DIR))
        print(np.shape(X_train))
        X_train = np.reshape(X_train, (-1, self.img_cols , 1, self.channels))
        print(np.shape(X_train))
        # Adversarial ground truths
        real = np.ones((batch_size, 1))
        fake = np.zeros((batch_size, 1))

        for epoch in range(epochs):

            # ---------------------
            #  Train Discriminator
            # ---------------------

            # Select a random half of images
            idx = np.random.randint(0, X_train.shape[0], batch_size)
            imgs = X_train[idx]

            # Sample noise and generate a batch of new images
            noise = np.random.normal(0, 1, (batch_size, self.latent_dim))
            gen_imgs = self.generator.predict(noise)

            # Train the discriminator (real classified as ones and generated as zeros)
            self.discriminator.trainable = True
            self.generator.non_trainable = False
            d_loss_real = self.discriminator.train_on_batch(imgs, real)
            d_loss_fake = self.discriminator.train_on_batch(gen_imgs, fake)
            d_loss = 0.5 * np.add(d_loss_real, d_loss_fake)

            # ---------------------
            #  Train Generator
            # ---------------------

            # Train the generator (wants discriminator to mistake images as real)
            self.discriminator.trainable = False
            self.generator.non_trainable = True
            g_loss = self.combined.train_on_batch(noise, real)

            # Plot the progress
            print ("%d [D loss: %f, acc.: %.2f%%] [G loss: %f]" % (epoch, d_loss[0], 100 * d_loss[1], g_loss))

            # If at save interval => save generated image samples
            if epoch % save_interval == 0:
                self.show_text(epoch)
    
    def show_text(self, epoch):
        r, c = 5, 5
        noise = np.random.normal(0, 1, (r * c, self.latent_dim))
        gen_text = self.generator.predict(noise)
        word_model = joblib.load('{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
        r, c = 5, 5
        noise = np.random.normal(0, 1, (r * c, self.latent_dim))
        gen_text = self.generator.predict(noise)
        word_model = joblib.load('{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
        with open(WRITE_TEXT_RESULT_DIR + '/generated_epoch{0}.txt'.format(epoch), "w+") as f:
            f.write('Epoch {0}rn'.format(epoch))
            for vector in gen_text:
                word = ''
                s = vector.flatten().reshape(self.img_cols, self.img_rows, self.channels)
                for w in s:   
                    if self.additional_feature_count > 0:
                        word = word + word_model.most_similar(w[:,0:-1*self.additional_feature_count])[0][0]
                    else:
                        word = word + word_model.most_similar(w)[0][0]
                print(word + 'rn')
                f.write(word)
    
    def save_imgs(self, epoch):
        r, c = 5, 5
        noise = np.random.normal(0, 1, (r * c, self.latent_dim))
        gen_imgs = self.generator.predict(noise)

        # Rescale images 0 - 1
        gen_imgs = 0.5 * gen_imgs + 0.5

        fig, axs = plt.subplots(r, c)
        cnt = 0
        for i in range(r):
            for j in range(c):
                axs[i, j].imshow(gen_imgs[cnt, :, :, 0], cmap='gray')
                axs[i, j].axis('off')
                cnt += 1
        fig.savefig("images/mnist_%d.png" % epoch)
        plt.close()


class ProcessWord():

    def __init__(self, word_max, vector_size, do_use_pos = True):
        self.mecab = MeCab()
        self.word_max = word_max
        self.vector_size = vector_size
        self.do_use_pos = do_use_pos
        
    def DcganTrigger(self, epochs=4000, batch_size=64, save_interval=10):
        if self.do_use_pos:
            dg = DCGAN(self.vector_size, self.word_max, 1)
        else:
            dg = DCGAN(self.vector_size, self.word_max, 0)
        dg.train(epochs=epochs, batch_size=batch_size, save_interval=save_interval)

    def ExtractWords(self):
        with open(WRITE_JSON_DIR + '/' + WRITE_JSON_FILE_NAME, 'r', encoding='utf_8') as jw: 
            json_data = json.load(jw)
        
        all_words = [[0]] * len(json_data) 
        with MeCab("-Owakati") as mecab:
            for i, j in enumerate(json_data):
                if i % 100 == 0:
                    print(i)
                text = j['text']
                
                all_words[i] = re.sub(r'[!-~]', "", mecab.parse(text))  # 記号、英語除去
                
        joblib.dump(all_words, '{0}/all_sentences.pkl'.format(WRITE_JOBLIB_DIR), compress=True)
        
    def CreateWord2Vec(self):
        all_sentences = joblib.load('{0}/all_sentences.pkl'.format(WRITE_JOBLIB_DIR))
        sentences = [[self.AnalyzeWord(word)[1] for word in document.lower().split()] for document in all_sentences]
        
        print("Building Word2Vec")
        word_model = Word2Vec(sentences, size=self.vector_size, min_count=1, window=5)
        joblib.dump(word_model, '{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
        # Code tried to prepare LSTM model for word generation
        
    def VectorizeWord(self):
        all_sentences = joblib.load('{0}/all_sentences.pkl'.format(WRITE_JOBLIB_DIR))
        sentences = [[word for word in document.lower().split()] for document in all_sentences]
        word_model = joblib.load('{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
               
        n_words = self.word_max
        # converted_sentences = [] 
        # converted_sentence = np.zeros(word_model.syn0.shape[0])
        input_data = []
        pos_labels = [] 
        for s in sentences:
            vectorized_sentence = []
            pos = []
            word_count = 0
            #print(s)
            for w in s:
                posid, surface = self.AnalyzeWord(w)
                vector = word_model.wv[surface]
                if self.do_use_pos:
                    vector = np.append(vector, float((posid - 34) / 68))  # POSID正規化
                pos.append(float((posid - 34) / 68))
                vectorized_sentence.append(vector)
                word_count += 1
                if w == '。':
                    if word_count < n_words:
                        for i in range(n_words - word_count):
                            vector = word_model.wv['。']
                            if self.do_use_pos:
                                vectorized_sentence.append(np.append(vector, -1))
                            else:
                                vectorized_sentence.append(vector)
                            pos.append(-1)
                            word_count += 1
                    if word_count == n_words:
                        input_data.append(vectorized_sentence)
                        pos_labels.append(pos)
                        vectorized_sentence = []
                        pos =  []
                        word_count = 0
                    elif word_count > n_words:
                        vectorized_sentence = []
                        pos = []
                        word_count = 0
                        
        print(np.shape(input_data))
        print(n_words * (word_model.layer1_size + 1))
        if self.do_use_pos:
            input_data = np.reshape(input_data, (-1, n_words * (word_model.layer1_size + 1)))  # 品詞IDの分がたされる
        else:
            input_data = np.reshape(np.array(input_data), (-1, n_words * (word_model.layer1_size)))
        print(np.shape(input_data))
        print(n_words * len(sentences))
        word_model = joblib.dump(input_data, '{0}/vectorized_words.pkl'.format(WRITE_JOBLIB_DIR))
        joblib.dump(input_data, '{0}/vectorized_words.pkl'.format(WRITE_JOBLIB_DIR))
        joblib.dump(pos_labels, '{0}/pos_labels.pkl'.format(WRITE_JOBLIB_DIR))
        print(1)
        
    def CheckWord2Vec(self, word1, word2, ope):
        word_model = joblib.load('{0}/word2vec.pkl'.format(WRITE_JOBLIB_DIR))
        if ope == '+':
            w = word_model.wv[word1] + word_model.wv[word2]
        if ope == '-':
            w = word_model.wv[word1] - word_model.wv[word2]
        if ope == 'all' or ope == 'ALL':
            self.CheckWord2Vec(word1, word2, '+')
            self.CheckWord2Vec(word1, word2, '-')
            self.CheckWord2Vec(word2, word1, '-')
            return    
        print(word_model.most_similar([w]))
        
    def AnalyzeWord(self, word):
        gen = self.mecab.parse(word, as_nodes=True)
        for w in gen:
            return w.posid, w.surface

if __name__ == '__main__':
    begin = timer()
    pw = ProcessWord(8, 32, False)    
    #pw.ExtractWords()
    #pw.CreateWord2Vec()    
    pw.VectorizeWord()
    pw.DcganTrigger(4000, 256, 10)
    print(timer() - begin)
    pw.CheckWord2Vec('学院', '悪', 'ALL')

GPUを使って無線LANをクラックする話 Pythonのプロファイルを取る回

近所を散歩していたら、もう梅が咲いていました。今日は立春だそうです。早いなー。

わたしは梅のつぼみが大好きなので、これからはしばらく目が離せない日々が続きそうです。

前回のおさらい

前回得た結果は、「CPUとGPUを同時に使うより、GPUだけを使う方が実は倍ぐらい速い!」という、これまた逆説的な結果でした。

未改造のPyritでは「CPUとGPUを同時」に使うか、「CPUだけ」を使うかなので、なまじ「同時」が「CPUだけ」よりは速いが故に、そんな単純な事にも気づかなかった、と。

うーん…なんでこんな事に…。GPUを追加で使うぶんだけ、せめてちょびっとでも性能が向上してほしかったんですが…。

Pythonのコードのプロファイルを取ろう

これまでにCUDAコードのプロファイルを二回ほど取りました(一回目二回目)が、今日はPythonのコードのプロファイルを取りましょう。もちろんPythonはCPUで実行されるわけですが、さらにPythonは高々1つのスレッドでしか同時に動作しないのでした。

今回Pythonのプロファイルを取ろうと画策するのは、次の観察からです:

  • 2CPU、48コアを全部ぶん回すと48倍の性能がでない。htopを観察する限り、かといってCPUを全部使っているわけでもなさそう
    • Pythonの1コアが各コアに仕事を分配しきれていないと考えると説明できます(わかんないけどね)
  • GPUだけを使うにしても、GPUは使い切っていなさそう
    • ここがよくわからない。GPUが消費しきる分のパスワードは生成できそうなので。
  • ダミーパスワードの生成に何割くらい掛けているのかが気になる

というわけでモリモリ取っていきまっしょい。プロファイルのためのツールは公式で提供されていて、cProfileというのを使えば良さそうです。

% python -m cProfile -s tottime <実行ファイル名> <args...>

とやると、関数の内部の実行に掛かった順番でソートされた結果が帰ってきます。tottimeは、関数の中で他の関数を呼び出した時間は除いた、その関数の純粋な実行時間。cumtimeと入れると、その関数から呼び出した関数の時間も含まれます。

例えば、CPUを1コアだけ使った時の表示はこんな感じになります:

(venv) server01@server01:~/src/Pyrit$ time python2.7 -m cProfile -s tottime ./pyrit benchmark
Pyrit 0.5.1 (C) 2008-2011 Lukas Lueg - 2015 John Mora
https://github.com/JPaulMora/Pyrit
This code is distributed under the GNU General Public License v3+

Running benchmark (985.4 PMKs/s)... -  

Computed 985.36 PMKs/s total.
#1: 'CPU-Core (SSE2/AES)': 1045.2 PMKs/s (RTT 2.8)
         79276 function calls (78919 primitive calls) in 69.771 seconds

   Ordered by: internal time

   ncalls  tottime  percall  cumtime  percall filename:lineno(function)
     1604   68.890    0.043   68.890    0.043 {time.sleep}
     1963    0.514    0.000    0.514    0.000 {method 'acquire' of 'thread.lock' objects}
        1    0.158    0.158    0.197    0.197 cpyrit.py:29(<module>)
        1    0.094    0.094   69.571   69.571 pyrit_cli.py:1184(benchmark)
       51    0.028    0.001   69.438    1.362 threading.py:309(wait)
    54392    0.011    0.000    0.011    0.000 {method 'random' of '_random.Random' objects}

...(略)...

real	1m9.902s
user	1m9.087s
sys	0m0.302s

ログの一番上の行を見ることで、time.sleepで68.9秒とたくさん時間を使っていること、次に使っているのがacquireメソッドであること、がわかります。

さらに、

   ncalls  tottime  percall  cumtime  percall filename:lineno(function)
        1    0.094    0.094   69.571   69.571 pyrit_cli.py:1184(benchmark)
real	1m9.902s

この二行を見ると、このベンチマークを実際に実行するbenchmark関数が69.5秒使っていること、処理時間が69.9秒で、その前後の初期化などに0.4秒ぐらい使われていること(まぁ、予想通りですよね)、などがわかりますよー、と。

この2つの情報を合計すると、CPUを1コアだけ使った時は、69.5秒中68.9秒、ほとんどtime.sleepで時間を潰していること(1コアのCPUが処理するのを待っていること)がなんとなくわかります。

以下前回みたく条件を変えて実行してみましたが、だらだらログを並べても中々わかりにくいので、表にしてみました:

条件1番時間を使っている関数
(全処理に占める割合)
2番時間3番時間4番時間
1cpu/0gpu{time.sleep}
(99.0%)
{method ‘acquire’ of ‘thread.lock’ objects}
(0.7%)
cpyrit.py(<module>)
(0.2%)
pyrit_cli.py(benchmark)
(0.1%)
48cpu/0gpu{time.sleep}
(92.0%)
{method ‘acquire’ of ‘thread.lock’ objects}
(4.5%)
pyrit_cli.py(benchmark)
(2.21%)
cpyrit.py(<module>)
(0.22%)
0cpu/4gpupyrit_cli.py(benchmark)
(61.3%)
{method ‘flush’ of ‘file’ objects}
(14.6%)
{method ‘acquire’ of ‘thread.lock’ objects}
(12.4%)
cpyrit.py(dequeue)
(4.04%)
48cpu/4gpu{method ‘acquire’ of ‘thread.lock’ objects}
(65.2%)
pyrit_cli.py:1184(benchmark)
(16.5%)
cpyrit.py(dequeue)
(5.0%)
{time.sleep}
(3.8%)

だいたい上位に並ぶのはどれも同じ顔ぶればかりなのですが、下のほうに行けばいくほど、つまり、たくさんのCPUやGPUが仕事を要求するようになればなるほど、sleepよりも{method ‘acquire’ of ‘thread.lock’ objects}の処理時間がガンガン増えていく事がわかります。一番下の一番処理を取り合っているところではなんと6割もロックにつぎ込んでいます。これじゃあ、GPUのパスワードのクラックではなく、ロックを取り合うプログラムを実行していたと言っても過言ではありませんな。

そしてdequeue関数の処理内容が増えていくのも気になりますねぇ。この関数は、生成したパスワードをCPUやGPUが受け取るためのもので、threading.Conditionを使ってマルチスレッドの調停を行っております。

この2つから予想されることは、…おそらくこのthreading.Conditionオブジェクトの中にあるロックを取り合ってるんでしょうね…。

あまりにも闇が深そうなので今日はこのぐらいにしておきましょう。

GPUだけを使っているケース(三行目)については追加でコメントさせてください。benchmark関数が一番時間を使っていることから、一生懸命処理するためのパスワードを用意していることはなんとなく察されるのですが、パスワードの生成に使っているはずのrandom関数がリストアップされてこなくて、そのかわりにfile.flushが二番目に食い込んでいます。このflushはCUDA関係っぽい気がしますが、謎です。もうちょっと追いかける価値があると思います。

まとめ

  • 仕事を要求するスレッドが増えれば増えるほど、仕事を用意するスレッドはsleepではなくlockの取り合いで時間を潰すようになる
    • 1つのロックを取り合ってそう
  • GPUだけで処理しているときはパスワードの生成を頑張るみたいだが、実態は要調査
  • 外へ出ろ、梅の花を見ろ!

今月は低電力モードとなっております。寒いからね、仕方ないね(ごめんなさい)。

Numbaを使ってpyhonコードを高速化する方法

こんにちは、Link-Uの町屋敷です。

今回は、機械学習を行う際にほぼ必ず行わなければならない前処理を、
GPUを使ってやったら早く終わったので、メモがてらに書いていきます。
今回はpythonでの話です。

pythonでcudaを使ったGPU演算をする方法として、
NVIDIAの公式でも紹介されているnumbaを使いたいと思います。

@Vectorizeを使った方法

pipでも入るそうですが、anacondaでもインストールできるのでcondaを使用します。

Anaconda,Minicondaのインストール方法はバックナンバーで、

conda install accelerate

このコマンドを使うとnamba以外にも必要なパッケージが入るので楽です。

早速試してみましょう。

先程の公式ページの動画内に書かれていたコードを使っててみます。

import numpy as np

from timeit import default_timer as timer
from numba import vectorize
from numba import cuda, float32

@vectorize(['float32(float32, float32)'], target='cpu')
def VectorAdd(a,b):
    return a + b

@vectorize(['float32(float32, float32)'], target='cuda')
def GpuAdd(a,b):
    return a + b

def NormalAdd(a,b,c):
    for i, _ in enumerate(a):
        c[i] = b[i] + a[i]    

def main():
    
    for N in [1e2, 1e3, 1e4, 1e5, 1e6, 1e7, 1e8]:
        vec_a = np.ones(int(N), dtype=np.float32)
        vec_b = np.ones(int(N), dtype=np.float32)
        vec_c = np.zeros(int(N), dtype=np.float32)
        
        start = timer()
        NormalAdd(vec_a, vec_b, vec_c)
        normal_time = timer() - start
        
        print('When [{0}] vector, NormalAdd took {1} seconds'.format(N, normal_time))
        
        vec_c = np.zeros(int(N), dtype=np.float32)
        
        start = timer()
        vec_c = VectorAdd(vec_a, vec_b)
        cpu_time = timer() - start
        #print(C)
        
        print('When [{0}] vector, CpuAdd took {1} seconds'.format(N, cpu_time))
        
        start = timer()
        vec_c = GpuAdd(vec_a, vec_b)
        gpu_time = timer() - start
        
        print('When [{0}] vector, GpuAdd took {1} seconds'.format(N, gpu_time))
        print('Normal speed / CPU speed = {0}'.format(normal_time/cpu_time))       
        print('Normal speed / GPU speed = {0}'.format(normal_time/gpu_time))    
        print('CPU speed / GPU speed = {0}'.format(cpu_time/gpu_time))

numbaでは高速化を行いたい関数にデコレータを付けます。

この例では@vectorize([‘float32(float32, float32)’], target=’cuda’)の部分が該当します。

@vectorizeの第一引数には、C言語の関数の宣言ように返り値と引数の型を記述します。

第二引数はをcudaにするとGPUをcpuにするとcpuを使って最適化します。

GPUサーバーで動かした結果はこの通り。

何もしていないときと比べて遥かに早くなりました。

ただ、GPUを使用し始めるのに3.5秒ほど準備期間が必要な模様で、最初の処理に時間がかかっています。

また、CPUのほうが早いという結果に。

桁数がもう少し増えれば逆転しそうだがあまり降らしすぎるとメモリーエラーになります。

公式の動画でtarget=cpuをやっていないのはそういうことがややこしいからなのか……

このままではCPUのほうが良いんじゃないか説が出てきてしまうので、少し処理を追加した関数でやってみましょう。

@vectorize(['float32(float32)'], target='cuda')
def RadToDegGpu(a):
    return (360 + 180 * a / np.pi) % 360

@vectorize(['float32(float32)'], target='cpu')
def RadToDegCpu(a):
    return (360 + 180 * a / np.pi) % 360

def main():
    
    for N in [1e2, 1e3, 1e4, 1e5, 1e6, 1e7, 1e8]:
        A = np.ones(int(N), dtype=np.float32) * np.pi
        B = np.zeros(int(N), dtype=np.float32)
        
        start = timer()
        B = RadToDegCpu(A)
        cpu_time = timer() - start
        
        print('When [{0}] vector, RadToDegCpu took {1} seconds'.format(N, cpu_time))
        
        start = timer()
        B = RadToDegGpu(A)
        gpu_time = timer() - start
         
        print('When [{0}] vector, RadToDegGpu took {1} seconds'.format(N, gpu_time))
        print('x{0} speed'.format(cpu_time/gpu_time))

ラジアンを°に変える簡単な関数です。

nambaはnumpyに最適化されているため関数を使う時は極力numpyを使いましょう(この例だと使ってるのはただの定数だけど)。

引数の数が変わっているので、@vectorizeの第一引数も変わっています。

x~speedの~が1以上だとgpuのほうが早いです。

桁数が小さい時はまだCPUのほうが早いですが、大きくなってくるとGPUを使ったほうが約3倍早くなっています。

ただし最初の準備期間3.5秒は一律でかかるので、この処理単体だけをやるとGPUのほうが遅くなります。

1回の実行では使う関数が変わっても1回しか準備期間はないようなので、封数の処理を一気にやると良いでしょう。

@cuda.jitを使った方法

次に行列を使った計算を高速化する方法をここを参考にして書いていきます。

例としてデータの前処理によく使う各行に対して平滑化を行うプログラムを作ります。

まず、普通の実装

def NormalSmooth(inputMat, outputMat):
    for i, row in enumerate(inputMat):
        for j, v in enumerate(row):
            if j == 0:
                outputMat[i, j] = (inputMat[i, j] + inputMat[i, j] + inputMat[i, j + 1])/3
            elif j == outputMat.shape[1] - 1:
                outputMat[i, j] = (inputMat[i, j - 1] + inputMat[i, j] + inputMat[i, j])/3
            else:
                outputMat[i, j] = (inputMat[i, j - 1] + inputMat[i, j] + inputMat[i, j + 1])/3

そして、@cuda.jitを使った実装

@cuda.jit
def GpuSmooth(inputMat, outputMat):
    row, col = cuda.grid(2)
    if row < outputMat.shape[0] and col < outputMat.shape[1]:
        if col == 0:
            outputMat[row, col] = (inputMat[row, col] + inputMat[row, col] + inputMat[row, col + 1])/3
        elif col == outputMat.shape[1] - 1:
            outputMat[row, col] = (inputMat[row, col - 1] + inputMat[row, col] + inputMat[row, col])/3
        else:
            outputMat[row, col] = (inputMat[row, col - 1] + inputMat[row, col] + inputMat[row, col + 1])/3

cuda.grid(2)を使うことで簡単に2重のfor文の処理を書き換えることが出来ます、

cuda.grid(2)が返す値はx,yともに正の値ですが、上限が行列の形と違うのでif文を使って行列の外になった場合処理をしないことが必要になります。

def main(): 
     
    for R in [10,100,1000,10000]:
        for C in [10,100,1000,10000]:
    
            A = np.ones((R, C))
            B = np.ones((R, C))
            
            start = timer()
            NormalSmooth(B, A)
            normal_time = timer() - start
            
            print('When [{0}, {1}] matrix, NormalSmooth took {2} seconds'.format(R, C, normal_time))
            
            start = timer()
            GpuSmooth(B, A)
            gpu_time = timer() - start
            
            print('When [{0}, {1}] matrix, GpuSmooth took {2} seconds'.format(R, C, gpu_time))
            print('x{0} speed'.format(normal_time/gpu_time))

こちらも行列が小さい時はCPUのほうが優位ですが大きくなってくると最大で129倍と圧倒的にGPUのほうが早くなります。

ちなみに手元のノートパソコンの結果が[10000, 10000]のとき121.75秒かかっていたので100倍以上高速化されたことになります。

まとめ

numbaを使うことでpytonスクリプトを高速化する手法を書いた

大きいデータを扱う時はかなり高速化されるので積極的に使っていきたい。

AIを端末側で処理するチップ達

以前NVIDIAの新しいGPUを紹介しましたが、NVIDIAのGPUはPCで主に利用されると思います。

しかし、恐らくAI(正確には機械学習の推論)を行う回数で言うと圧倒的にスマートフォンの方が多いと思います。

今回はスマートフォン等に搭載されるチップがAIをどう処理するかを紹介していきましょう。今回はどのチップもPVがあったので全編PV付きでお送りします。

Snapdragon 845

スマホ用チップの雄、QUALCOMMの最上位チップです。Androidのハイエンド機は一部の例外を除いてほぼこのチップを使っていると言っても過言ではないかと。

ただし、最近のチップの中では比較的AI系の性能は控えめで、どちらかというと従来からのCPUやGPUといった部分の強い正統派のチップです。

AIに対する対応はCPU+GPU+DSPで行い専用回路は持ちません。

この対応だと、AI用の回路がない分、チップに余裕ができるので、CPUやGPUといった汎用的な回路を強化できる利点があります。

一方特にCPUなどは汎用的な処理ができる分、大量のAI向け演算などは処理量の点でも、電力効率の点でも苦手と言えます。

このあたりAIに関連する処理が、全スマホ利用時間のうちどの程度を占めるか、メーカーによる考え方の違いとなって表れてきます。

QUALCOMMとしては音声認識など局所的にAIを使うことはあっても、全利用時間に占める割合はさほど高くないと踏んでいるのではないでしょうか。

とはいえ、AI性能は全世代のSnapdragon 835比で3倍をうたっていますので、CPUやGPUの性能が30%程度の向上であることを加味すると力を入れてきているなというのはあります。

ところでSnapdragon 845にはWiFiの接続が16倍高速になるという機能向上もあり、性能面で言うとそちらが一番の底上げであったり。

そちらも個人的には気になります。

Kirin 980

最近登場した中国Huawei製のチップ。恐らくHuawei製の端末以外採用されていない・・・はず。

QUALCOMMとは真逆で積極的にAI系の機能を盛り込んでくるメーカーで、昨年のKirin 970からAI系の回路を盛り込んできています。

Kirin 980では回路規模を倍増させており、AI系の性能はKirin 970の2倍以上になっていると思われます。

とはいえ、CPU性能も75%向上したらしく、AI系の性能を特別に向上させた、というよりは全般的に2倍程度になるようにチップを設計したようです。

中国メーカーの技術力の伸びを示すような、かなりアグレッシブな性能のあげ方です。

Apple A12 Bionic

日本で一番普及しているiPhoneの最新版に搭載されているチップ。

Appleは他の会社と違ってちゃんと素の演算回数を公表してくれており毎秒5超回の演算ができます。

全世代からAI系の機能をチップに入れてきていますが、前世代と比較して9倍の性能になっています。

こうして他のチップと比較してみるとAppleが全力でAI系を増強してきていることがわかります。

実はAppleは自社生産かつ大量生産なので、他の会社よりもチップにお金をかけることができます。

なので、チップに他社より多くの機能を入れることができ、AI系の機能を強化する余裕があるのかもしれません。

Movidius NCS

IntelがAIチップベンチャーを買収して作ったUSBデバイス。1Wで100GFlopsです。

専用設計で100GFlopsというのは若干微妙な気もしないではないですが、USBインターフェイスとかも入っていることも考えるとこんなものでしょうか。

このデバイス、USBで刺すと使えるので後付けできるのが最大の魅力です。しかも複数台刺して性能向上なんてこともできるようです。

Raspberry Piに刺して使えるなどD.I.Yでデバイスが作れそうな予感が大変します。

というかIntelもCPUにそういう系の機能つければいいのにと思わなくもないです。

まとめ

各メーカーで取り組み方に差はみられるものの、基本的にどのチップメーカーもAI強化に舵を切っていることがわかります。

現在はカメラ、音声認識などでの利用が多いようですが、AIを使えば便利になるシーンは多いと思います。

TuringアーキテクチャのGPUが発表されました

NVIDIAが新しいGPUシリーズ、Turingシリーズを発表しました。

弊社にあるGPU、GeForce 1080TiはPascalアーキテクチャ、Tesla V100はVoltaアーキテクチャです。

VoltaもTuringもPascalの進化系ではあるのですが、VoltaはGPGPUなどのコンピューティング系、Turingはグラフィックス系です。

GPUの使われ方が二極化してきたので設計を分けたのでしょうか。

まずは性能を比較してみましょう

とりあえず、CPUメーカーの長たるIntelの最上位モデルとNVIDIAのゲーム向け、サーバー向けの現行最上位モデル、そしてTuringのゲーム向け最高峰の浮動小数点演算能力を比較してみましょう。

 Xeon Platinum 8180GeForce GTX 1080TiTesla V100GeForce RTX 2080Ti
倍精度(64bit)1.12 TFLOPS0.33TFLOPS6.38TFLOPS0.37TFLOPS
単精度(32bit)2.24 TFLOPS10.61TFLOPS12.75TFLOPS11.75TFLOPS
半精度(16bit)2.24 TFLOPS10.61TFLOPS25.5TFLOPS11.75TFLOPS
消費電力205W250W250W250W
お値段$10009$699$8000~$9000$999

* Turingは未発売のGPUであるため、実際のスペックは異なる可能性があります。
* 定格の能力で比較しています。実際には温度等の環境によって性能が上下します。

これだけを見ると、10~20%しか性能が向上していないのに、しれっと価格は40%増しになっています。

これしか違いがないかというと、さすがにそんな馬鹿なことはなくて、追加の機能としてレイトレーシング向けのRTコアとTensorコアの搭載があります。

逆に、この2つの新機能が必要ないのであれば、発売から時間もたって価格も手ごろなPascal系のGPUを買うのが正解ではないかと思います。

Tensorコアとは?

Tensor:テンソル、とは何でしょうか?Wikipedia先生に聞いてみたところ、

テンソル: tensor, : Tensor)とは、線形的なまたは線形的な幾何概念を一般化したもので、基底を選べば、多次元の配列として表現できるようなものである。

全くをもって意味不明ですね。

要は行列なんですが、行列演算は機械学習の学習/推論の両方で多用されます。

そこで、普通に演算器を大量に積むのではなく、行列演算用のアクセラレータを用意した、そのアクセラレータがTensorコアという感じです。

実はTensorコア自体はVoltaにも入っています。VoltaとTuringのTensorコアを比較してみます。

 Tesla V100GeForce RTX 2080Ti
半精度(16bit)101.99T Tensor FLOPS99.53T Tensor FLOPS
消費電力250W250W
お値段$8000~$9000$999

* Turingは未発売のGPUであるため、実際のスペックは異なる可能性があります。
* 定格の能力で比較しています。実際には温度等の環境によって性能が上下します。

何故半精度だけで比較しているかというと、どうもVoltaのTensorコアは学習よりで、TuringのTensorコアは推論寄りっぽいので、どちらでも使うFP16しか被らないっぽいのです。

自信がなくてすいません、間違っていたら弊社お問い合わせフォームよりご指摘くださいませ。

比較してみて

エンタープライズとゲーム向けという差はあれど、Turingのゲーム向けGPUの圧倒的なコスパの良さが光ります。

残念ながら利用規約によりデータセンターでは使えないものの、データセンター以外で使うならTuringでいいんじゃないかな。。。

単精度演算はVoltaのみのサポートのようですが、機械学習の学習もFP16でやってもそこまで問題があるわけではないですし・・・。

RTコアについて

一応紹介したので、今回初登場のRTコアについてご説明します。

RTコアとは一言で言うと「レイトレーシングアクセラレーター」です。

Nvidiaのデモを見ていただけるとよくわかると思うのですが、今まで難しかった光の反射のリアルタイム描画を可能にする光線処理のアクセラレーターになります。

機械学習に興味がなくても、これだけでも買い替える価値があるやもしれませんし、逆にこの機能にも興味がないとなると、大して性能が違わないので前世代のGeForceでもいいかもしれません。

GPUを使って無線LANをクラックする話・ボトルネック発見編

データセンターで偶然出会った、キーボード・No.1さんです。ナンバーワンですって。

まさか、こんなところでキーボード界の頂点に出会えるとは。偶然てのは、すごいですね。

キーボードを見る目のないわたしには、この方を見てもその凄さ−《王者の風格》、とでも言うべきものでしょうか−そういったものは、感じ取れませんでした。「能ある鷹は爪を隠す」ってやつなんでしょうか。いやぁ、御見逸れ致しました。わたくしにはキーボードの才能が無いのでしょう。わたしはせいぜい、指でキーボードぺちぺちする身で一生甘んじていようと思います。

キーボードのNo.1ってどうやって決めるんだろうとか、次の世界大会はいつなのかとか、聞けば良かったな。サインももらえば良かったかも。

…えー、もちろん嘘で、ただの通し番号です。データセンターにはたくさん共用のキーボードがあって空いてるものを番号の小さいものから貸してもらえるんですが、同時に作業する人は居ないのでいつもだいたいこのキーボードを貸してもらって作業しています。でも不思議かな、みんなこのキーボードばっかり使ってるはずなのに、「キーボードNo.10」とかの方がボロボロな感じがします。…もしかして、やっぱ世界1位《ザ・トップ・オブ・キーボード》だったのか?

いや…まさかね。

前回までのPyritなのですよ

えー、何の話でしたっけ。あーそうそう。無線LANのパスワード認証方式WPA2-PSKをPyritでクラックしてたんだけど、WPA2の仕様書が読めないから諦めて直接Pyritのソースを読んでいたんでしたね。

えー、何の話でしたっけ。あーそうそう。無線LANのパスワード認証方式WPA2-PSKをPyritでクラックしてたんだけど、WPA2の仕様書が読めないから諦めて直接Pyritのソースを読んでいたんでしたね。

簡単におさらいです:

  • IteratorCrackerというのがいて
  • 辞書ファイルから読んだパスワードのリストを、Iteratorがresultsにして
  • Crackerがresultsを検証して「そのパスワード、あたりっ!」「はずれっ!」と判定する
  • CUDAを使っているのはIteratorだけ。CrackerはCPUでしか動かない。

というのを前回ソース読んで確認したんでした。

そもそも何を測っているのかを把握するのですよ

太古の昔に書いた最初の記事では、Pyritの「benchmark」というコマンドを実行してCPUと比較してみたり、GPU同士で比較して「Tesla V100はGTX1080よりめっちゃ高くて性能いいはずなのに大差ない…」とか嘆いていたわけですが、そもそもこれは何を測っていたのでしょうか。前回多少なりともソースコードを読んだ事により、このような多少深まった疑問が生まれるわけですな。前回見た限り、Crackerも二種類あって、データやコマンドラインフラグによって変わるんですよね。その辺の条件は?

…というわけで、早速読んでみましょう。pyritコマンドのサブコマンドの実装は全部ルートにあるpyrit_cli.pyに置いてありまして、コマンド名と関数名はすべて一致しています。つまり、benchmarkコマンドもbenchmark関数に置いてあります。大変わかりやすい。

ちょっと長いですが、多少はしょりつつ載せます:

    def benchmark(self, timeout=45, calibrate=10):
# ... 略 ...
# ダミーの仕事をなげまくる仕事
            t = time.time()
            perfcounter = cpyrit.util.PerformanceCounter(timeout + 5)
            while time.time() - t < timeout:
                pws = ["barbarbar%s" % random.random() for i in xrange(bsize)]
                cp.enqueue('foo', pws)
                r = cp.dequeue(block=False)
# ... 略 ...
# CPUの結果を表示する
            for i, core in enumerate(cp.cores):
                if core.compTime > 0:
                    perf = core.resCount / core.compTime
                else:
                    perf = 0
                if core.callCount > 0 and perf > 0:
                    rtt = (core.resCount / core.callCount) / perf
                else:
                    rtt = 0
                self.tell("#%i: '%s': %.1f PMKs/s (RTT %.1f)" % \
                            (i + 1, core.name, perf, rtt))
# ... 略 ...
# GPUの結果を表示する
            for i, CD in enumerate(cp.CUDAs):
                if CD.compTime > 0:
                    perf = CD.resCount / CD.compTime
                else:
                    perf = 0
                if CD.callCount > 0 and perf > 0:
                    rtt = (CD.resCount / CD.callCount) / perf
                else:
                    rtt = 0
                self.tell("#%i: '%s': %.1f PMKs/s (RTT %.1f)" % \
                          (i + 1, CD.name, perf, rtt))
# ... 略 ...

cpyrit.cpyrit.CPyrit() as cp というのはIteratorが依存している、パスワードを入れるとresultsというリストが出て来るモジュールでした(で、これをCrackerがチェックする)。

すぐに気がつくのは、CPUやCUDAごとにベンチマークしてるわけではないことです。cp.enqueueとして「パスワード計算してくれや」とお願いすると、cpくんが代理店となって、実際にCPU、CUDA(、OpenCL)のどれで計算するかを決めるようになっています。で、あとで集計するときにcpくんに「ところでTesla使ったのってどれくらい?」って聞いて、その結果をベンチの結果としてまとめる、と。

うーん。APIとしてはバックエンドが隠蔽されてて、いい感じだと思います。でも、これで本当に「公平」にベンチマークできているのかは、若干怪しい感じがするんですが…。

とりあえず。前回、「CPUだけで走るCrackerがボトルネックとなってbenchmarkでTeslaがあんまり速くなかったのでは?」という仮説を立てましたが、それは否定されました。だって、見ての通り、Crackerは走らないからです。もちろん、「このbenchmarkでは」であって、実際にパスワードを検索する時はCrackerも走らせて正しいパスワードかどうかチェックするので、Crackerがボトルネックとなる可能性は残ります。

綺麗なCUDAプロファイラーでも眺めるのですよ

なんか手詰まり感があるので、ここで一旦ソースコードの調査を打ち切ってプロファイラーで遊んでみましょう。

というのも。

ベンチマーク取ってる間、nvidia-smiコマンドを打つとこんな感じで消費電力が見れるわけですが:

+-----------------------------------------------------------------------------+
| NVIDIA-SMI 390.30                 Driver Version: 390.30                    |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  Tesla V100-PCIE...  Off  | 00000000:3D:00.0 Off |                    0 |
| N/A   55C    P0    43W / 250W |      0MiB / 16160MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   1  Tesla V100-PCIE...  Off  | 00000000:3E:00.0 Off |                    0 |
| N/A   57C    P0    45W / 250W |      0MiB / 16160MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   2  GeForce GTX 108...  Off  | 00000000:B1:00.0 Off |                  N/A |
| 28%   50C    P0    60W / 250W |      0MiB / 11178MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   3  GeForce GTX 108...  Off  | 00000000:B2:00.0 Off |                  N/A |
| 32%   56C    P0    57W / 250W |      0MiB / 11178MiB |      2%      Default |
+-------------------------------+----------------------+----------------------+

+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID   Type   Process name                             Usage      |
|=============================================================================|
|  No running processes found                                                 |
+-----------------------------------------------------------------------------+

pyrit benchmarkを実行している間も、このアイドル時の画面と正直消費電力があんまり変わらなくて、ごくごくたまにGPUのどれか1つがちょっと電気を使うようになって、また戻る…そんな感じなんですよ。

+-----------------------------------------------------------------------------+
| NVIDIA-SMI 390.30                 Driver Version: 390.30                    |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  Tesla V100-PCIE...  Off  | 00000000:3D:00.0 Off |                    0 |
| N/A   54C    P0    48W / 250W |    427MiB / 16160MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   1  Tesla V100-PCIE...  Off  | 00000000:3E:00.0 Off |                    0 |
| N/A   56C    P0    51W / 250W |    427MiB / 16160MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+
|   2  GeForce GTX 108...  Off  | 00000000:B1:00.0 Off |                  N/A |
| 44%   74C    P2   198W / 250W |    167MiB / 11178MiB |     36%      Default | ←一時的に使用率が上がった(でもたった36%??)
+-------------------------------+----------------------+----------------------+
|   3  GeForce GTX 108...  Off  | 00000000:B2:00.0 Off |                  N/A |
| 48%   79C    P2    85W / 250W |    161MiB / 11178MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+

+-----------------------------------------------------------------------------+
| Processes:                                                       GPU Memory |
|  GPU       PID   Type   Process name                             Usage      |
|=============================================================================|
|    0     96620      C   /home/psi/src/Pyrit/.env/bin/python          416MiB |
|    1     96620      C   /home/psi/src/Pyrit/.env/bin/python          416MiB |
|    2     96620      C   /home/psi/src/Pyrit/.env/bin/python          157MiB |
|    3     96620      C   /home/psi/src/Pyrit/.env/bin/python          151MiB |
+-----------------------------------------------------------------------------+

CPU(htop)もこんな感じで全コア使い切ってない感が否めない…。

でも、それっておかしいですよねぇ?上で読んだように、CPUやGPUごとじゃあなくて、全部に対して仕事を投げまくってるので、理想的な状況としては全GPUと全CPUがTDPギリギリに張り付いてないとおかしいはずなんです。で、せっかく静かにしたファンが全力で回るようになると。ちなみにわたくしリモートで勤務するバーチャル実在する架空のプログラマですので、音のほうは確かめられません。フヒヒ。

というわけで、まずは本当にGPUが遊んでいるのかどうか、プロファイラーで調べられないか試してみます。

PythonとC言語とCUDAが組み合わさった現代のバベルの塔としか言いようがないこの複雑極まりないシステムですが、DeepLearningで非常によくある組み合わせだからかツール自体はかなり整備されておりまして(そしてDeepLearningの場合はさらに統計的機械学習の難しい理論や巨大なデータがその塔の上に積み上がっていく)、こんな感じで非常に簡単にプロファイルできるようになっています:

# nvidia-profilerは別パッケージらしい
sudo apt install nvidia-profiler

# 測定したいコマンドの前にnvprof -o <出力>.nvvpをつけるだけ。すっごいかんたん!
nvprof -o profile.nvvp pyrit benchmark_long

しかもこのprofile.nvpの結果を見るためのVisual Profilerは、なんとmacでも動きます!リモートのLinuxでももちろん動きます!GPUがあってもなくても!…資本主義ってこえぇなぁ。

nVIDIAのサイトからCUDAのmac版をインストールして、DriverとToolkitとSampleのうち、Toolkitだけインストールします。DriverもインストールできるってことはGeForce積んでるmacもあるんだな…。

で、/Developer/NVIDIA/CUDA-9.2/bin/nvvp (macの場合。CUDAバージョンによって違う)を実行して、openからscpか何かで持ってきたprofile.nvvpを開きます。結果はこんな感じになりましたとさ:

めっちゃスカスカやんけー!

この茶色とか青い棒が出てるところが「計算したで」「CPUとGPUの間でデータのやりとりしたで」を表しているので、このスカスカ具合は、…要するに「GPU全然使えてません」ということになります。つまり、nvidia-smiから雑に考えた仮説は正しかったってこと。

もうすこし拡大して見て行くと、だいたい1秒に一回ぐらいGPUが叩き起こされて、0.06〜0.08秒だけ起こされてまた寝てます。正直ちょっと羨ましい!…いや、そんな頻繁にちょっとだけ起こされるのもやだな。

最初、GPUとCPUの間でパスワードのデータをやりとりするところがボトルネックになるのかなぁ、と予想したりもしていましたが、それは否定されました。計算している80ミリ秒とかに対して、データの転送は0.1ミリ秒程度です。誤差。

ベンチというのは分母と分子が大事なのですよ

まだあわてるような時間じゃない

仮にほとんどGPUを使っていないとしても、ベンチマークの値の信憑性にただちに影響があるわけではありません。

元のソースコードに戻りましょう:

            for i, CD in enumerate(cp.CUDAs):
                if CD.compTime > 0:
                    perf = CD.resCount / CD.compTime
                else:
                    perf = 0

「ベンチマーク」と英語で横文字を使ってみるとなんとなく知的でイケてる感じですが、要するに「仕事率」を測っておるわけです。「やった仕事」を「掛かった時間」で割ってるだけ。

この「掛かった時間」が実際にはどこからどこまでなのかが問題なわけです。

仮に、GPUが1秒間に80ミリ秒しか仕事してないとしても、「掛かった時間」も80ミリ秒なら、その「ベンチマーク」は「GPUの計算の仕事率」を測れていることになります。もし仮にそうなら、GTX1080とTeslaはパスワードクラックという仕事では、そんなに実力に差がないという事になるでしょうね。Pyritに対しては「もっとGPUに仕事割り当ててあげて」と言う事になります。

じゃあ、そうじゃなかったら?…それは、GPUを働かせるために必要な他の仕事の方にばかり時間を取られていて、GPUをうまく動かせていない、という結論が出てきます。例えていうなら、30分の大学院のミーティングに出るために2時間電車乗ってるような。…はぁ…。もし仮にそうなら、ベンチマークの結果はGPUそのものの性能差はあまり反映できていない、という事になるでしょうね。もっといえば、その「余分な仕事」はCPUを使ってるはずなので、このCPUとGPUを同時に走らせるベンチマークは、公平でないことになります。

ソースコードをちょっと書き換えて、この分母であるCD.compTime も表示するようにしましょう。

した結果がこんな感じ:

#1: 'CPU-Core (SSE2/AES)': 528.7 PMKs/s (RTT 2.7) compTime=24
#2: 'CPU-Core (SSE2/AES)': 585.8 PMKs/s (RTT 2.8) compTime=31
#3: 'CPU-Core (SSE2/AES)': 559.6 PMKs/s (RTT 2.9) compTime=35
#4: 'CPU-Core (SSE2/AES)': 537.1 PMKs/s (RTT 2.7) compTime=32
#5: 'CPU-Core (SSE2/AES)': 566.9 PMKs/s (RTT 3.0) compTime=27
#6: 'CPU-Core (SSE2/AES)': 488.4 PMKs/s (RTT 3.0) compTime=33
#7: 'CPU-Core (SSE2/AES)': 569.0 PMKs/s (RTT 3.1) compTime=31
#8: 'CPU-Core (SSE2/AES)': 506.9 PMKs/s (RTT 2.9) compTime=29
#9: 'CPU-Core (SSE2/AES)': 520.3 PMKs/s (RTT 2.8) compTime=19
#10: 'CPU-Core (SSE2/AES)': 579.7 PMKs/s (RTT 2.8) compTime=28
#11: 'CPU-Core (SSE2/AES)': 543.4 PMKs/s (RTT 3.1) compTime=31
#12: 'CPU-Core (SSE2/AES)': 580.9 PMKs/s (RTT 3.0) compTime=33
#13: 'CPU-Core (SSE2/AES)': 569.9 PMKs/s (RTT 2.9) compTime=37
#14: 'CPU-Core (SSE2/AES)': 496.3 PMKs/s (RTT 2.9) compTime=23
#15: 'CPU-Core (SSE2/AES)': 546.1 PMKs/s (RTT 2.8) compTime=34
#16: 'CPU-Core (SSE2/AES)': 564.5 PMKs/s (RTT 2.8) compTime=36
#17: 'CPU-Core (SSE2/AES)': 538.7 PMKs/s (RTT 2.9) compTime=31
#18: 'CPU-Core (SSE2/AES)': 621.0 PMKs/s (RTT 2.8) compTime=25
#19: 'CPU-Core (SSE2/AES)': 527.8 PMKs/s (RTT 2.8) compTime=31
#20: 'CPU-Core (SSE2/AES)': 492.7 PMKs/s (RTT 2.8) compTime=34
#21: 'CPU-Core (SSE2/AES)': 532.0 PMKs/s (RTT 2.8) compTime=28
#22: 'CPU-Core (SSE2/AES)': 558.4 PMKs/s (RTT 2.9) compTime=29
#23: 'CPU-Core (SSE2/AES)': 548.8 PMKs/s (RTT 2.8) compTime=31
#24: 'CPU-Core (SSE2/AES)': 528.6 PMKs/s (RTT 2.9) compTime=32
#25: 'CPU-Core (SSE2/AES)': 594.7 PMKs/s (RTT 2.8) compTime=28
#26: 'CPU-Core (SSE2/AES)': 508.1 PMKs/s (RTT 2.8) compTime=33
#27: 'CPU-Core (SSE2/AES)': 534.2 PMKs/s (RTT 2.9) compTime=32
#28: 'CPU-Core (SSE2/AES)': 545.9 PMKs/s (RTT 3.1) compTime=31
#29: 'CPU-Core (SSE2/AES)': 543.0 PMKs/s (RTT 3.0) compTime=39
#30: 'CPU-Core (SSE2/AES)': 550.1 PMKs/s (RTT 2.6) compTime=29
#31: 'CPU-Core (SSE2/AES)': 595.8 PMKs/s (RTT 3.1) compTime=31
#32: 'CPU-Core (SSE2/AES)': 581.9 PMKs/s (RTT 2.7) compTime=30
#33: 'CPU-Core (SSE2/AES)': 537.9 PMKs/s (RTT 2.9) compTime=29
#34: 'CPU-Core (SSE2/AES)': 566.2 PMKs/s (RTT 3.0) compTime=30
#35: 'CPU-Core (SSE2/AES)': 563.1 PMKs/s (RTT 2.6) compTime=24
#36: 'CPU-Core (SSE2/AES)': 590.0 PMKs/s (RTT 2.7) compTime=32
#37: 'CPU-Core (SSE2/AES)': 581.8 PMKs/s (RTT 2.7) compTime=27
#38: 'CPU-Core (SSE2/AES)': 566.3 PMKs/s (RTT 2.7) compTime=33
#39: 'CPU-Core (SSE2/AES)': 516.1 PMKs/s (RTT 3.2) compTime=28
#40: 'CPU-Core (SSE2/AES)': 561.8 PMKs/s (RTT 2.9) compTime=35
#41: 'CPU-Core (SSE2/AES)': 542.7 PMKs/s (RTT 2.9) compTime=29
#42: 'CPU-Core (SSE2/AES)': 577.9 PMKs/s (RTT 2.9) compTime=29
#43: 'CPU-Core (SSE2/AES)': 586.3 PMKs/s (RTT 3.0) compTime=33
#44: 'CPU-Core (SSE2/AES)': 523.4 PMKs/s (RTT 3.0) compTime=30
#45: 'CPU-Core (SSE2/AES)': 526.7 PMKs/s (RTT 2.9) compTime=32
#46: 'CPU-Core (SSE2/AES)': 505.5 PMKs/s (RTT 2.9) compTime=32
#47: 'CPU-Core (SSE2/AES)': 567.8 PMKs/s (RTT 2.7) compTime=24
#48: 'CPU-Core (SSE2/AES)': 521.5 PMKs/s (RTT 2.9) compTime=32
CUDA:
#1: 'CUDA-Device #1 'Tesla V100-PCIE-16GB'': 113061.9 PMKs/s (RTT 0.4) compTime=11.3
#2: 'CUDA-Device #2 'Tesla V100-PCIE-16GB'': 120956.8 PMKs/s (RTT 0.4) compTime=13.1
#3: 'CUDA-Device #3 'GeForce GTX 1080 Ti'': 94154.2 PMKs/s (RTT 0.4) compTime=13.3
#4: 'CUDA-Device #4 'GeForce GTX 1080 Ti'': 96664.7 PMKs/s (RTT 0.5) compTime=16.4

45秒のベンチに対して、10秒程度。たしかにCPUに比べてかなり短いですが、1秒中80ミリ秒、つまり8%しか使ってないことを考えるとあまりに長いです。Pyritくんが口ではGPUで実行している!といいつつ、実はそうでない部分がかなり占めていて、GPUを全然使い切れていない可能性はやはりかなり高い、と考えてよいでしょう。

もう一つ。CPUもGPUも、どっちも45秒のベンチマークのうち、けっこうの部分遊んでいます。htopであんまりコアを使い切ってない気がしたのもここで説明がつきました。これはどう解釈すればいいのかといえば…たぶん、ダミーのパスワードの生成が、追いついてないんじゃないでしょうか。CPUやGPUがすぐパスワードを処理して「もっとくれや!」って言うけれど、

["barbarbar%s" % random.random() for i in xrange(bsize)]
としてパスワードをどんどん作る処理が追いついてない。

月刊Pyrit、今月のまとめなのですよ

  • benchmarkはパスワードクラックの一部の処理しか測ってない
  • GPUもCPUも遊びまくっている
  • nVIDIAのプロファイラきれい(明らかにEclipse製)

さぁて、来月の月刊Pyritは何なのですよ?

いきあたりばったりで書いてるのでどうなるかはわかんないんですけど、このボトルネックたちと遊んでいきましょう。いいっすね、なんかやっと技術ブログっぽくなってきた。…台風の低気圧で今日はテンションがひくいのですよ。みなさまも強い風に飛ばされないように気をつけて〜。

そう言えば先月書いた正解のパスワード、検閲されちゃいましたね。「xxxxxxxxxxx」ですって、いったいどんな文字列だったんでしょう、前後の文脈から想像をかき立てられちゃいますねぇ。「えっちなことば」って昔こんな感じで生まれたのかな。フヒヒ。

GPUサーバーを静かにする話

はじめに

弊社にテスト用のGPUマシンがいらっしゃったわけですが、この評判がすこぶる・・・悪い。

とてーもお高いサーバーであるはずなのに、とても評判が悪い。

このサーバーはNVIDIAさんのライセンスのおかげでデータセンターに置けずオフィスに置いてあるのです。

があまりにもうるさすぎて、執務スペースを追い出されて廊下というか玄関?に置いてあります。

今回はGPUサーバーをどうにかしてみんなに受け入れてもらった話をしましょう。

現状把握をしよう

騒音を計測してみましょう。

1.起動した瞬間

BAKUON!

起動時はファンが前回になるため一番うるさいです。にしても100デシベルと言えば、自動車のクランクション等に相当し、聴覚機能に異常をきたすレベルだそうな。起動後落ち着いた状態(アイドル時)

2. 起動後落ち着いた状態(アイドル時)

騒々しい街頭や掃除機に相当するレベルだそうな。

3. 起動後落ち着いた状態@会議室

扉を1枚隔てた会議室での測定結果。会議室も会議ができないほどではないですが、余裕で音が聞こえます。

4. 執務スペースの社長の机

執務スペースは扉1枚隔てているものの、余裕で聞こえる、というかうるさい。社長の席がGPUサーバーに一番近いので、私の席だともう少し静か。60デシベルで、時速40kmで走行する車の車内に相当する騒音だそうな。

5. 起動後落ち着いた状態@オフィスの外の廊下

外でも余裕で聞こえる。これはそのうちビルの管理会社に怒られそう。

6. 起動後落ち着いた状態@オフィスの外のエレベーターホール

エレベーターホールまで行けばあまり気にならない。

が相変わらず聞こえる。

やはりなんとかせねば。

サーバーの設定を見直して静かにしよう

ところがこのサーバー、BIOSやUEFIにFANに関する設定がないんだなー。

じゃあどうするんだというとIPMIから設定する感じ。

サーバーらしく、大量のセンサーが存在。これらのセンサー値とファンを連動させることができます。

初期設定は最初から50%ぐらいの力でファンが回っていて、負荷がかかるとすぐに75%ぐらいまで上がるように設定されています。

が、廃棄はかなり涼しく、正直過剰冷却感否めず。

そこで以下のように変更してみます。

ある瞬間までは頑張って耐えますが、温度が上がってくると突然ファンを全開にして温度を下げる方針です。

本当はもう少し緩やかにファン回転数を上げて言ってもいいんですが、熱で壊れたりしたら嫌なのでコンサバに行きます。

さて結果はどうなりましたでしょうか。

結果発表

起動した瞬間はどちらも全力で動くので割愛します。

起動後落ち着いた状態

おおお、劇的に静かになりました。

起動後落ち着いた状態@会議室

ここまで来るとほぼ聞こえません。

起動後落ち着いた状態@オフィス街の廊下

こちらもほとんど聞こえません。というかサーバー自体より騒音値が上なので、廊下自体の騒音の方がうるさい感じです。

エレベーターホールは聞こえなかったのと、執務室は空気清浄機の音の方が支配的だったので割愛します。

番外編

弊社のデータセンターにも同じ機種のGPUがないものがあったので、ファンのチューニングをしてみました。

騒音はデータセンター自体元々うるさいので測定していませんが、アイドル時の消費電力が420W→160Wに減少しました。。。この機種の初期設定がひどすぎる説・・・。

結論

サーバーも設定次第で静かにすることが可能です。諦めずにチューニングしましょう。

GPUを使って無線LANをクラックする話・読解編・序

<VTuber語法>
はいどうもー!実在する架空のバーチャルプログラマの平藤でーす。のじゃのじゃー。

世の中世知辛い!

あの〜本職(?)が〜データセンターおじさんなんですけど〜、
無線LANのクラックやってもお金にならない!やってもあんま意味ない!(社会的に)

え〜そんなことよりも〜はるかに〜あの〜
データセンターとかで〜
重いサーバー二人掛かりで設置して腰痛めたり〜
あの〜機嫌が悪いとパケットが通らなくなるルータくんの設定したり〜
あのぉ…新しいサーバくんにSSDを刺したのに無視されたりすることなどのほうがぁ…
人生においてぇ…重要なことぉ…なのじゃあ。のじゃーーーーー。
</VTuber語法>

えー、何の話でしたっけ。あーそうだ。WPA2のパスワードをクラックしてたんでした。今回はその続きをやって、実際にクラックするのを目指していきましょう。

写真は今年の紫陽花です。

そういえば。そろそろ紫陽花の季節も終わりですけど、紫陽花の「花」って散らないんですよ。知ってました?…「フォトジェニック」ではないからか、街中の紫陽花の「花」はだいたい人間に刈り取られちゃうんですが。せちがらーーーーい!

WPA2でこっそり堂々と鍵を交換するヒミツ:有料

前回がんばってスマホが無線LANルータと「つながる」までのパケットをキャプチャーしたんでしたね。

今回は、このパケットのうち、「認証」をするために送りあっているKeyメッセージの部分をWPA2の仕様に沿ってもう少し詳しく読み解きましょう。

…うーん…。

…うーん…??

よーし、パパ気の効いた解説でも書いちゃうぞー、と思っていたのですが、ネット上に参考になる記事がありません。しょうがないにゃあ・・、(元の仕様書読むから)いいよ。

…と思ったところ、WPA2はIEEE802.11iという規格が元になっているらしいというところまでは突き止めたですが、IEEE802.11iの仕様書はなんと有料らしいのです。あと最近出たWPA3の仕様書ダウンロードフォームを発見したのですが(WPA2はどこだよ??)、なんか簡単にダウンロードできる気配じゃない!

世の中WiFiデバイスばっかなのに仕様書が地味にどこにあるのかわからない!あっても有料!

世の中、世知辛いのじゃー!

秘伝のソースからWPA2を理解するプログラマーの会

…えー。というわけで、作戦を変更して、逆にPyritのソースコードを解読してWPA2の「パスワード認証」に迫っていこうと思います。

えっ、そんな場当たり的でいいの?なんて声が聞こえてきたような気がしましたが、いいんです。無線LANのハックそのものが目的ではないですし(「目的」は強いて言うならGPUと遊ぶベンチマーク取る事だったんだよなぁ)、Pyritのソースで遊んでいるわけですから、WPA2の仕様書よりPyritのソースと仲良くなる方がきっとこの先楽しいと思います。たぶん。

雑に状況を要約:ゼアイズパスワードオンリー、合言葉しかない

こうなると我々には頼るべきものはもう何もありません。我々に残されたのは、パケットダンプとPyritのソース、そしてWiFiデバイスのみ。…まぁ、よくあるシチュエーションだなぁ。

無線LANルータに接続するために必要な情報は、WPA2-PSKを使っている場合、パスワードだけです。それ以外には正確に何もいりません。これが漏れるとオシマイです。FreeWiFiになります。

何を今更、当たり前なことを何度も繰り返しやがって、と思われるかもしれませんが、これはもう一度じっくり観察しておいてもよい事実だと思います。

WiFiの仕様を決めているWiFiアライアンスはいわゆるライセンス商法をやっておりまして、WiFiアライアンスの言う通りに製品を作って、WiFiアライアンスの試験に合格しないと「WiFi」は名乗れないようになっています。そうやって縛ることで互換性を高めて「とりあえず『WiFi』ってシール貼ってあるやつ買えば無線でネット繋がるんでしょ?」って消費者に安心して買ってもらおうとしてるわけですね。USBもまぁだいたい似たような感じなんですが、USBは最近USB Type-C™ Authenticationというのが入りまして、USBでお互いに通信する前に、相手の機械が本当にUSBの業界団体から認証をもらっているのかどうかのチェックをする事が可能になったらしいのです。今までは「USBのロゴを使ってよい」という商標の仕組みを使って許可してれば十分だったのが、それを守らない人が出てきたからそうなったんでしょうね〜。

さて。もし仮に、WiFiにもUSB Authenticationのような仕組みが導入されていた場合、Pyritはクラックする時に、パスワードだけでなく、WiFiアライアンスからもらった証明書も考慮してクラックしないといけなかったりするかもしれません(そうじゃないかもしれません)。そういうややこしい仕組みはない、パスワードだけ考えればいい、そういう事を確認したかったのです。なんで仕様書もソースも読まずにそんな事を断言できるのかというと、…WiFiモジュールの一切入っていないGPUサーバでPyritが動いてクラックできているからです。えぇ、我々にはもはや…こんな推論しか、できる事はないのですよ。

言い換えると、WPA2-PSKで認証できたかどうかのチェックをする関数は、次のような関数シグネチャに集約される事が予想されます:

// あくまで予想。
bool checkWPA2Password(byte[] captureData, string password);

うーん、とってもかんたん。もちろん、パケットキャプチャデータの中からどのデータをどう使うかが大事なわけですし、もっと言えば並列化してるのでもっともっと複雑になっていて、関数1つに纏まっているはずはないと思います。が、こういうのを多少なりとも頭に浮かべながら読むのって大事だと思うんよ。

とりあえず正解のデータを突っ込んで「解けた」と言ってもらう

Pyritのソースもそれなりにデカいので、まず一旦「解けた」と言うコードパスを走ってもらいましょう。言い換えると、正解を突っ込んで「クラックできました!」とPyritに叫んでもらいましょう。マッチポンプ…はちょっと違うか。

えーと。パスワードの正解をズバリ言ってしまうと「xxxxxxxxx」です。今すぐダウンロー
。データセンターに置いてあるWiFiのパスワードもこれです。

(ブログを検閲する偉い人によって文章の一部が改変されました)

…あーっ、ちょっと、悪用しないでくださいよ。まぁ「おらけんま」ができなくなった今、わざわざオランダヒルズの面倒な警備を乗り越えてオフィスにFreeWiFiしにやってくる輩も居ないだろうと思って書いてるわけですが。FreeWiFiなんて今日び、駅からファミレスからコンビニ、スタバまで、そこらじゅうにあるからね。…あとは隣のオフィスの人がこれを読んで無いのを祈るだけだな。

さて。この正解のパスワードをPyritに教えて「解けた!!!」って叫んでもらうには次の呪文を唱えればよいです(zoi.capはリポジトリに入っています):

~/src/Pyrit$ pyrit -r zoi.cap -i passwords.txt -b '00:1a:eb:ac:51:30' --all-handshakes attack_passthrough
Pyrit 0.5.1 (C) 2008-2011 Lukas Lueg - 2015 John Mora
https://github.com/JPaulMora/Pyrit
This code is distributed under the GNU General Public License v3+

Parsing file 'zoi.cap' (1/1)...
Parsed 4055 packets (4055 802.11-packets), got 104 AP(s)

Attacking 4 handshake(s).
Tried 1 PMKs so far; 0 PMKs per second. xxxxxxxxx

The password is 'xxxxxxxxx'.

passwords.txtにパスワードの羅列が書いてあります。今回は正解の1行だけを入力しました:

$ cat passwords.txt
xxxxxxxxx

しかし。そんな雑なパスワードでいいのかよ!という声がいかにも聞こえてきそうですが、Pyritも使っている辞書攻撃はだいたい英語を前提にしているので、ローマ字の日本語と英語を混ぜたこのパスワードは意外と当てられないんじゃ無いかなぁというのが個人的な見解です。

Pyritがデフォルトで持ってる辞書がtest/dict.gz以下にあるのですが、例えばそれでは破れません:

$ pyrit -r zoi.cap -i test/dict.gz -b '00:1a:eb:ac:51:30' --all-handshakes attack_passthrough
Pyrit 0.5.1 (C) 2008-2011 Lukas Lueg - 2015 John Mora
https://github.com/JPaulMora/Pyrit
This code is distributed under the GNU General Public License v3+

Parsing file 'zoi.cap' (1/1)...
Parsed 4055 packets (4055 802.11-packets), got 104 AP(s)

Attacking 4 handshake(s).
Tried 4094 PMKs so far; 1552 PMKs per second. abrogate

Password was not found.

あとWiFiって近くに居ないと繋げないですから、攻撃を受ける回数はネット越しで攻撃できるsshと比べれば段違いに少ないはずですし、「ネットをつなぐための鍵」なのでできるだけ覚えやすい方が嬉しいかなぁと思ってこれに設定してます。スマホの暗証番号が4桁なのとだいたい同じ理論武装です。

嘘です。ジャバがすき(きらい)だからで、残りは全部後付け。Javaとは何で、なぜ必要ですか。

…はい。

WiFiクラックって(今まさにやってるような)スクリプトキディ的ハッキングの題材としてはありがちで、ネット上でも雑誌上でもよく見るから、みんな好きなんでしょうけど(Pyritクラスター作ってるやつらもおるぞ)、「WiFiのパスワードってそこまでクラックする価値あるの?sshとかTLSの穴探した方が『コスパ』いいんじゃない?」みたいなのが正直な感想ですな。己の記事全否定か〜?って感じですが、いや、まぁ、今回やってるのはベンチマークであって、そういう「価値」とはまた違う話ですから。…それに、公開IPにやってくる大量のsshログイン試行ログを見ていると、明らかに「業務」「ビジネス」としてやってる感じが見て取れて、それはそれで、気が滅入るのですよ。その点、Pyritクラスターは見てるだけでも楽しそうでしょ?

ソースコードの海に飛び込むのです

だいぶ話が逸れました。Pyritがパスワードを発見するのに成功した場合と失敗した場合のメッセージを得たので、これを足がかりにソースコードを読解していきます。Pyritもそれなりに大きいので、全部読むのは正直大変だなーという判断です。

するとサクッとヒットします。二箇所あるんですが、コマンド名と関数が同じこちらでしょう。

    def attack_passthrough(self, infile, capturefile, essid=None, \
                           bssid=None, outfile=None, all_handshakes=False, \
                           use_aes=False):

# ...

        if not all_handshakes:
            crackers.append(cpyrit.pckttools.AuthCracker(auths[0], use_aes))
        else:
            self.tell("Attacking %i handshake(s)." % (len(auths),))
            for auth in auths:
                crackers.append(cpyrit.pckttools.AuthCracker(auth, use_aes))
        with cpyrit.util.FileWrapper(infile) as reader:
            with cpyrit.cpyrit.PassthroughIterator(essid, reader) as rstiter:
                for results in rstiter:
                    for cracker in crackers:
                        cracker.enqueue(results)
# ...
        for cracker in crackers:
            cracker.join()
            if cracker.solution is not None:
                self.tell("\nThe password is '%s'.\n" % cracker.solution)
                if outfile is not None:
                    with cpyrit.util.FileWrapper(outfile, 'w') as writer:
                        writer.write(cracker.solution)
                break
        else:
            errmsg = "\nPassword was not found."

これの上の方を手繰っていくと、cpyrit.pckttools.AuthCrackerというのがクラッカーをまとめてる親玉らしいことがわかります:

class AuthCracker(object):

    def __init__(self, authentication, use_aes=False):
        self.queue = Queue.Queue(10)
        self.workers = []
        self.solution = None
        if authentication.version == "HMAC_SHA1_AES" \
         and authentication.ccmpframe is not None \
         and use_aes:
            self.cracker = CCMPCrackerThread
        else:
            self.cracker = EAPOLCrackerThread
        for i in xrange(util.ncpus):
            self.workers.append(self.cracker(self.queue, authentication))

    def _getSolution(self):
        if self.solution is None:
            for worker in self.workers:
                if worker.solution is not None:
                    self.solution = worker.solution
                    break

    def enqueue(self, results):
        self.queue.put(results)
        self._getSolution()

    def join(self):
        self.queue.join()
        for worker in self.workers:
            worker.shallStop = True
        self._getSolution()

早速CCMPCrackerThreadとEAPOLCrackerThreadという二種類がいることがわかるわけですが、これをprintfデバッグしてどっちなのか確かめると(こういうのを何の躊躇いもなく「さっ」とやるために前回わざわざ自前でビルドしたんやぞ)、EAPOLCrackerThreadを使っているらしいです。が、–aesってオプションつけてやるとCCMPCrackerThreadを使うようになります。でも、どっちでもちゃんとパスワードを解読できます。…ど、どういう事じゃ…?

ちなみに実装はすぐ上にありまして…

class CrackerThread(threading.Thread):

    def __init__(self, workqueue):
        threading.Thread.__init__(self)
        self.workqueue = workqueue
        self.shallStop = False
        self.solution = None
        self.numSolved = 0
        self.setDaemon(True)
        self.start()

    def run(self):
        while not self.shallStop:
            try:
                results = self.workqueue.get(block=True, timeout=0.5)
            except Queue.Empty:
                pass
            else:
                solution = self.solve(results)
                self.numSolved += len(results)
                if solution:
                    self.solution = solution[0]
                self.workqueue.task_done()


class EAPOLCrackerThread(CrackerThread, _cpyrit_cpu.EAPOLCracker):

    def __init__(self, workqueue, auth):
        CrackerThread.__init__(self, workqueue)
        _cpyrit_cpu.EAPOLCracker.__init__(self, auth.version, auth.pke,
                                          auth.keymic, auth.keymic_frame)


class CCMPCrackerThread(CrackerThread, _cpyrit_cpu.CCMPCracker):

    def __init__(self, workqueue, auth):
        if auth.version != "HMAC_SHA1_AES":
            raise RuntimeError("CCMP-Attack is only possible for " \
                               "HMAC_SHA1_AES-authentications.")
        CrackerThread.__init__(self, workqueue)
        s = str(auth.ccmpframe[scapy.layers.dot11.Dot11WEP])
        msg = s[8:8+6]
        counter = (s[0:2] + s[4:8])[::-1]
        mac = scapy.utils.mac2str(auth.ccmpframe.addr2)
        _cpyrit_cpu.CCMPCracker.__init__(self, auth.pke, msg, mac, counter)

_cpyrit_cpuパッケージは、C言語で書かれたモジュールです。あれっ、CUDAは?まぁいいや。とりあえずここでPythonの外に出る事が確認できました。そして、initに渡されている情報が、それぞれのCrackerで必要な情報らしいことがわかります(そしてそれはEAPoLとCCMPで違うこともわかります)。

それぞれのCrackerはsolve()という(C言語で書かれた)メソッドを持ってて、ここへresultsというのが降ってくるのがわかりますが、このresultsというのは具体的に何なのだろう…と思って”The password is “のところまで戻ってもっかい手繰っていくと、cpyrit.cpyrit.PassThroughIteratorというところへ行き着き、これはさらにその名もCPyritというクラスに処理を投げていることが見て取れます。…そして、こちらにはCUDAとかなんとか書いてありますね。

class CPyrit(object):
    """Enumerates and manages all available hardware resources provided in
       the module and does most of the scheduling-magic.
       The class provides FIFO-scheduling of workunits towards the 'host'
       which can use .enqueue() and corresponding calls to .dequeue().
       Scheduling towards the hardware is provided by _gather(), _scatter() and
       _revoke().
    """

    def __init__(self):
        """Create a new instance that blocks calls to .enqueue() when more than
           the given amount of passwords are currently waiting to be scheduled
           to the hardware.
        """
        self.inqueue = []
        self.outqueue = {}
        self.workunits = []
        self.slices = {}
        self.in_idx = self.out_idx = 0
        self.cores = []
        self.CUDAs = []
        self.OpCL = []
        self.all = []
        self.cv = threading.Condition()

        # CUDA
        if config.cfg['use_CUDA'] == 'true' and 'cpyrit._cpyrit_cuda' in sys.modules and config.cfg['use_OpenCL'] == 'false':

            CUDA = _cpyrit_cuda.listDevices()

            for dev_idx, device in enumerate(CUDA):
                self.CUDAs.append(CUDACore(queue=self, dev_idx=dev_idx))

# ...

        # CPUs
        for i in xrange(util.ncpus):
            self.cores.append(CPUCore(queue=self))

# ...

むっ、じゃあさっきのCrackerThreadっていういかにも「クラックしてます!!!!!」的な空気を纏わせてた連中はCUDAでは走らないってこと?

・・・今日のところは、この辺で勘弁しといてやりますか・・・。

ソース読解のまとめ

まぁとりあえず。CrackerとIteratorの2つの「重い処理」があって、IteratorはCUDAでも処理されているが、CrackerはCPUにしか移譲されてなさそうな所は見えてきました。そしてパスワードのリストを読んでIteratorが作ったresultsをCrackerがsolveして「パスワードあたり!」って叫んだり叫ばなかったりすることも。

これはすごい雑な仮説ですが、もしCUDAのカードをいっぱいぶち込んで5000兆フロップス並列に処理できたとしても、Iteratorは速くなるけど、CPUで動くCrackerがボトルネックになってGPUが遊んでしまうでしょう。すると、Teslaでも大して速くなかった前回の結果を擁護できるかもしれません。Teslaくんは悪く無い、CPUで動いてるCrackerが遅いのが悪いんだ!って。ありそうなボトルネックとしてテキストファイルを読む所や、GPUへパスワードを運ぶメモリ帯域を想像していたんですが、実際にソースを読む事でもう一つ見つけられたわけです。もちろん、ベンチマークとか取ったり実験してみないと、どこがボトルネックかなんてわからんわけですが。

<VTuber語法>
よくわかん事、Pyrit…CUDA…よくわからん…のじゃー(突然戻ってくる狐)

…でも、でもでもでも!
こうやって、こう ちょっとつ、ステップアップする所のぉー

この…CUDAとかC言語とPythonが組み合わさったやつのプロファイルもどうやって取ればいいのかわからないし。
ほんとはこう…いろいろやりたいけど。
(まだ)できないっ!

世知辛い!
世の中世知辛い!
世の中世知辛いのじゃーーー!!!!のじゃー!
ふーんふーんふーん(元ネタを知らないがとりあえず踊っておく狐)
</VTuber語法>

プロファイルの取り方を調べるのは次回までの宿題にしておきます。

GPUを使って無線LANをクラックする話・準備編

<Youtuber語法>ブンブン・ハロー、Link-U。どうも、平藤です。

えーそれでは今日はね、無線LANのパスワードをクラックするPyrit、これのほうを、CPUとかGPUで動かしてみようと思いまーす。</Youtuber語法>

はい。というわけで、弊社の業務用無線LANルータ、SSIDがその名も「zoi」のWPA2-PSKのパスワードをCPUやGPUでクラックすることで、機械学習とはまた違った「アプリケーション」でのパフォーマンスを実測してみたいと思います。

パケットをキャプチャする

ちなみに写真はオフィスにある会議室の机の下です。いつもと見慣れない風景で、ちょっとかっこよくない?…そうでもないか。

なにはなくとも、まずは攻撃するために、ルータがスマホをパスワードで認証して通信を開始する「ハンドシェイク」の瞬間のパケットをキャプチャしましょう。以下、わたしの使っているMacBook Pro + OSX HighSierraでWPA2-PSKのパケットをキャプチャする方法を説明します。Linuxでの方法は他の人に譲ることにします。餅は餅屋なので。

まずは次のコマンドで、どんな電波が飛んでるかをチェックします:

sudo /System/Library/PrivateFrameworks/Apple80211.framework/Versions/Current/Resources/airport -s

結果はこんな感じです:

% sudo
/System/Library/PrivateFrameworks/Apple80211.framework/Versions/Current/Resources/airport -s
Password:
SSID BSSID RSSI CHANNEL HT CC SECURITY (auth/unicast/group)
LU-MOBILE 00:1a:eb:b2:c3:b6 -66 112,-1 Y JP WPA2(PSK/AES,TKIP/TKIP)
PAI-AP2.4G1 34:76:c5:4d:a7:c8 -76 10,-1 Y JP WPA(PSK/AES,TKIP/TKIP) WPA2(PSK/AES,TKIP/TKIP)
PAI-AP2.4G2 06:76:c5:4d:a7:c8 -75 10,-1 Y JP NONE zoi 00:1a:eb:ac:51:20 -44 9 Y JP WPA2(PSK/AES/AES)
LU-MOBILE 00:1a:eb:b2:c3:b5 -54 6 Y JP WPA2(PSK/AES,TKIP/TKIP)
RADIX01 4c:e6:76:ae:00:1e -86 6 Y -- WPA(PSK/AES,TKIP/TKIP) WPA2(PSK/AES,TKIP/TKIP)
Eviry24G c0:25:a2:89:b2:72 -61 7 Y JP WPA2(PSK/AES/AES)
GW-office_G1 34:3d:c4:86:83:d0 -90 3 Y -- WPA2(PSK/AES/AES)
zoi 00:1a:eb:ac:51:30 -54 40 Y JP WPA2(PSK/AES/AES)

まぁ色々と出てきますが、今回注目なのはこの一番最後の行です:

zoi 00:1a:eb:ac:51:30 -54  40      Y  JP WPA2(PSK/AES/AES)

ここから、SSID「zoi」のBSSID(MACアドレス)が「00:1a:eb:ac:51:30」なこと、チャンネルが40chなことが読み取れますので、これをメモっておきます。

おっと、ここで質問をいただきました。

「zoiは上にもチャンネル9のやつがあるけど?」

良い質問です。おっしゃる通りです。いやまぁ今わたしが考えた質問なんですが。

チャンネルは13番までがいわゆる「2.4GHz」帯で、それ以降は「5GHz」帯になります。 2.4GHz帯の方をキャプチャしても、もちろん構いません。当たり前ですが、パスワードはどの周波数でも共有なので、クラックできればどちらも同じパスワードが得られます。ただ、これも当たり前ですが、ハンドシェイクが行われている場面をキャプチャできなければクラックには使えません。見てない家政婦は事件も解決できないってやつです。手元にあるどんな一般的なご家庭にもありそうな最近のスマホは5GHz帯をデフォルトで使うようになっているようでしたので、これからは40chだけに着目していきます。

で、次に、実際にキャプチャするコマンドがこちらになります(en0は内臓無線LANアダプタのデバイス名です):

% sudo /System/Library/PrivateFrameworks/Apple80211.framework/Versions/Current/Resources/airport en0 sniff 40
Password:
Capturing 802.11 frames on en0.
^CSession saved to /tmp/airportSniffGUDQbe.cap.

最後に表示されたのがキャプチャされた結果のファイルです。これをPyritに解析してもらう事になります。

ああ、一つ注意があります。これを使うとMacの無線LANが一切使えなくなって、ネットも使えなくなります。ダウロードとかsshの接続が切れても困らない時にやってください。あと、この理由があるので、MacBook Proだけではキャプチャはできません。スマホか何か他のデバイスが必要です。

あともう一つ注意としては、スマホの「設定」画面でWiFiをon-offしただけではハンドシェイクは行われません。再起動してください。今回は4回ほど再起動しました。…回数に「おまじない」以上の意味はありませんので、ハンドスピナー的に飽きるまで再起動してください。

これももう一つ注意ですが。tcpdumpはWPA2のハンドシェイクが終わった後のパケットしか見れません。つまり、クラックには利用できません。なんでかというと、それはそもそもtcpdumpのキャプチャするレイヤーがWPA2より一つ上なのと、tcpdumpの使ってBarkley Packet Filterの制限もあるような気がするのですが、OSXのカーネルに立ち入るのは今日は流石に勘弁しといてやりましょうぜ。でもフォーマットはtcpdumpの読み書きできるpcapフォーマットです。なんかちょっと不思議な気分。

WPA2-PSKで「おててつなぐ」ところを観察

さて。この無線LANルータ(厳密にはブリッジ)「zoi」は、WPA2-PSKによって隣のオフィスの人や、たまにやってくる「おらけんま」ことテイルゲーターから守られております。このWPA2-PSKを倒すには、まずこの敵のことをよく観察して理解しなければなりません。これはゲームの攻略だけでなく無線LANクラックにおいても同じです。たぶん。というわけで、さきほどキャプチャした実際のやりとりを眺めながら、WPA2-PSKでパスワード認証する仕組みを実際に見てみましょう。

すぐ上に書いた通り、キャプチャしたファイルはtcpdumpで読めますが、細かい情報に関しては表示してくれないようでしたので、GUIでパケットの情報が見れるWiresharkを使います。鮫好きにもやさしいブログなのです。嘘です。

ビーコン

WiFiに接続しようって時に、無線LANのルーターのリストがスマホに表示されますよね?そのためにルーターが虚空に向かって己の存在をシャウトしている様子がこちらです:

上でハイライトされてる行を見ると、SSID=zoiの「ビーコン」フレームであることや、「zoi」がアライド・テレシス製な事も書いてあります。だいたい1/10秒ごとに射出されているようなので、人間の時間スケール的には「常時叫びっぱなし」といったところでしょうか。ちなみに、さらにProbe-requestというメッセージをスマホから叫ぶと、「誰かいるか〜!?」と聞くことができます。機械にとっては、1/10秒っていうのは「永遠」なんでしょうね。

さて、中身を見ていきましょう。300バイトの中に密に詰まったたくさんの情報があって情報の大洪水に流されそうになりますが、淡々と情報を見ていくと「RSN Information」というところにAESと書いてあります。もっというと、PSKというのも下の方に書いてあります。ここで「zoiでつなぐ時はAESで暗号化してPSKで認証するんだぞい!」と叫んでいるわけです。秒間10回も。ええ、無線LANルータの仕事も大変なのですよ。

アソシエーション手続き

無線LANルータを見つけたら「繋がせていただけないでしょうか?」とまずは声を掛ける「アソシエーション手続き」が始まります。ただ、ここでは「繋ぎたいんですが」と申し出るだけです。玄関で言うならピンポンを押して「佐川急便でーす」という行為に近いでしょうか。

秒間10回ビーコンが叫んでいることもあって、これを手動で探すのは難しいです。そんな時は、上の「表示フィルタ」ってところに「wlan.fc.type_subtype eq 11」と打ち込んでみてください

Authentication Request

「佐川急便でーす」に相当するように、SONYのMACアドレス(スマホ)からアライド・テレシスのMACアドレス(zoi)へのメッセージである旨が記載されています。

Authenticationリクエストは、Authenticationリクエストといいつつ、単に名乗るだけです。パスワードに関係する内容をやりとりしたりはしません。もっというと、歴史的事情があって今は形骸化しているらしいです。まさに「あいさつ」ですね。

パスワードをクラックするのには関係ないので一切無視してもいいんですが、この辺も「そういえばそもそも無線LANってどうやって繋がってんだ!?」という見学気分で流れるように見ていきまます。こういう経験の積み重ねがいざという時のネットワークデバッグに…役立つといいね。

Association Request

形骸化したAuthenticationリクエストの代わりに色々な情報が詰まっているのが、Associationリクエストです。パケットを眺めてみると、図のようにデータレートが書かれています。「zoi」は54MBit/secまでしか対応してないらしいです。知らなかった…。

ここには接続するために必要な細々とした情報はたくさん乗っていますが、パスワードとか認証に関係する情報は乗っていません。なので、これも流していきますよ。

Keyメッセージ

次に続く「Key」メッセージがパスワードで認証するために使われるメッセージです。4つあるので、「4-way handshake」と言います。玄関の例えでいうなら、窓から見える佐川の制服、声(先週と同じ人だ!とか)、手元にある再配達表なんかを見て、本当にドアの向こうにいるのが佐川のお兄さんであって、全然関係ない消火器を売る人とかでないことを確信してドアを開けるためのやりとりです。

今はざっくり見ていただくだけで良いのですが、「パスワード」を直接送っているわけではないことを観察しましょう。当たり前です。パスワード、つまり「あいことば」を直接喋ってしまったら、こうやってずっと「聞いて」いれば、すぐ「あいことば」がバレてしまいます。もっと「ふしぎな方法」で、こっそり、ひっそり、でも堂々と、こうやってやりとりしています。

これに関しては今後の記事でもっとくわしく書く予定ですが、お互いに事前に共有したパスワード(PSK)を共有している同士で、この”Nonce”とか”IV”とかなんとかついてる、4つのメッセージを使って「ある計算」をすれば、「あいことば」を直接喋らなくとも、「我々は同じ合言葉を共有しあった仲間だ!」と確信しあえる…らしいことは覚えておきましょう。まぁもっとも、我々はこれからその「連帯感」「信頼感」を全力でぶち壊す側に回るわけですが。ふふっ、いい気味だ。

Pyritをビルドする

ここから話は手元のMacBook Proから、リモートのGPUやCPUのいっぱい積まれたUbuntuサーバへと話が変わります。

次はPyritのビルドです。無線LANクラック道は、長く険しい坂なのです。打ち切りにならない程度に頑張っていきましょう。

Pyritは実はビルドなんかしなくてもaptで入ります。

$ apt search pyrit
ソート中... 完了
全文検索... 完了
pyrit/xenial 0.4.0-5 amd64
  GPGPU-driven WPA/WPA2-PSK key cracker

pyrit-opencl/xenial 0.4.0-1 amd64
  OpenCL extension module for Pyrit

pyrite-publisher/xenial 2.1.1-11 amd64
  Convert html and text documents to palm DOCformat

wifite/xenial,xenial 2.0r85-1 all
  Python script to automate wireless auditing using aircrack-ng tools

が、見てもらえばわかるとおり、OpenCLバージョンは入ってもCUDAバージョンは入らないのと、これから改造して遊ぼう〜とか考えるとビルドする方がいいので、ビルドしていきます。

必要なライブラリのインストール

粛々とインストールしましょう。環境によっては足りないやつとかもあるかもしれませんが、その辺は柔軟にお願いします:

sudo apt install clang libpcap-dev openssl-dev python-virtualenv libpython-all-dev

ソースのclone

今後改造することもあるかもしれないのでわたしの個人リポジトリにforkしてしまいましたので、そこからcloneしてください。

git clone https://github.com/ledyba/Pyrit
cd Pyrit

CPU版のビルド

CPUだけで解析するやつを一旦ビルドしましょう。

そのための呪文は次に置いておくので唱えてください。

# venv環境を作って有効にします。Pyritは2.7じゃないと動かない…。
$ virtualenv .env --python=/usr/bin/python2.7
$ . .env/bin/activate

# Pythonからcapファイルを見るためのライブラリです。
# 古い2.3.3でないと動かないので、バージョンを指定してインストール。
$ pip install "scapy==2.3.3"

# ビルド。
$ python setup.py build

# 己の作ったvenv環境にいることを確認(しないとたまに間違えるので…)
$ which python
/home/psi/src/Pyrit/.env/bin/python

# インストール
$ python setup.py install

…「ビルド方法」って出てくる結果は単純なのにどうしてこうも試行錯誤が必要なんだろう?

ちなみに上記の呪文でMacでもビルドして実行できるはずなのですが、「scapyが見つからない」という謎のエラーメッセージがでるのでMacの人は注意してください(うそつけちゃんとバージョン指定して入れたやんけ)。わたしは今日のところは諦めました。

なにはなくともビルドできたはずですので、次の呪文を唱えてCPUが48個あることを確認します。

$ pyrit list_cores
Pyrit 0.5.1 (C) 2008-2011 Lukas Lueg - 2015 John Mora
https://github.com/JPaulMora/Pyrit
This code is distributed under the GNU General Public License v3+

The following cores seem available...
#1:  'CPU-Core (SSE2/AES)'
#2:  'CPU-Core (SSE2/AES)'
#3:  'CPU-Core (SSE2/AES)'
#4:  'CPU-Core (SSE2/AES)'
#5:  'CPU-Core (SSE2/AES)'
#6:  'CPU-Core (SSE2/AES)'
#7:  'CPU-Core (SSE2/AES)'
#8:  'CPU-Core (SSE2/AES)'
#9:  'CPU-Core (SSE2/AES)'
#10:  'CPU-Core (SSE2/AES)'
#11:  'CPU-Core (SSE2/AES)'
#12:  'CPU-Core (SSE2/AES)'
#13:  'CPU-Core (SSE2/AES)'
#14:  'CPU-Core (SSE2/AES)'
#15:  'CPU-Core (SSE2/AES)'
#16:  'CPU-Core (SSE2/AES)'
#17:  'CPU-Core (SSE2/AES)'
#18:  'CPU-Core (SSE2/AES)'
#19:  'CPU-Core (SSE2/AES)'
#20:  'CPU-Core (SSE2/AES)'
#21:  'CPU-Core (SSE2/AES)'
#22:  'CPU-Core (SSE2/AES)'
#23:  'CPU-Core (SSE2/AES)'
#24:  'CPU-Core (SSE2/AES)'
#25:  'CPU-Core (SSE2/AES)'
#26:  'CPU-Core (SSE2/AES)'
#27:  'CPU-Core (SSE2/AES)'
#28:  'CPU-Core (SSE2/AES)'
#29:  'CPU-Core (SSE2/AES)'
#30:  'CPU-Core (SSE2/AES)'
#31:  'CPU-Core (SSE2/AES)'
#32:  'CPU-Core (SSE2/AES)'
#33:  'CPU-Core (SSE2/AES)'
#34:  'CPU-Core (SSE2/AES)'
#35:  'CPU-Core (SSE2/AES)'
#36:  'CPU-Core (SSE2/AES)'
#37:  'CPU-Core (SSE2/AES)'
#38:  'CPU-Core (SSE2/AES)'
#39:  'CPU-Core (SSE2/AES)'
#40:  'CPU-Core (SSE2/AES)'
#41:  'CPU-Core (SSE2/AES)'
#42:  'CPU-Core (SSE2/AES)'
#43:  'CPU-Core (SSE2/AES)'
#44:  'CPU-Core (SSE2/AES)'
#45:  'CPU-Core (SSE2/AES)'
#46:  'CPU-Core (SSE2/AES)'
#47:  'CPU-Core (SSE2/AES)'
#48:  'CPU-Core (SSE2/AES)'

うーん。壮観かな。

PyritはGPGPUで解析できるのがウリなのですが、CPUの時もご覧の通り、SSE2とAES系の命令を使ってアクセラレーションしてくれるらしいです。

GPU版のビルド

CUDA版とOpenCL版があるのですが、今日はCUDA版をインストールします。

$ cd modules/cpyrit_cuda/
$ python build install

手元の環境では一発でうまくいったのでよかったのですが、うまく行かなかった人は、この何語かもよくわからない本家のIssueを参考に頑張ってください。コマンドが読めればきっとなんとかなる。はず。

GPU付きで実行する

以上でインストールは完了ですが、設定ファイルでCUDAをONにしないとCUDAは有効になりません

$ cat ~/.pyrit/config
default_storage = file://
limit_ncpus = 0
rpc_announce = true
rpc_announce_broadcast = false
rpc_knownclients =
rpc_server = false
use_CUDA = true
use_OpenCL = false
workunit_size = 75000

CUDAを使うために、use_CUDA = trueとしてください。

virtualenv使ってるのに、これだけ$HOME/.pyritにあるのすごい変な感じがするんですが、まぁ、置いておきましょう。

以上で有効にしたら、以前と同じコマンドを打った時にCUDAデバイスが表示されるようになるはずです:

$ pyrit list_cores
Pyrit 0.5.1 (C) 2008-2011 Lukas Lueg - 2015 John Mora
https://github.com/JPaulMora/Pyrit
This code is distributed under the GNU General Public License v3+

The following cores seem available...
#1:  'CPU-Core (SSE2/AES)'
(略)
#48:  'CPU-Core (SSE2/AES)'

The following CUDA GPUs seem aviable...
#1:  'CUDA-Device #1 'Tesla V100-PCIE-16GB''
#2:  'CUDA-Device #2 'Tesla V100-PCIE-16GB''
#3:  'CUDA-Device #3 'GeForce GTX 1080 Ti''
#4:  'CUDA-Device #4 'GeForce GTX 1080 Ti''

いぇーい。

実測する前にベンチマークをする

ここまでだいぶ長かったので、いちばん美味しい解析はまた今度にとっておくことにして、Pyrit付属の「ベンチマーク」を実行して遊んで終わりとしましょう。…美味しいものを最初に食べる派のみなさん、すいません。

$ pyrit benchmark
Pyrit 0.5.1 (C) 2008-2011 Lukas Lueg - 2015 John Mora
https://github.com/JPaulMora/Pyrit
This code is distributed under the GNU General Public License v3+

Running benchmark (134716.6 PMKs/s)... /

Computed 134716.61 PMKs/s total.
#1: 'CPU-Core (SSE2/AES)': 558.2 PMKs/s (RTT 2.8)
#2: 'CPU-Core (SSE2/AES)': 525.5 PMKs/s (RTT 3.1)
#3: 'CPU-Core (SSE2/AES)': 576.1 PMKs/s (RTT 2.8)
#4: 'CPU-Core (SSE2/AES)': 582.1 PMKs/s (RTT 2.9)
#5: 'CPU-Core (SSE2/AES)': 516.2 PMKs/s (RTT 2.9)
#6: 'CPU-Core (SSE2/AES)': 575.0 PMKs/s (RTT 2.9)
#7: 'CPU-Core (SSE2/AES)': 635.2 PMKs/s (RTT 2.8)
#8: 'CPU-Core (SSE2/AES)': 513.1 PMKs/s (RTT 2.9)
#9: 'CPU-Core (SSE2/AES)': 500.3 PMKs/s (RTT 2.8)
#10: 'CPU-Core (SSE2/AES)': 546.0 PMKs/s (RTT 2.7)
#11: 'CPU-Core (SSE2/AES)': 502.4 PMKs/s (RTT 3.1)
#12: 'CPU-Core (SSE2/AES)': 549.2 PMKs/s (RTT 3.1)
#13: 'CPU-Core (SSE2/AES)': 588.1 PMKs/s (RTT 3.0)
#14: 'CPU-Core (SSE2/AES)': 487.7 PMKs/s (RTT 2.9)
#15: 'CPU-Core (SSE2/AES)': 612.9 PMKs/s (RTT 2.9)
#16: 'CPU-Core (SSE2/AES)': 615.7 PMKs/s (RTT 2.7)
#17: 'CPU-Core (SSE2/AES)': 548.0 PMKs/s (RTT 2.9)
#18: 'CPU-Core (SSE2/AES)': 572.3 PMKs/s (RTT 2.9)
#19: 'CPU-Core (SSE2/AES)': 568.8 PMKs/s (RTT 2.9)
#20: 'CPU-Core (SSE2/AES)': 589.1 PMKs/s (RTT 2.8)
#21: 'CPU-Core (SSE2/AES)': 546.5 PMKs/s (RTT 2.9)
#22: 'CPU-Core (SSE2/AES)': 570.1 PMKs/s (RTT 2.9)
#23: 'CPU-Core (SSE2/AES)': 552.7 PMKs/s (RTT 3.1)
#24: 'CPU-Core (SSE2/AES)': 571.3 PMKs/s (RTT 3.0)
#25: 'CPU-Core (SSE2/AES)': 518.7 PMKs/s (RTT 2.8)
#26: 'CPU-Core (SSE2/AES)': 544.2 PMKs/s (RTT 2.8)
#27: 'CPU-Core (SSE2/AES)': 560.4 PMKs/s (RTT 2.8)
#28: 'CPU-Core (SSE2/AES)': 569.5 PMKs/s (RTT 2.9)
#29: 'CPU-Core (SSE2/AES)': 566.7 PMKs/s (RTT 3.0)
#30: 'CPU-Core (SSE2/AES)': 552.7 PMKs/s (RTT 2.9)
#31: 'CPU-Core (SSE2/AES)': 518.4 PMKs/s (RTT 3.0)
#32: 'CPU-Core (SSE2/AES)': 613.1 PMKs/s (RTT 3.1)
#33: 'CPU-Core (SSE2/AES)': 597.8 PMKs/s (RTT 2.6)
#34: 'CPU-Core (SSE2/AES)': 572.2 PMKs/s (RTT 3.1)
#35: 'CPU-Core (SSE2/AES)': 577.3 PMKs/s (RTT 2.9)
#36: 'CPU-Core (SSE2/AES)': 597.6 PMKs/s (RTT 3.0)
#37: 'CPU-Core (SSE2/AES)': 558.5 PMKs/s (RTT 3.0)
#38: 'CPU-Core (SSE2/AES)': 560.2 PMKs/s (RTT 3.0)
#39: 'CPU-Core (SSE2/AES)': 523.3 PMKs/s (RTT 2.8)
#40: 'CPU-Core (SSE2/AES)': 566.2 PMKs/s (RTT 3.0)
#41: 'CPU-Core (SSE2/AES)': 594.0 PMKs/s (RTT 2.8)
#42: 'CPU-Core (SSE2/AES)': 544.6 PMKs/s (RTT 3.0)
#43: 'CPU-Core (SSE2/AES)': 512.6 PMKs/s (RTT 3.0)
#44: 'CPU-Core (SSE2/AES)': 558.7 PMKs/s (RTT 3.2)
#45: 'CPU-Core (SSE2/AES)': 596.8 PMKs/s (RTT 2.9)
#46: 'CPU-Core (SSE2/AES)': 579.3 PMKs/s (RTT 2.8)
#47: 'CPU-Core (SSE2/AES)': 562.6 PMKs/s (RTT 2.8)
#48: 'CPU-Core (SSE2/AES)': 605.8 PMKs/s (RTT 3.1)
CUDA:
#1: 'CUDA-Device #1 'Tesla V100-PCIE-16GB'': 125707.5 PMKs/s (RTT 0.4)
#2: 'CUDA-Device #2 'Tesla V100-PCIE-16GB'': 118158.6 PMKs/s (RTT 0.4)
#3: 'CUDA-Device #3 'GeForce GTX 1080 Ti'': 113215.9 PMKs/s (RTT 0.4)
#4: 'CUDA-Device #4 'GeForce GTX 1080 Ti'': 97661.2 PMKs/s (RTT 0.4)

PMKs/sっていうのが、「パスワードを一秒間に何回試したか」です。Pyritは、パスワードを総当たり(ブルートフォース)でクラックします。銀行の暗証番号とかロッカーや自転車の鍵で例えると、0000から9999まで全部ためす感じです。アホな戦略を計算力で全力でごり押していくのが、Pyritスタイル。たぶん。

じっと数字を眺めていると、こんなところに気が付きます:

  • CPU48コアを全部束るとだいたい3万パスワード/secで、12万パスワード/secのTesla1台の1/4くらい

まぁCPUはかなり頑張った方なんじゃないでしょうか?「みんなで同じこと(クラック)をデータ(パスワード)だけ差し替えて計算する」のはGPUのもっとも得意とするところなので、健闘していると思います。

  • GTXとTeslaでほとんど差がない

これ結構ショックなんですが。Teslaはとってもたかいんだぞー。ただ、ひょっとすると、これは「最適化のしがいがある」と捉える事も可能かもしれません。Teslaの方が潜在的な能力があるのは確かなはずなので、ここまで差がないのは、なんか違和感があります。もう少し詳しくプロファイリングする甲斐があるというものでしょう。技術ブログってやつにはこういうのが必要なんですよ。たぶん。

無論、単純にベンチマークが実態と乖離してるだけで、実データを使うとそうでもない可能性もあります。

  • このサーバが全力を出すと、zoiのパスワード([a-z]が11文字)を総当たりクラックするのに200年くらい掛かる

まぁ200年も業務データがクラックされなければ十分なんじゃないでしょうか。誰もクラックされたデータがなんだったのか、思い出せなくなってる頃だと思います。

ただPyritも一番アホな「総当たり攻撃」だけではなく「辞書攻撃」もするので、場合によってはすぐクラックできる可能性はあります。

ああ。あと。最初の方に出てる「Computed 134716.61 PMKs/s total.」ってのがなんか変な気がするので、次までにその辺はデバッグしておきますね。

<Youtuber語法>ご静聴、ありがとう!
気に入ってくれたみんなはRSSリーダーに登録したり↓↓のソーシャルボタンとかで拡散してね</Youtuber語法>