scryptがGPUに破られる時

（推定読了時間：27分00秒）

一般的によく知られている SHA-256 や MD5 などのハッシュ関数は非常に単純な設計となっており、非力なパソコンや組み込み機器、スマフォなどでも高速に計算できます。しかしながらその一方で、ハッシュ関数を手当たり次第に計算し、もとの入力値を復元するいわゆる「ブルートフォース攻撃」が容易であるというデメリットがあります。

特にこのような SHA-256 や MD5 といったハッシュ関数は、GPU を用いるか、もしくは専用のハードウェア (FPGA もしくは ASIC) を製作することで非常に高い効率で計算（攻撃）ができてしまうことが知られています。そのため、GPU ないし専用ハードウェアを用いたとしても、攻撃効率の改善が難しくなるような新たなハッシュ関数がいくつか提案されています。その中で比較的古く (2012年ごろ) に開発され、他のハッシュ関数にも影響を与えている「scrypt^[1]」と呼ばれるハッシュ関数^[2]があります。

scryptは、その計算に利用するメモリ領域が非常に大きく、またランダムアクセスを行う必要もあることから GPU や専用ハードウェアで効率化を図ることは難しいとされていました。しかしながら、オープンソースの電子マネープロジェクト「Bitcoin」をフォークして作られた「Litecoin」と呼ばれるオープンソースプロジェクトで scrypt が採用されて以来、様々な人が様々な最適化を行った結果、ハイエンド GPU を用いることでハイエンド CPU と比べて 10 倍程度の速度^[3]で計算が実行できるようになってしまいました。また、現在 (2014年ごろより) では Litecoin の scrypt 計算に特化した専用ハードウェアが製作されており、GPU で計算を行うよりもさらに何十倍もの電力効率で計算が行えてしまいます。

GPU や専用ハードウェアで効率を上げにくいことを謳って開発された scrypt ですが、このように登場より数年足らずで設計者の思惑は外れてしまったことになります。それではなぜ scrypt は「破られて」しまったのでしょうか？その原因を探るために、安直な方法で実装された計算コードに対し少しづつ最適化を行いながら、なぜそのような最適化で効率が上がるのかなどの検証および解析を行いましたので本記事にて報告させていただきます。

なお、本記事を読むにあたって、最適化に関する知識は仮定しませんが、OpenCL (言語仕様は C 言語とほぼ一緒です) およびコンピュータアーキテクチャに関する基礎知識を仮定します。

本記事で解説されている最適化手法やその考え方は、ハッシュ関数の計算に限らず、メモリを多用する一般的なアプリケーションにも当てはまる部分も多々あります。現在 GPU などを用いた高速な計算 (GPGPU 計算) を行っている、もしくは検討中の方には参考になる部分もあることと思いますので、最適化の方法がわからない方や最適化に行き詰まった方は是非ご一読いただければと思います。

scryptとは？

まずは敵を知る、ということで scrypt の設計思想やアルゴリズムについて簡単にふれます。

GPU耐性、専用ハードウェア耐性

重要性

SHA-256 や MD5 のような伝統的なハッシュ関数では GPU や専用ハードウェアを用いることで非常に効率よく計算ができてしまうために、「ブルートフォース攻撃」が容易であるため scrypt が考えだされた、というのは冒頭で述べたとおりですが、もう少し具体的にどのような攻撃が存在しうるのかをみてみましょう。

ハッシュ関数には、効率的に目的のデータを検索することのできるデータ構造である「ハッシュテーブル」、配列の添字に数値以外の値を用いることのできる「連想配列」、既知の集合の中に与えられた要素が含まれるかどうかを非常に高速に判定することのできる「Bloom フィルタ」など、実に多くの利用用途があり、どれもブルートフォース攻撃の標的となりうる^[4]ものではありますが、最も影響の大きいのはパスワードのハッシュ値を計算する場合でしょう。

ユーザをパスワードで認証する方式のうち最も単純なやり方として、サーバ側にユーザのパスワードを保存しておき、ユーザがパスワードを入力した際にあらかじめサーバ側に記録されていたパスワードと一致するかどうか確かめるという方法が考えられます。しかしながらこの方式ですと、万が一サーバに攻撃者が侵入し保存していたパスワードの情報が盗まれてしまいますと、ユーザが全く同じパスワードを他のサービスで使いまわしていた場合には、攻撃者はそれら全てのサービスにログインできてしまうことになります。そのため、現在ではユーザのパスワードをそのまま保存するのではなく、パスワードのハッシュ値を計算し、その値のみをサーバに保存しておく、という方法^[5]をとるのが一般的です。この方法で認証を行う場合には、サーバ側でユーザの入力したパスワードのハッシュ値を計算し、これをあらかじめ保存しておいたハッシュ値と比較をすればOKです。一方、万が一攻撃者にサーバへ侵入され、パスワードのハッシュ値が盗まれてしまったとしても、元のパスワードは容易には復元できませんので、ユーザがパスワードを複数のサービスで使いまわしていた場合でも比較的安全であるといえます。

しかしながら、攻撃者が非常に効率よく大量のハッシュ値の計算を行えるとしたら、ユーザがパスワードとして入力しうる値に対してしらみ潰しにハッシュ値を計算し、サーバから盗んできたハッシュ値と比較することで元のパスワードを復元することができてしまいます。このように総当りで何らかの答えを見つけるといった攻撃を総称して「ブルートフォース攻撃」といいます。

また、似たようなハッシュ関数の使い方として、ユーザの入力したパスワードのハッシュ値を計算し、そのハッシュ値を共通鍵として用いてデータを暗号化する、というものがあります。具体的な方式はいくつかあるのですが、最も有名なのは PBKDF2 (Password-Based Key Derivation Function 2) と呼ばれるもので、これは SHA-256 などを含めた任意のハッシュ関数と組み合わせることでパスワードから共通鍵を得る方式です。こちらの場合も上記のパスワードチェックの場合と同様に、パスワードとして取りうる値をしらみ潰しに計算していくことで共通鍵を復元できてしまいますので、ハッシュ値の計算が効率よくできてしまうと都合が悪いということが解ると思います。

他にも、たとえば電子署名を作成する際には電子署名対象のデータそのものではなく、署名対象のデータのハッシュ値に対して署名処理を行うことが多いです。このような場合でも、GPU や専用ハードウェアを駆使することで効率よくハッシュ値の計算を行い、ブルートフォース攻撃を仕掛ければ、データのハッシュ値を変えることなくデータのみを都合よく書き換える攻撃が容易となってしまうでしょう。^[6]

一般的には、簡単で、速く、効率よく計算できるアルゴリズムのほうが優秀であるといえますが、こと暗号セキュリティに関しては高速に計算ができることが必ずしも優れた方式であるとは言えないことがお分かりいただけたかと思います。このような事情から、あまり計算効率をあげることができない^[7]ハッシュ関数として scrypt が作られた^[8]のです。

scryptのとった戦略

以上で GPU や専用ハードウェアの利用により効率よく計算ができてしまうと困ったことになることはお分かりいただけたかと思いますが、それでは scrypt がどのようにして GPU 耐性や専用ハードウェア耐性を向上させた^[9]のかをみてみましょう。

scrypt の解説スライドによれば、scrypt の基本的な設計思想は、専用ハードウェアの製作コストを上げるためには使用するメモリ容量を増やし、必要なトランジスタ数を増やせばよい、というところにあります。

メモリを増やせばよい、と聞くと「最近は1万円くらい出せば16GBくらいのメモリを買えるくらい低価格化が進んでいるじゃないか！」というツッコミがきそうですが、一般的に「メモリ」といっているものは「DRAM」と呼ばれる種類のもので、確かに価格は安いのですが、CPU の動作速度と比べるとさほど速くはありません。そのため、いくら安価で大量に利用することができたとしても読み書き速度が演算速度に追いつかず、逆に遅くなってしまいます。

そのため CPU では、DRAM メモリの遅さを補うために CPU 内部にいわゆる「キャッシュメモリ」と呼ばれる機構を搭載しています。 CPU のカタログに記載されている L1 (一次) キャッシュとか L2 (二次) キャッシュとかいうものですね。これらのキャッシュメモリはいわゆる「SRAM」と呼ばれるものであり、SRAM は演算回路と同様にトランジスタで構成されていますので、CPU の動作周波数に追従することのできる非常に高速なメモリとなっています。ところがその一方で DRAM と比べると非常に高価な部品（トランジスタ）を使用しているためにあまり容量をあげることができないという欠点があります。昔と比べればかなり容量は上がってきていますが、それでも L3 キャッシュで 8MiB とか 16MiB とかいう容量しかありません。

ですので現在のコンピュータでは、速度は速いが容量を上げられない SRAM メモリと、容量は大きいが速度の遅い DRAM メモリをうまく組み合わせて速度向上とメモリ容量向上を同時に図っているのです。

最新の CPU ですら SRAM の容量はせいぜい 10MiB 前後であることを考えると、演算速度に追従できる速度を持つ SRAM を大量に確保することがいかに難しいかお分かりいただけるでしょう。特に GPU や専用ハードウェアでは多数の計算コアや回路が協調して並列的に計算を行いますので、その演算コアに比例した数だけメモリが必要となってしまい、コスト的にはかなり厳しいものとなります。

アルゴリズム概要

この節では、前節で説明したように「GPU や専用ハードウェアを作りにくいように、たくさんのメモリを利用する」という設計思想を前提としながら scrypt の具体的なアルゴリズムについて簡単に見ていきましょう。

なお、本記事で利用するのは後半の「簡易版」のみですので、前半の「完全版」の方は読み飛ばしてもらっても構いません。

完全版の擬似コード

以下に scrypt の完全な処理の擬似コードを示します。以下の擬似コードはこちらのソースコードを参考にしています。

N: メモリ使用量および計算量を調整するためのパラメータ。
r: メモリ使用量を調整するためのパラメータ。
p: 計算量を調整するためのパラメータ。
k: scryptの処理単位である「ブロック」の大きさ（ビット単位）。scryptの実装では512ビット=64バイト。

HASH(data): kビットの出力を持つ簡易なハッシュ関数。scryptの実装ではSalsa20/8が用いられる。
AS_LE(data): 入力されたバイト列をリトルエンディアンの整数として解釈し、返却する。
PBKDF2(password, salt, iteration, output_length_in_bits): scryptの実装ではPBKDF2+SHA256を用いている。

// 入力されたデータを簡易的なハッシュ関数に通しながら撹拌し、入力と全く同じ長さのデータを出力する。
// Input:
//   Z: kビット✕2r個の配列。
// Output:
//   kビット✕2r個の配列。
function BlockMix(Z[0..2r]):
  // 作業用のkビットのデータを初期化。
  X := Z[2r-1]
  // Xの値を順次アップデートしながら、その値を2r個の出力Yとする。
  for j in (0..2r-1):
    X &lt;= HASH(X xor Z[j])
    Y[j] := X
  // Yのデータを適切に並び替えて最終結果とする。
  return (Y[0] || Y[2] || ... || Y[2r-2]) || (Y[1] || Y[3] || ... || Y[2r-1])

// scryptのコア・ルーチン。大量のメモリを作業領域として用いながら入力データを撹拌し、入力と全く同じ長さのデータを出力する。
// Input:
//   Bi: kビット✕2r個の配列。
// Output:
//   kビット✕2r個の配列。
function ROMix(Bi[0..2r-1]):
  // BiにBlockMixをj回適用したものをV[j]とし、N個の配列Vを初期化する。
  for j in (0..N-1):
    V[j][0..2r-1] := BlockMix^j(Bi[0..2r-1])
  // BiにBlockMixをN回適用したものをXとして初期化。
  X[0..2r-1] := BlockMix^N(Bi[0..2r-1])
  // N回以下の処理を実行
  for N times:
    // X[2r-1]をリトルエンディアン(LE)の整数と解釈し、Nで割った余りをkとする。kは擬似乱数として振る舞う。
    k &lt;= AS_LE(X[2r-1]) mod N
    // XおよびV[k]の排他的論理和をとり、BlockMixする。kのランダム性より、ランダムな場所からVのデータを読みだす必要がある。
    X[0..2r-1] &lt;= BlockMix(X[0..2r-1] xor V[k][0..2r-1])
  return X[0..2r-1]

// 入力された配列Bをシリアライズし、ひとつのバイト列として返却する。
// Input:
//   B: kビット✕p個✕2r個の配列。
// Output:
//   k✕p✕2rビットのバイト列。
function serialize(B[0..p-1][0..2r-1]):
  return (B[0][0]) || (B[0][1]) || ... || (B[0][2r-1]) || (B[1][0]) || (B[1][1]) || ... || (B[p-1][r-1])

// 与えられたパスワードおよびsaltからscryptハッシュ値を計算する。
// Input:
//   password: ユーザの入力したパスワードのバイト列。
//   salt: saltのバイト列。
//   outlen: 出力されるハッシュ値の長さ。
// Output:
//   outlenビットのバイト列（ハッシュ値）。
function scrypt(password, salt, outlen):
  // PBKDF2を使い、入力のパスワードおよびsaltからkビット×p個×2r個の大きさを持つ配列Bを初期化。
  serialize(B[p][2r]) := PBKDF2(password, salt, iteration=1, output_length_in_bits=k*p*2r);
  // scryptのコア・ルーチンであるROMixを各B[i]に適用していく。この処理は自明に並列化可能である。
  for i in (0..r-1):
    ROMix(B[i])
  // B を salt とし、PBKDF2 を用いて最終結果を計算する。
  return PBKDF2(password, salt=serialize(B), iteration=1, output_length_in_bits=outlen);

N: メモリ使用量および計算量を調整するためのパラメータ。

r: メモリ使用量を調整するためのパラメータ。

p: 計算量を調整するためのパラメータ。

k: scryptの処理単位である「ブロック」の大きさ（ビット単位）。scryptの実装では512ビット=64バイト。

HASH(data): kビットの出力を持つ簡易なハッシュ関数。scryptの実装ではSalsa20/8が用いられる。

AS_LE(data): 入力されたバイト列をリトルエンディアンの整数として解釈し、返却する。

PBKDF2(password, salt, iteration, output_length_in_bits): scryptの実装ではPBKDF2+SHA256を用いている。

// 入力されたデータを簡易的なハッシュ関数に通しながら撹拌し、入力と全く同じ長さのデータを出力する。

// Input:

// Z: kビット✕2r個の配列。

// Output:

// kビット✕2r個の配列。

function BlockMix(Z[0..2r]):

// 作業用のkビットのデータを初期化。

X := Z[2r-1]

// Xの値を順次アップデートしながら、その値を2r個の出力Yとする。

for j in (0..2r-1):

X <= HASH(X xor Z[j])

Y[j] := X

// Yのデータを適切に並び替えて最終結果とする。

return (Y[0] || Y[2] || ... || Y[2r-2]) || (Y[1] || Y[3] || ... || Y[2r-1])

// scryptのコア・ルーチン。大量のメモリを作業領域として用いながら入力データを撹拌し、入力と全く同じ長さのデータを出力する。

// Input:

// Bi: kビット✕2r個の配列。

// Output:

// kビット✕2r個の配列。

function ROMix(Bi[0..2r-1]):

// BiにBlockMixをj回適用したものをV[j]とし、N個の配列Vを初期化する。

for j in (0..N-1):

V[j][0..2r-1] := BlockMix^j(Bi[0..2r-1])

// BiにBlockMixをN回適用したものをXとして初期化。

X[0..2r-1] := BlockMix^N(Bi[0..2r-1])

// N回以下の処理を実行

for N times:

// X[2r-1]をリトルエンディアン(LE)の整数と解釈し、Nで割った余りをkとする。kは擬似乱数として振る舞う。

k <= AS_LE(X[2r-1]) mod N

// XおよびV[k]の排他的論理和をとり、BlockMixする。kのランダム性より、ランダムな場所からVのデータを読みだす必要がある。

X[0..2r-1] <= BlockMix(X[0..2r-1] xor V[k][0..2r-1])

return X[0..2r-1]

// 入力された配列Bをシリアライズし、ひとつのバイト列として返却する。

// Input:

// B: kビット✕p個✕2r個の配列。

// Output:

// k✕p✕2rビットのバイト列。

function serialize(B[0..p-1][0..2r-1]):

return (B[0][0]) || (B[0][1]) || ... || (B[0][2r-1]) || (B[1][0]) || (B[1][1]) || ... || (B[p-1][r-1])

// 与えられたパスワードおよびsaltからscryptハッシュ値を計算する。

// Input:

// password: ユーザの入力したパスワードのバイト列。

// salt: saltのバイト列。

// outlen: 出力されるハッシュ値の長さ。

// Output:

// outlenビットのバイト列（ハッシュ値）。

function scrypt(password, salt, outlen):

// PBKDF2を使い、入力のパスワードおよびsaltからkビット×p個×2r個の大きさを持つ配列Bを初期化。

serialize(B[p][2r]) := PBKDF2(password, salt, iteration=1, output_length_in_bits=k*p*2r);

// scryptのコア・ルーチンであるROMixを各B[i]に適用していく。この処理は自明に並列化可能である。

for i in (0..r-1):

ROMix(B[i])

// B を salt とし、PBKDF2 を用いて最終結果を計算する。

return PBKDF2(password, salt=serialize(B), iteration=1, output_length_in_bits=outlen);

一見するとなにやら複雑なことをやっている印象を受けるかも知れませんが、全体的な処理の流れとしては

PBKDF2を用いて、任意長の入力値（password, salt）から固定長のバイト列を得る。
次の処理（ROMix）を r 回行う。
1. N 個の巨大配列 V を作成する。
2. V からランダムに値を読み出しながらデータを撹拌（BlockMix）していく処理を N 回行う。
以上で計算した結果をPBKDF2で処理し、ユーザの望む出力長（outlen）にする。

といった感じです。

簡易版の擬似コード

前節の完全版のコードを対象として最適化の解析を行っても良いのですが、メモリやCPUなどのリソース消費量を調整するパラメータ N, r, p のうち、 r, p は同じ処理を繰り返すことでリソース消費量を調整しているだけで、本質的には今回の目的とは余り関係がないので、簡単のため以降では r=1, p=1 として解析を行うことにします。

r=1, p=1 としてソースコードを簡略化すると以下のような擬似コードになります。

N: メモリ使用量および計算量を調整するためのパラメータ。
k: scryptの処理単位である「ブロック」の大きさ（ビット単位）。scryptの実装では512ビット=64バイト。

HASH(data): kビットの出力を持つ簡易なハッシュ関数。scryptの実装ではSalsa20/8が用いられる。
AS_LE(data): 入力されたバイト列をリトルエンディアンの整数として解釈し、返却する。
PBKDF2(password, salt, iteration, output_length_in_bits): scryptの実装ではPBKDF2+SHA256を用いている。

// 入力されたデータを簡易的なハッシュ関数に通して撹拌し、入力と全く同じ長さのデータを出力する。
// Input:
//   Z: kビット✕2個の配列。
// Output:
//   kビット✕2個の配列。
function BlockMix(Z[0..1]):
  X := HASH(Z[0] xor Z[1])
  Y := HASH(Z[1] xor X)
  return (X || Y)

// scryptのコア・ルーチン。大量のメモリを作業領域として用いながら入力データを撹拌し、入力と全く同じ長さのデータを出力する。
// Input:
//   B: kビット✕2個の配列。
// Output:
//   kビット✕2個の配列。
function ROMix(B[0..1]):
  // BにBlockMixをj回適用したものをV[j]とし、N個の配列Vを初期化する。
  for j in (0..N-1):
    V[j][0..1] := BlockMix^j(B[0..1])
  // BにBlockMixをN回適用したものをXとして初期化。
  X[0..1] := BlockMix^N(B[0..1])
  // N回以下の処理を実行
  for N times:
    // X[1]をリトルエンディアン(LE)の整数と解釈し、Nで割った余りをkとする。kは擬似乱数として振る舞う。
    k &lt;= AS_LE(X[1]) mod N
    // XおよびV[k]の排他的論理和をとり、BlockMixする。kのランダム性より、ランダムな場所からVのデータを読みだす必要がある。
    X[0..1] &lt;= BlockMix(X[0..1] xor V[k][0..1])
  return X[0..1]

// 与えられたパスワードおよびsaltからscryptハッシュ値を計算する。
// Input:
//   password: ユーザの入力したパスワードのバイト列。
//   salt: saltのバイト列。
//   outlen: 出力されるハッシュ値の長さ。
// Output:
//   outlenビットのバイト列（ハッシュ値）。
function scrypt(password, salt, outlen):
  // PBKDF2を使い、入力のパスワードおよびsaltからkビット×p個×2r個の大きさを持つ配列Bを初期化。
  (B[0] || B[1]) := PBKDF2(password, salt, iteration=1, output_length_in_bits=2k);
  // scryptのコア・ルーチンであるROMixをBに適用する。
  ROMix(B)
  // B を salt とし、PBKDF2 を用いて最終結果を計算する。
  return PBKDF2(password, salt=(B[0]||B[1]), iteration=1, output_length_in_bits=outlen);

N: メモリ使用量および計算量を調整するためのパラメータ。

k: scryptの処理単位である「ブロック」の大きさ（ビット単位）。scryptの実装では512ビット=64バイト。

HASH(data): kビットの出力を持つ簡易なハッシュ関数。scryptの実装ではSalsa20/8が用いられる。

AS_LE(data): 入力されたバイト列をリトルエンディアンの整数として解釈し、返却する。

PBKDF2(password, salt, iteration, output_length_in_bits): scryptの実装ではPBKDF2+SHA256を用いている。

// 入力されたデータを簡易的なハッシュ関数に通して撹拌し、入力と全く同じ長さのデータを出力する。

// Input:

// Z: kビット✕2個の配列。

// Output:

// kビット✕2個の配列。

function BlockMix(Z[0..1]):

X := HASH(Z[0] xor Z[1])

Y := HASH(Z[1] xor X)

return (X || Y)

// scryptのコア・ルーチン。大量のメモリを作業領域として用いながら入力データを撹拌し、入力と全く同じ長さのデータを出力する。

// Input:

// B: kビット✕2個の配列。

// Output:

// kビット✕2個の配列。

function ROMix(B[0..1]):

// BにBlockMixをj回適用したものをV[j]とし、N個の配列Vを初期化する。

for j in (0..N-1):

V[j][0..1] := BlockMix^j(B[0..1])

// BにBlockMixをN回適用したものをXとして初期化。

X[0..1] := BlockMix^N(B[0..1])

// N回以下の処理を実行

for N times:

// X[1]をリトルエンディアン(LE)の整数と解釈し、Nで割った余りをkとする。kは擬似乱数として振る舞う。

k <= AS_LE(X[1]) mod N

// XおよびV[k]の排他的論理和をとり、BlockMixする。kのランダム性より、ランダムな場所からVのデータを読みだす必要がある。

X[0..1] <= BlockMix(X[0..1] xor V[k][0..1])

return X[0..1]

// 与えられたパスワードおよびsaltからscryptハッシュ値を計算する。

// Input:

// password: ユーザの入力したパスワードのバイト列。

// salt: saltのバイト列。

// outlen: 出力されるハッシュ値の長さ。

// Output:

// outlenビットのバイト列（ハッシュ値）。

function scrypt(password, salt, outlen):

// PBKDF2を使い、入力のパスワードおよびsaltからkビット×p個×2r個の大きさを持つ配列Bを初期化。

(B[0] || B[1]) := PBKDF2(password, salt, iteration=1, output_length_in_bits=2k);

// scryptのコア・ルーチンであるROMixをBに適用する。

ROMix(B)

// B を salt とし、PBKDF2 を用いて最終結果を計算する。

return PBKDF2(password, salt=(B[0]||B[1]), iteration=1, output_length_in_bits=outlen);

検証環境

今回の記事で検証に用いたマシンの構成は以下の通りです。

GPU: AMD Radeon R9 280X GAMING 3G (MSI)
CPU・M/B: Asus C60M1-I (AMD APU C-60, 2-core, 1GHz)
メモリ: 4GB x2

GPU は非常に高性能ですが、CPU はマザーボード一体型のものであり、ブラウジングですら厳しいくらいの貧弱なものとなっております。計算はほとんど GPU で行いますのでこれでも全く問題ありません。

ちなみに記事冒頭のアイキャッチ画像は、今回検証に用いた上記マシンの写真となっています。

参照実装

コード

最適化を行うためには最適化対象のコードがないと何も始まりませんので、 scrypt のアルゴリズムを「素直に」（上記の擬似コード通りに）実装したコード（参照実装）を作成しました。

OpenCL を起動するための処理なども含まれておりますのでこの記事ではコードすべての紹介はしませんが、 scrypt のメイン処理部分のみを抜き出すと以下のようになります (reference.cl)。

void scrypt_core(uint4 X[8], __global void *scratchpad) {
  const size_t offset = get_global_id(0) - get_global_offset(0);
  __global uint4 *V = &amp;((__global uint4*)scratchpad)[offset * (8 * SCRYPT_N)];
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // V_i &lt;= X
    for(size_t k=0; k&lt;8; k++) {
      V[8*i+k] = X[k];
    }
    // X &lt;= H(X)
    salsa(X);
  }
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // j &lt;= Integerify(X) % N
    size_t j = X[7].x &amp; (SCRYPT_N - 1);
    // X &lt;= H(X xor V_j)
    for(size_t k=0; k&lt;8; k++) {
      X[k] ^= V[8*j+k];
    }
    salsa(X);
  }
}

void scrypt_core(uint4 X[8], __global void *scratchpad) {

const size_t offset = get_global_id(0) - get_global_offset(0);

__global uint4 *V = &((__global uint4*)scratchpad)[offset * (8 * SCRYPT_N)];

for(size_t i=0; i<SCRYPT_N; i++) {

// V_i <= X

for(size_t k=0; k<8; k++) {

V[8*i+k] = X[k];

}

// X <= H(X)

salsa(X);

}

for(size_t i=0; i<SCRYPT_N; i++) {

// j <= Integerify(X) % N

size_t j = X[7].x & (SCRYPT_N - 1);

// X <= H(X xor V_j)

for(size_t k=0; k<8; k++) {

X[k] ^= V[8*j+k];

}

salsa(X);

}

このコードはこちらから確認・ダウンロードができます。

なお、本コードの作成にあたって Kolivas 氏作成の scrypt 用カーネルを一部流用させていただきました。

パフォーマンス

最適化前の上記コードのパフォーマンスは以下のとおりでした。

95 kH/s

この速度は現時点でのハイエンド CPU と同程度の速度であり、GPU で計算した割にはまださほど速くありません。

最適化方針

既に解説したとおり、巨大なデータを高速な SRAM にすべて乗っけることは非常に難しいです。実際、上の scrypt の参照実装のコードでも巨大配列 V は SRAM に収まりきらないため、比較的速度の遅いグローバル (DRAM) メモリ上に配置しているためにさほど速度が出ていないと考えられます。ですので、本記事では速度の遅いメモリアクセス部分の処理をいかにして改善し、メモリアクセスのレイテンシを隠蔽するのかということを主なターゲットとして最適化していくこととします。

最適化を行う

それでは上記のコードに対して様々な最適化技法を適用し、それによりどの程度のパフォーマンス向上が見られるのかをテストしてみましょう。

ループ・アンローリング

for 文や while 文などでループを行う場合には一般的には以下のような処理がされます。

例えば

for(i=0; i&lt;100; i++) {
  func();
}

for(i=0; i<100; i++) {

func();

}

というコードであれば、

ループ変数の初期化 (i=0)
ループ内部の処理を実行するかどうか判定 (i<100)
ループ内部処理を実行 (func())
ループ変数をインクリメント (i++)
2. へ戻る

という処理（機械語）に翻訳されます。

ループ・アンローリング（ループ展開）と呼ばれる最適化技法は、ループ内部の処理をまとめて複数回実行することで、2. の終了判定処理および 4. のループ変数の更新処理を減らすというものです。先ほどのコードの場合には例えば以下のようにします。

for(i=0; i&lt;100; i+=2) {
  func();
  func();
}

for(i=0; i<100; i+=2) {

func();

}

このようにループ内で同じ処理を二回行うと、終了判定処理とループ変数の書き換え処理が半分になりますのでその分の高速化が期待できます。ちなみに「アンロール」という言葉は、ループによって100回分「巻かれて」いる処理を解きほぐし、50回分の「巻き数」に変換する、というところから来ているのだと思います。

なお、当然

for(i=0; i&lt;100; i+=4) {
  func();
  func();
  func();
  func();
}

for(i=0; i<100; i+=4) {

func();

}

のように四回アンロールしたりと、アンロール回数を増やすことでさらなる高速化が期待できます。

ですので極端には100回アンロールしたくなるところですが、さすがに100回も同じコードをコピペすると、コピペする回数を間違えそうですし、可読性も死にますのでちょっとつらいですね。そのため、多くのコンパイラではコンパイラに対してループ・アンローリングを指示することができます。 C 言語および OpenCL では以下のように記述します。

#pragma unroll 100
for(i=0; i&lt;100; i++) {
  func();
}

#pragma unroll 100

for(i=0; i<100; i++) {

func();

}

unroll の次の 100 はアンロール回数を表していて、省略された場合には 4 が仮定されることが多いようです^[10]。

では、さっそくループ・アンローリングを scrypt のコードに適用してみましょう (unroll.cl)。

void scrypt_core(uint4 X[8], __global void *scratchpad) {
  const size_t offset = get_global_id(0) - get_global_offset(0);
  __global uint4 *V = &amp;((__global uint4*)scratchpad)[offset * (8 * SCRYPT_N)];
  #pragma unroll 32
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // V_i &lt;= X
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      V[8*i+k] = X[k];
    }
    // X &lt;= H(X)
    salsa(X);
  }
  #pragma unroll 32
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // j &lt;= Integerify(X) % N
    size_t j = X[7].x &amp; (SCRYPT_N - 1);
    // X &lt;= H(X xor V_j)
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      X[k] ^= V[8*j+k];
    }
    salsa(X);
  }
}

void scrypt_core(uint4 X[8], __global void *scratchpad) {

const size_t offset = get_global_id(0) - get_global_offset(0);

__global uint4 *V = &((__global uint4*)scratchpad)[offset * (8 * SCRYPT_N)];

#pragma unroll 32

for(size_t i=0; i<SCRYPT_N; i++) {

// V_i <= X

#pragma unroll 8

for(size_t k=0; k<8; k++) {

V[8*i+k] = X[k];

}

// X <= H(X)

salsa(X);

}

#pragma unroll 32

for(size_t i=0; i<SCRYPT_N; i++) {

// j <= Integerify(X) % N

size_t j = X[7].x & (SCRYPT_N - 1);

// X <= H(X xor V_j)

#pragma unroll 8

for(size_t k=0; k<8; k++) {

X[k] ^= V[8*j+k];

}

salsa(X);

}

計算速度の測定結果は次のとおりです。

75 kH/s

うーん、変わらないどころか逆に遅くなってますね……。というかコンパイル終わるの遅い。。。

ループ・アンローリングは非常に基本的な最適化技法ですので、コンパイラが勝手に判断をして勝手にループ・アンローリングをしてくれることがあります。ですので、明示的にアンロールをしようとしたとしても既にコンパイラが勝手にやってくれていて、速度はなにも変わらなかった、という可能性があります。

また、ループ・アンローリングをすることでその分コード量が増えてしまいますので、命令キャッシュから溢れてしまったり、終了判定の分岐先がメモリ番地的に遠くなってしまうなどの理由により逆に遅くなってしまうこともあります。

そのため、ループ・アンローリングはループ内処理が小さく、終了判定処理およびループ変数変更処理がループ内の処理に比べて相対的に大きくなってしまうような場合にのみ適用するように心がけた方がよいでしょう。

上記の点を鑑みて、外部ループのアンロールを廃止し、以下のようなコードに変更してみました (unroll2.cl)。

void scrypt_core(uint4 X[8], __global void *scratchpad) {
  const size_t offset = get_global_id(0) - get_global_offset(0);
  __global uint4 *V = &amp;((__global uint4*)scratchpad)[offset * (8 * SCRYPT_N)];
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // V_i &lt;= X
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      V[8*i+k] = X[k];
    }
    // X &lt;= H(X)
    salsa(X);
  }
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // j &lt;= Integerify(X) % N
    size_t j = X[7].x &amp; (SCRYPT_N - 1);
    // X &lt;= H(X xor V_j)
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      X[k] ^= V[8*j+k];
    }
    salsa(X);
  }
}

void scrypt_core(uint4 X[8], __global void *scratchpad) {

const size_t offset = get_global_id(0) - get_global_offset(0);

__global uint4 *V = &((__global uint4*)scratchpad)[offset * (8 * SCRYPT_N)];

for(size_t i=0; i<SCRYPT_N; i++) {

// V_i <= X

#pragma unroll 8

for(size_t k=0; k<8; k++) {

V[8*i+k] = X[k];

}

// X <= H(X)

salsa(X);

}

for(size_t i=0; i<SCRYPT_N; i++) {

// j <= Integerify(X) % N

size_t j = X[7].x & (SCRYPT_N - 1);

// X <= H(X xor V_j)

#pragma unroll 8

for(size_t k=0; k<8; k++) {

X[k] ^= V[8*j+k];

}

salsa(X);

}

97 kH/s

うーん、最適化前とほとんど変わらないですね。。。

メモリの連続アクセスを意識する

利用している GPU のアーキテクチャにも依存しますが、GPU 上のメモリに対する読み書きは連続的に行う方が高速になることが知られています。

これは多くのアプリケーションではメモリを連続的に読み書きすることが多いため、一箇所のメモリを参照した場合にその周囲のデータも一括してキャッシュをするという機構が多くの CPU や GPU に搭載されているためです。

ですのでここではメモリをなるべく連続的に読むべく、グローバルメモリ上に配置されている配列 V のメモリ上の配置を工夫してみます。

まず、参照実装での取り扱いはどうかというと、V[n][i] を「n 番目のワークアイテム（CPU コア）の保有しているデータのうち、i 番目の要素」を表すとすると、メモリ上では以下のように並んでいます。

V[0][0] V[0][1] … V[0][N-1] V[1][0] V[1][1] … V[1][N-1] V[2][0] …

ところが、前半の

for(size_t k=0; k&lt;8; k++) {
  V[8*i+k] = X[k];
}

for(size_t k=0; k<8; k++) {

V[8*i+k] = X[k];

}

という処理では、V[0][i], V[1][i], …, V[MAX_WORK_ITEM-1][i] というデータに対して同時に書き込みを行っていますので、

V[0][0] V[1][0] … V[MAX_WORK_ITEM-1][0] V[0][1] V[1][1] … V[MAX_WORK_ITEM-1][1] V[0][2] …

というふうに配置したほうが良さそうだと分かります。

そこで次のように書き換えてみましょう。

// n番目の要素のアドレス番地を返却する。
#define addr(n, i) (n*(8*GLOBAL_WORK_SIZE) + 8*offset + i)

void scrypt_core(uint4 X[8], __global void *scratchpad) {
  const size_t offset = get_global_id(0) - get_global_offset(0);
  __global uint4 *V = (__global uint4*)scratchpad;
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // V_i = X
    #pragma unroll
    for(size_t k=0; k&lt;8; k++) {
      V[addr(i, k)] = X[k];
    }
    // X = H(X)
    salsa(X);
  }
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // j = Integerify(X) % N
    size_t j = X[7].x &amp; (SCRYPT_N - 1);
    // X = H(X xor V_j)
    #pragma unroll
    for(size_t k=0; k&lt;8; k++) {
      X[k] ^= V[addr(j, k)];
    }
    salsa(X);
  }
}

// n番目の要素のアドレス番地を返却する。

#define addr(n, i) (n*(8*GLOBAL_WORK_SIZE) + 8*offset + i)

void scrypt_core(uint4 X[8], __global void *scratchpad) {

const size_t offset = get_global_id(0) - get_global_offset(0);

__global uint4 *V = (__global uint4*)scratchpad;

for(size_t i=0; i<SCRYPT_N; i++) {

// V_i = X

#pragma unroll

for(size_t k=0; k<8; k++) {

V[addr(i, k)] = X[k];

}

// X = H(X)

salsa(X);

}

for(size_t i=0; i<SCRYPT_N; i++) {

// j = Integerify(X) % N

size_t j = X[7].x & (SCRYPT_N - 1);

// X = H(X xor V_j)

#pragma unroll

for(size_t k=0; k<8; k++) {

X[k] ^= V[addr(j, k)];

}

salsa(X);

}

288 kH/s

おお、これだけで三倍程度改善しました！

CPU の場合にはキャッシュ機構が非常に優秀であり、連続アクセスにそんなに気をつけなくてもそこまで処理落ちはしないことが多い^[11]ですが、 GPU の場合には演算コアが多く一度に大量のデータを必要にすることや、キャッシュ機構が CPU と比べれば貧弱であることなどから、メモリの連続アクセス性を気にしたコーディングを行うことで大幅な性能向上が見られることが多々あります。

GPU で計算を行う際には必ずメモリは連続的にアクセスするようにしましょう。

時間と空間のトレードオフを利用する

scrypt には「時間と空間（メモリ）のトレードオフ (Time Memory Trade Off; TMTO) が容易である」という特徴があり、大きな欠点のひとつであると見られることが多いです。

多くの計算問題では、メモリを通常より多く消費することでより高速に計算する方法や、逆に計算処理が通常より増えてしまうもののメモリをより節約できる方法があります。例えばデータベースを検索する際にインデックスを作ることでメモリの消費量は増えてしまいますが、検索は非常に高速になることが知られています。また、例えばデータ圧縮を用いると圧縮や展開に計算は必要となってしまいますが、ディスク使用量（ある種のメモリ）は削減することができます。

前者の場合には「空間（メモリ）を増やして、時間を減らす」後者の場合には「時間を増やして、空間（メモリ）を減らす」ようなものになっていますので、時間と空間（メモリ）は互いにトレードオフの関係になっています。このような計算問題には「時間と空間（メモリ）のトレードオフ」がある、といいます。

最近の CPU やとりわけ GPU などでは、演算コアの速度に比べてメモリの転送速度（帯域幅）の方が遥かに小さく、実際に計算を行っている時間よりもメモリからのデータの転送待ちをしている時間の方が遥かに長い、という状況が起こることが多々あります。例えば今回の実験に用いた AMD Radeon R9 280X は倍精度浮動小数点演算性能 1TFLOPS^[12] ですから、一秒間に $2 \times \text{sizeof}(\text{double}) \times 10^{12} = 16\text{TB}$ 分の小数の計算ができることになります。一方、メモリの帯域幅は 288GB/s ですから、演算速度に比べてメモリの帯域幅が非常に小さいことが分かります。そのためとりわけ巨大なデータ同士の演算（行列・ベクトル積など）を行う際には、メモリから毎回データを取ってくるのではなく、いかに効率よく同じデータを使いまわしながら計算を行うのかが極めて重要になります。また、メモリ帯域幅の制限とは別に、利用するメモリ量が減ることによるキャッシュヒット率（欲しいデータがキャッシュ上に乗っており、わざわざメインメモリからデータを取ってくる必要がなくなる確率）の向上なども期待できるでしょう。

そのため、この時間と空間（メモリ）のトレードオフを利用することで scrypt の計算に必要なメモリ使用量を抑えることができれば、ある程度の効率化が見込めると期待できます。

scrypt で登場する巨大なデータは配列 V[] ですが、まずはこの V[] の作り方に注目してみましょう。

V[i] = (BlockMix^i)(B);

1 2	V[i] = (BlockMix^i)(B);

でした^[13]ので、極端には B の値のみを保存しておき、V[i] の値が欲しくなったら上記のコードを実行して「その場で」計算してしまえば全くメモリを消費せずに計算ができそうです。

しかしながら、毎回この計算をするのではさすがに演算器の性能のほうが追いつかなくなってしまいますから、中間をとって例えば 4 個おきに「セーブポイント」を設けて次のようにしてみましょう。

// 作業用配列の初期化
Vtmp[0] = (BlockMix^0)(B);
Vtmp[1] = (BlockMix^4)(B);
Vtmp[2] = (BlockMix^8)(B);
...
Vtmp[SCRYPT_N/4-1] = (BlockMix^(SCRYPT_N-4))(B);

// V[i] を返却する関数
function getV(i):
  V = Vtmp[floor(i/4)];
  (BlockMix^(i%4))(V);
  return V;

// 作業用配列の初期化

Vtmp[0] = (BlockMix^0)(B);

Vtmp[1] = (BlockMix^4)(B);

Vtmp[2] = (BlockMix^8)(B);

...

Vtmp[SCRYPT_N/4-1] = (BlockMix^(SCRYPT_N-4))(B);

// V[i] を返却する関数

function getV(i):

V = Vtmp[floor(i/4)];

(BlockMix^(i%4))(V);

return V;

このようにすると、計算量は getV(i) 内の BlockMix の分だけ増えてしまうものの、メモリ使用量はちょうど $1/4$ に抑えられることが分かります。

以上の考え方を用いてメモリ使用量の削減を行う最適化を行ったコードを以下に示します。 TMTO_FACTOR で何個おきにセーブポイントを設けるのかを指定しています。

#define TMTO_FACTOR (2)

// Returns the address of n-th element on scratchpad.
#define addr(n, i) (n*(8*GLOBAL_WORK_SIZE) + 8*offset + i)

void scrypt_core(uint4 X[8], __global void *scratchpad) {
  const size_t offset = get_global_id(0) - get_global_offset(0);
  __global uint4 *Vtmp = (__global uint4*)scratchpad;
  for(size_t i=0; i&lt;SCRYPT_N; i+=TMTO_FACTOR) {
    // Vtmp[i/TMTO_FACTOR] &lt;= X = (BlockMix^i)(B)
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      Vtmp[addr(i/TMTO_FACTOR, k)] = X[k];
    }
    // X &lt;= (BlockMix^(TMTO_FACTOR))(X) = (BlockMix^(i+4))(B)
    #pragma unroll
    for(size_t k=0; k&lt;TMTO_FACTOR; k++) {
      salsa(X);
    }
  }
  for(size_t i=0; i&lt;SCRYPT_N; i++) {
    // j &lt;= Integerify(X) % N
    size_t j = X[7].x &amp; (SCRYPT_N - 1);
    // V &lt;= Vtmp[floor(j/TMTO_FACTOR)]
    uint4 V[8];
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      V[k] = Vtmp[addr(j/TMTO_FACTOR, k)];
    }
    // V &lt;= (BlockMix^(j%TMTO_FACTOR))(V) = (BlockMix^(j))(B);
    #pragma unroll
    for(size_t k=1; k&lt;TMTO_FACTOR; k++) {
      if(j%TMTO_FACTOR &gt;= k) salsa(V);
    }
    // X &lt;= H(X xor V)
    #pragma unroll 8
    for(size_t k=0; k&lt;8; k++) {
      X[k] ^= V[k];
    }
    salsa(X);
  }
}

#define TMTO_FACTOR (2)

// Returns the address of n-th element on scratchpad.

#define addr(n, i) (n*(8*GLOBAL_WORK_SIZE) + 8*offset + i)

void scrypt_core(uint4 X[8], __global void *scratchpad) {

const size_t offset = get_global_id(0) - get_global_offset(0);

__global uint4 *Vtmp = (__global uint4*)scratchpad;

for(size_t i=0; i<SCRYPT_N; i+=TMTO_FACTOR) {

// Vtmp[i/TMTO_FACTOR] <= X = (BlockMix^i)(B)

#pragma unroll 8

for(size_t k=0; k<8; k++) {

Vtmp[addr(i/TMTO_FACTOR, k)] = X[k];

}

// X <= (BlockMix^(TMTO_FACTOR))(X) = (BlockMix^(i+4))(B)

#pragma unroll

for(size_t k=0; k<TMTO_FACTOR; k++) {

salsa(X);

}

for(size_t i=0; i<SCRYPT_N; i++) {

// j <= Integerify(X) % N

size_t j = X[7].x & (SCRYPT_N - 1);

// V <= Vtmp[floor(j/TMTO_FACTOR)]

uint4 V[8];

#pragma unroll 8

for(size_t k=0; k<8; k++) {

V[k] = Vtmp[addr(j/TMTO_FACTOR, k)];

}

// V <= (BlockMix^(j%TMTO_FACTOR))(V) = (BlockMix^(j))(B);

#pragma unroll

for(size_t k=1; k<TMTO_FACTOR; k++) {

if(j%TMTO_FACTOR >= k) salsa(V);

}

// X <= H(X xor V)

#pragma unroll 8

for(size_t k=0; k<8; k++) {

X[k] ^= V[k];

}

salsa(X);

}

696 kH/s

かなり速くなりましたね。計算量は平均して $1.5$ 倍程度に増えますので一見すると非効率ではありますが、メモリへの書き込み回数が半減するため大幅な速度向上が見られていることがわかると思います。

最適化限界

さて、以上で $95\text{kH/s} \rightarrow 696\text{kH/s}$ と約7倍程度高速化ができたことになりますが、これ以上の最適化はもう不可能なのでしょうか？このことを検討するために、配列 V への読み書き部分の処理をコメントアウトし除去した（正しくない）コードを実行してみましょう(no-mem-access.invalid.cl)。

メモリアクセスの遅さを完全に隠蔽できればこのコードの速度に近い値になるはずです。

1,100 kH/s

一見するとまだまだ改善できそうな気がしますが、時間と空間（メモリ）のトレードオフを利用したために計算量が平均して $1.5$ 倍になったことを考えると $1100\times\frac{2}{3} = 733 \text{kH/s}$ 程度が限界であると考えられます。ですので、BlockMix (salsa) 部分の処理の改善などを除いた、メモリアクセス部分の最適化だけではこれ以上の性能改善は非常に難しいということが分かります。逆に言うと、上記の最適化によって GPU では改善が難しいとされていたメモリアクセスの遅さは、ほぼ完璧に隠蔽することができたと言うことができるでしょう。

おわりに

以上でいかにして「GPU や専用ハードウェアでの計算は難しい」とされてきた scrypt が GPU で高速に計算できるようになってしまったのかお分かりいただけたと思います。

GPU に限らず最適化にはコンピュータの動作原理に関する知識、コンパイラや言語の癖、GPU ごとにことなるアーキテクチャに関する知識など、ハードウェア・ソフトウェア両面での深い知識と洞察が必要になります。「高速化」は読んで字の如くただ単に「速く動くようにすること」ではありますが、純粋に目的通りの動作をするプログラムを書くことに比べて最適化作業は遥かに高い難易度を持っています。しかしながら最適化技法を学ぶうちにハードウェアやソフトウェアのアーキテクチャに触れ、理解が進んでいくことではじめから高速かつ効率よく動作するようにコーディングが行えるようになりますので、是非これを機会に学んでみていただければと思います。

参考情報

書籍

これから OpenCL による最適化を学ばれる方におすすめなのがこちらのフィックスターズさんの書籍です。フィックスターズは GPU や Cell B.E. などのマルチコア・メニーコアをターゲットとして主に企業向けに高速化ソリューションを提供している企業です。

GPU プログラミングは「正しく動くこと」よりも「（場合によっては、多少間違っていても）高速に動くこと」を目的とするものですので、通常とは少し違った考え方が必要となります。この書籍はそのような着眼点をもって書かれておりますので、初心者から中・上級者に至るまで非常に有用だと思いますのでぜひご一読ください。

リンク

Tarsnap – The scrypt key derivation function and encryption utility (scrypt 公式(?)ページ)
Colin Percival, Stronger Key Derivation via Sequential Memory-Hard Functions, presented at BSDCan’09, May 2009. (英語論文)
scrypt: A new key derivation function (scrypt 解説スライド)
Mining Hardware Comparison – Litecoin Wiki (CPU や GPU ごとに scrypt の計算スピードが載っている)
GitHub: sgminer-dev/sgminer (scrypt 計算用カーネルを含む)

脚注：

「えす・くりぷと」もしくは「すくりぷと」と読むのが一般的です。ただ、後者だと「scr i pt」と紛らわしいのと、おそらく Secure-CRYPT の略から来ていると思われますので、個人的には前者の「えす・くりぷと」と呼んでいます ↩
scrypt は正確にはユーザの入力したパスワードから共通鍵を得る方式であり、厳密にはハッシュ関数ではないのですが、動作的にはハッシュ関数とほぼ同じなので本記事では特に区別せず「ハッシュ関数」と言ってしまいます ↩
CPU では Intel Core i7 4770 で毎秒約90,000回 (90kH/s) の計算ができるのに対し、ほぼ同価格帯の GPU である AMD Radeon R9 280X ではおおよそ 700kH/s 超の速度で計算ができます ↩
上に挙げた利用例はどれも一見するとブルートフォース攻撃をしたとしてもあまり意味がなさそうなものばかりですが、例えばハッシュ値に意図的に衝突を起こさせることで性能低下を起こさせることができてしまうなど、何らかの攻撃が可能です ↩
ただし、単純にパスワードのハッシュ値をとっただけですと、辞書攻撃やレインボーテーブルなどの手法により元のパスワードを復元されてしまう可能性がありますので、salt と呼ばれるランダムな文字列を加えてからハッシュ値を取ることが推奨されています ↩
ご指摘いただいた @nasakawa 様、ありがとうございました ↩
もちろんたくさんパソコンを用意すればいくらでも計算能力は水増しできますので、根本的な対策にはなっていないのですが、専用ハードウェアによって一般的なコンピュータで行うのと比べて同じマシン代と電気代で 100 倍と言った効率で計算できてしまうと、その分攻撃に必要なコストが下がってしまいますから、計算効率をある水準より抑えることは非常に重要です ↩
なお、scrypt は設計者が開発・運営を行っているオンラインバックアップサービスの Tarsnap でデータを暗号化するために作られたようです ↩
この記事の主旨どおり、GPU や専用ハードウェアを用いることでパソコンより効率よく計算ができてしまうわけですが、それでも SHA-256 などの伝統的なハッシュ関数と比べれば遥かに効率を上げにくいものとなっています ↩
具体的なオプション省略時のデフォルト値については、お使いのコンパイラのドキュメントをご参照ください ↩
特に最近の CPU キャッシュは 16MiB といった容量がありますので、それ以下のメモリ領域しか使わないのであればランダムアクセスをしたとしてもさほど速度低下は起きません ↩
FLOPS = FLoating-point Operations Per Second。一秒間に何回の浮動小数点演算を行えるかを表す指標で、とりわけ数値計算の分野では重要とされています ↩
実際には V[i] や B は配列なので、配列としてコピーする操作が必要ですが、ここでは可読性のため省略します ↩

Master.D
- 2016年12月15日

はじめまして。Master.Dと申します。
最近GPU解析に強いハッシュ関数を勉強しておりまして、こちらにたどり着きました。

こちらの記事を読んだところ
「ハイエンドCPUに比べて7倍程度早くできた」との事ですが、
これはGPUで可能な限り並列計算させて7倍でしょうか、それとも1コア（1スレッド？）だけで7倍でしょうか？

返信
- びりある
  - 2017年2月2日
  
  返信遅くなってしまい申し訳ありません。
  これは当然「CPU一台と、GPU一台を比べて」ということになります。
  （もちろん、世代や価格帯によって処理能力が変わりますのでこの「7倍」というのはあくまで参考値となります）
  
  返信
  - Master.D
    - 2017年4月23日
    
    返信ありがとうございます。
    「GPU1台丸ごと使ってCPUの7倍」となると、確かに速くなったわけですが、
    これぐらいの範囲であれば依然としてScriptは並列解析しやすいPBKDF2よりはるかに解析に強いことになりますよね。
    
    未だに米国政府がPBKDF2を推奨ハッシュアルゴリズムにしてるそうなので、
    新しく出たArgon2も含めて、いったいどれがハッシュアルゴリズムに向いているのか
    専門家のまとまった記事を読んでみたいです。
    お暇がありましたらお願いします。
    
    返信
    - Master.D
      - 2017年4月23日
      
      すみません。
      酔った勢いでいきなり無礼なお願いをしてしまいました。
      酔払いの戯言として受け流して頂ければ幸いです。
      
      返信
匿名
- 2017年6月18日

メモリの連続アクセスってコンパイラがよしなにやってくれると思ってました、違ったんですね

返信
儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】 | 2ちゃんねるまとめ速報
- 2019年8月7日

[…] 説明が上に来るだろ？家ではそうだぞ。 ScryptのコンセプトはASIC耐性だよな。 https://blog.visvirial.com/articles/519 にはSRAMの話で説明してる。よってSHA256の演算の様には爆速にはならんだろ […]

返信
anycoin ~仮想通貨2chまとめ~
- 2019年9月7日

[…] 説明が上に来るだろ？家ではそうだぞ。 ScryptのコンセプトはASIC耐性だよな。 https://blog.visvirial.com/articles/519 にはSRAMの話で説明してる。よってSHA256の演算の様には爆速にはならんだろ […]

返信
儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】｜ビットコイン相場から今後オススメICOとは
- 2020年7月19日

[…] 説明が上に来るだろ？家ではそうだぞ。 ScryptのコンセプトはASIC耐性だよな。 https://blog.visvirial.com/articles/519 にはSRAMの話で説明してる。よってSHA256の演算の様には爆速にはならんだろ […]

返信
儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】 | 仮想通貨まとめ｜ビットコイン（BTC）、イーサリアム（ETH）、ネム（NEX/XEM）などの最新情報をお届け！
- 2020年8月7日

[…] 説明が上に来るだろ？家ではそうだぞ。 ScryptのコンセプトはASIC耐性だよな。 https://blog.visvirial.com/articles/519 にはSRAMの話で説明してる。よってSHA256の演算の様には爆速にはならんだろ […]

返信
儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】
- 2021年2月1日

[…] ScryptのコンセプトはASIC耐性だよな。 https://blog.visvirial.com/articles/519 にはSRAMの話で説明してる。 […]

返信

返信をキャンセル

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください。

scryptとは？

GPU耐性、専用ハードウェア耐性

重要性

scryptのとった戦略

アルゴリズム概要

完全版の擬似コード

簡易版の擬似コード

検証環境

参照実装

コード

パフォーマンス

最適化方針

最適化を行う

ループ・アンローリング

メモリの連続アクセスを意識する

時間と空間のトレードオフを利用する

最適化限界

おわりに

参考情報

書籍

リンク

コメント

Master.D

びりある

Master.D

Master.D

匿名

儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】 | 2ちゃんねるまとめ速報

anycoin ~仮想通貨2chまとめ~

儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】 ｜ ビットコイン相場から今後オススメICOとは

儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】 | 仮想通貨まとめ｜ビットコイン（BTC）、イーサリアム（ETH）、ネム （NEX/XEM）などの最新情報をお届け！

儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】

返信をキャンセル

儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】｜ビットコイン相場から今後オススメICOとは

儲けたお金の税金・確定申告、市況以外の全般26【仮想通貨】 | 仮想通貨まとめ｜ビットコイン（BTC）、イーサリアム（ETH）、ネム（NEX/XEM）などの最新情報をお届け！