第12回 配信講義　計算科学技術特論A（2025）

並列計算フレームワーク Kokkos入門の入門 2025-07-03 計算科学技術特論A 第12回 理化学研究所 似鳥 啓吾（にたどり けいご）

自己紹介 • 昨年までA64FXでのN体計算カーネルのチューニングの話をしていま した、専門はこちらなので今回扱う格子QCDは専門外となります • （Kokkosについても正式なデベロッパやエヴァンジェリストではく1 ユーザーとしての感想、個人的見解を含みます） • 2009年博士（理学）、2012年理研AICS研究員、2019年理研R-CCS技 師（現職） • 言語歴 • 2000--: C言語 • 2005--: SSE（言語か？） • 2006--: C++ • 2007--: CUDA • 2024--: Kokkos • 講義用ソースコード • https://github.com/nitadori/Tokuron2025/tree/main

• https://github.com/nitadori/Tokuron2025/tree/main

概要 • 講演者は2024年度からC++上の並列計算フレームワーク Kokkosへと向け格子QCDコードBridge++ベンチマークコード の移植を行ってきた。そのときの経験を元に、Kokkosの導入 時に特につまづきやすいCMAKE、必要最小限のAPI群、開発し たベンチマークコードについて簡潔に紹介する。 • オープンソースのKokkosは単一のユーザーコードからCPU向 けOpenMP、CUDA GPU、その他のアクセラレータ向けの コードを生成可能なライブラリ群であり、ソースコードを単一 かつベンダニュートラルに保ちながらディレクティブベースの ものよりも明示的にアクセラレータを活用したいときの候補と なる。

もくじ 1. 2. 3. 4. Kokkosとは CMAKEから 最小限のAPI 格子QCDでの実装例

はじめに：Kokkosとは • C++用の並列計算（アクセラレータ）開発ためのライブラリな いしフレームワーク • APIがあって呼び出して使うという意味ではライブラリといえるが「 Kokkosのやりかたでコード全体を書こう」という意味ではフレームワ ーク • 公式では「Performance Portability Ecosystem」 • 特徴： • Performance Portability（性能移植性）を最大限に強調してい • Kokkosでひとつのコードを書けばCPUでもGPUでも高速に動作すること • ベンダニュートラル • CUDAで書けばNVIDIAのGPUでは速いけど、特定ベンダに固定されたソフトに なってしまう • Kokkosで書いておけばC++標準とOSSのKokkosにのみ依存

https://kokkos.org/about/overview/ より • 米DOEの国立研究所から始まったが最近は仏CEAも参加している • Lunux FoundationプロジェクトでありHigh Performance Software Foundation (HPSF)の一部（理研もHPSFに入りました）

• https://kokkos.org/about/overview/

使ってみた感想など • Host/Deviceのメモリがはっきりと別れていてここは陽に書く必要が ある • cudaMalloc()やcudaMemcpy()がKokkos::Viewによって適度に抽象化され ており少ない記述で書ける • cudaFree()相当の部分は書かずに済む （スマートポインタ風） • CPU用バックエンドでは冗長なメモリ確保やコピーを揉み消してくれる仕 組みがある • どんな人向けか • 既存のコードを少しづつ書き換えるのには向かない、一気にKokkosで書 き換えるか新規プロジェクトにおすすめ • C++やCUDAの経験はあるに越したことはないが必須ではない • Templateやlambdaを多用しているがユーザーコードに複雑な記述は求め られていない • 強いて言うならcopy captureの挙動にだけ注意 • Fortran/FORTRANネイティブの人にも挑戦してみてよいのでは • C++とはいえ変なオブジェクト指向信仰に毒されていない • 世界は配列とループである（doの代わりにparallel_for）

CMAKEの壁（個人の感想です） • KokkosではHello, world!のような最小のコードであっても CMAKEでビルドするのが原則 • gcc hello.cやnvcc hello.cuのようには行かない • どうしてこうなった？ • ユーザーコードとKokkosのランタイムを同じ構成（CUDAやOpenMP ）でコンパイルしておく必要がある • バックエンドのコンパイラも様々：gcc, nvcc, hipcc... • 他のやり方がなかったとまでは言わないが、何らかのビルドシステム に頼るという判断になったよう • 問題点 • CMAKEが初めての人にはどんなCMakeLists.txtを書けばHello, world!が 動かせるのか見当もつかない • （ちょっと前まで）公式のドキュメントを読んでもどのように手を付 ければいいのかわからない状態

新版のQuick Start • https://kokkos.org/kokkos-core-wiki/get-started/quick-start.html • 最小限のCMakeLists.txtはコピペするだけ cmake_minimum_required(VERSION 3.16) project(MyProject) include(FetchContent) FetchContent_Declare( Kokkos URL https://github.com/kokkos/kokkos/archive/refs/tags/4.5.01.zip ) FetchContent_MakeAvailable(Kokkos) add_executable(HelloKokkos HelloKokkos.cpp) target_link_libraries(HelloKokkos Kokkos::kokkos) • Kokkosのランタイムもbuildディレクトリの中に作られる • あとは $ cmake -B build_openmp -DKokkos_ENABLE_OPENMP=ON -DKokkos_ARCH_ARMV9_GRACE=ON $ cmake -B build_cuda -DKokkos_ENABLE_CUDA=ON -DKokkos_ARCH_HOPPER90=ON $make -C build_openmp -j8 $make -C build_cuda -j8

• https://kokkos.org/kokkos-core-wiki/get-started/quick-start.html
• https://github.com/kokkos/kokkos/archive/refs/tags/4.5.01.zip

SpackでKokkosが入っている場合 • SpackはHPC向けのパッケージマネージャ • 見つけてきたKokkosの構成に合わせてbuildを作ってくれる cmake_minimum_required(VERSION 3.16) project(MyProject) find_package(Kokkos REQUIRED) add_executable(HelloKokkos HelloKokkos.cpp) target_link_libraries(HelloKokkos Kokkos::kokkos) • あとはbuild dirを作るときにspack loadされていればいい . ~/spack-alpha4/share/spack/setup-env.sh • ちょっとしたソースファイ ルの編集だけならmakeだ spack load kokkos+openmp cmake -B build_openmp_spack けすればいい spack unload kokkos+openmp • センターがSpackでKokkos spack load kokkos+cuda を提供してくれてるケース cmake -B build_cuda_spack spack unload kokkos+cuda などではこちらもいいかも

実行してみる • --helpや--kokkos-print-configurationが 利用可能 • 特に後者、本当にGPUを使う版になって いるかの確認に有用 [build_cuda]$ ./HelloKokkos Goodbye World [build_cuda]$ ./HelloKokkos --help --------------------------------------------------------------------------------------------Kokkos command line arguments--------------------------------------------------------------------------------------------------------------------This program is using Kokkos. You can use the following command line flags to control its behavior: Kokkos Core Options: --kokkos-help : print this message --kokkos-disable-warnings : disable kokkos warning messages --kokkos-print-configuration : print configuration --kokkos-tune-internals : allow Kokkos to autotune policies and declare tuning features through the tuning system. If left off, Kokkos uses heuristics --kokkos-num-threads=INT : specify total number of threads to use for parallel regions on the host. --kokkos-device-id=INT : specify device id to be used by Kokkos. ... Join us on Slack, visit https://kokkosteam.slack.com Report bugs to https://github.com/kokkos/kokkos/issues -------------------------------------------------------------------------------- Goodbye World [build_cuda]$ ./HelloKokkos --kokkos-print-configuration Kokkos Version: 4.6.1 Compiler: KOKKOS_COMPILER_NVCC: 1260 KOKKOS_COMPILER_NVHPC: 250100 Architecture: CPU architecture: none Default Device: Cuda GPU architecture: HOPPER90 platform: 64bit Atomics: Vectorization: KOKKOS_ENABLE_PRAGMA_IVDEP: yes KOKKOS_ENABLE_PRAGMA_LOOPCOUNT: no KOKKOS_ENABLE_PRAGMA_UNROLL: yes KOKKOS_ENABLE_PRAGMA_VECTOR: yes Memory: Options: KOKKOS_ENABLE_ASM: no KOKKOS_ENABLE_CXX17: yes KOKKOS_ENABLE_CXX20: no KOKKOS_ENABLE_CXX23: no KOKKOS_ENABLE_CXX26: no KOKKOS_ENABLE_DEBUG_BOUNDS_CHECK: no KOKKOS_ENABLE_HWLOC: no KOKKOS_ENABLE_LIBDL: yes Host Serial Execution Space: KOKKOS_ENABLE_SERIAL: yes Serial Runtime Configuration: Device Execution Space: KOKKOS_ENABLE_CUDA: yes Cuda Options: KOKKOS_ENABLE_CUDA_RELOCATABLE_DEVICE_CODE: no KOKKOS_ENABLE_CUDA_UVM: no KOKKOS_ENABLE_IMPL_CUDA_MALLOC_ASYNC: no Cuda Runtime Configuration: macro KOKKOS_ENABLE_CUDA : defined macro CUDA_VERSION = 12060 = version 12.6 Kokkos::Cuda[ 0 ] NVIDIA GH200 480GB : Selected Capability: 9.0 Total Global Memory: 95 GiB Shared Memory per Block: 48 KiB Can access system allocated memory: 1 via Address Translation Service: 1 Goodbye World

Break • ここまでが「入門の入門」ぐらい • CMAKE • 複数configuratioのbuild分け • ここからは実際のC++構文と最小限のAPI • lambdaとキャプチャ • parallel_for/parallel_reduce • マクロいくつか • View

関数オブジェクトとlambda • C++ lambda: 変数や即値に関数を格納する仕組み • C++11から • それ以前は関数オブジェクト作っていた • lambdaというキーワードはないけれど[]で記述できる • [=]でコピーキャプチャ、[&]で参照キャプチャ、外の変数にアクセスでき る • Kokkosで[=]に代えてマクロKOKKOS_LAMBDAを使う // C++11以前 struct GraterThanDouble{ bool operator()(double lhs, double rhs){ return lhs > rhs; } }; std::sort(a, a+10, GraterThanDouble()); // C++11以降 auto cmp = [](auto lhs, auto rhs){ return lhs > rhs; }; std::sort(a, a+10, cmp); // もしくは std::sort(a, a+10, [](auto lhs, auto rhs){ return lhs > rhs; }); (Kokkos)ではなくC++ 標準のソート 比較関数をユーザーが 渡せる

利用例
Kokkos::View<int*> v("v", 5);
// Fill view with sequentially increasing values v=[ 0,1,2,3,4]
Kokkos::parallel_for("fill", 5, KOKKOS_LAMBDA(int i) { v(i) = i; });
// Compute accumulated sum of v's elements r=0+1+2+3+4
int r;
Kokkos::parallel_reduce(
"accumulate", 5,
KOKKOS_LAMBDA(int i, int& partial_r) { partial_r += v(i); }, r);

• parallel_for(名前, 範囲, 関数);

• 構文としては、loop index iを受け取って、その番号に対して何かをする関数を渡し
ている
• やっているのは
for(int i=0; i<5; i++) v(i) = i;
というだけ、逐次コードから書き換えるのも簡単
• ViewへのアクセスはFortran風の(i)や(i,j)、C風の[i]や[i][j]ではない

• prallel_reduce(名前, 範囲, 関数, &積算値）

• Lambdaの方にも引数が増えている
• OpenMP reductio相当
• max/minは少し複雑になるけど、sumならそのままrを渡していい

開始と終了
#include <Kokkos_Core.hpp>
int main(int argc, char** argv) {
Kokkos::initialize(argc, argv);
{
// main関数内でViewを宣言するときは
// 必ずこの{}の内側でやること
}
Kokkos::finalize();
return 0;
}

• ほぼほぼ「おまじない」なのだが、、、
• finalizeに入る前に全てのViewの寿命が尽きている必要がある
• Viewはスマートポインタ風、自動で開放される

大事なマクロいくつか • CUDAではGPUで実行される関数に__device__を、更にホスト でも実行される関数に__host__も付けなければいけないのでこ れらをラップするマクロがある • OpenMP版が動いてCUDA版がコンパイルエラーのときは付け忘れを 疑ってみよう • KOKKOS_LAMBDA • 意味的には[=]、parallel_for等の引数に • KOKKOS_FUNCTION • 例えばparallel_forの内側から呼ばれる関数などの頭には付けておく必 要がある • KOKKOS_INLINE_FUNCTION • 特にinline関数にしたい場合、あるいはクラス内実装するメンバ関数 のように暗黙にinlineになっているもの

View
• 軽く抽象化された（多次元）配列
• Kokkosの理解はViewの理解が一番大事
• 原則データは全部Viewに置いてください

• 宣言の例
• Kokkos::View<double*> v1("v1", 5); // v1(i)=...
• Kokkos::View<double**> v2("v2", 5, 3); // v2(i,j)=...
• *は元々はポインタの意味だが形骸化していて、個数で多次元配列の次数
を示している
• 最初の引数で名前をつける（省略不可）

• デフォルトでデバイス(GPU)側のメモリを指す
• parallel_forの内側でしかアクセスできない
• ホストとのやりとりは後述

• 代入はshallow copyになる
• [=]で勝手にparallel_forにcopyキャプチャされる
• 関数の引数にするときも値渡でいい（むしろ参照渡 の方がトラブルの元）

コード例
#include <cstdio>
#include <Kokkos_Core.hpp>
int main(int argc, char** argv) {
Kokkos::initialize(argc, argv);
{
const int N = 5;
Kokkos::View<double*> v("v", N);
Kokkos::parallel_for(N, KOKKOS_LAMBDA(int i){
v(i) = double(i + 1);
});
auto v_host = Kokkos::create_mirror_view(v);
Kokkos::deep_copy(v_host, v);
for(int i=0; i<N; i++){
printf("%i %f\n", i, v_host(i));
}
}
Kokkos::finalize();
return 0;
}

• Viewを作ってデバイ
ス側でprallel_forで
書き込むのは先述の
とおり
• ホスト側で読み書き
できるように
mirror_viewを作り、
deep_copyをしてい
る
• シリアルやOpenMP
版だと冗長なメモリ
確保とコピーはスキ
ップされる

多次元Viewのレイアウト（Left/Right） • LayoutLeft • (i, j, k)ならiが最内側でリニアなアドレス • Fortran風（ただし0-based） • 行列a(i, j)に対してcolumn-major • GPU側でのデフォルトレイアウト • LayoutRight • (i, j, k)ならkが最内側でリニアなアドレス • C言語風（[i][j][k]と書いたとき） • 行列a(i, j)に対してrow-major • CPU側でのデフォルトレイアウト

コード例
int main(int argc, char** argv) {
Kokkos::initialize(argc, argv);
{
const int NI = 5;
const int NJ = 3;

11
21
31
41
51

Kokkos::View<double**>
v2d("v2d", NI, NJ);
// Kokkos::View<double**, Kokkos::LayoutLeft> v2d("v2d", NI, NJ);
// Kokkos::View<double**, Kokkos::LayoutRight> v2d("v2d", NI, NJ);

12
22
34
42
52

13
23
33
43
53

5×3行列の添字をN字
順(left)かZ字順(right)
でなぞった結果になる

Kokkos::parallel_for(
"fill 2D",
Kokkos::MDRangePolicy< Kokkos::Rank<2> > ({0,0}, {NI, NJ}),
KOKKOS_LAMBDA(int i, int j){
v2d(i, j) = 10.0 * double(i+1) + 1.0 * double(j+1);
}
);
auto v2d_h = Kokkos::create_mirror_view(v2d);
Kokkos::deep_copy(v2d_h, v2d);
// 生のメモリでの順序はどうなっているのか？
const double *h_ptr = v2d_h.data();
for(int ij=0; ij<NI*NJ; ij++){
printf("%2d %2.0f\n", ij, h_ptr[ij]);
}
}
Kokkos::finalize();
return 0;
}

• 単にスカラーのNの代わりにMDRangePolicyというのが
使える
• OpenMPのcollapse(2)相当

• std::vecto同様にView::data()で生ポインタを抜ける

生成AIを使ってみる（GPT 4o） • 一発でデモコードを 出してくれた • というかKokkos知っ ているのね • 目で追う限り正しい コードになってる • 私の作ったサンプル とそっくり

ついでにCMakeも • 特に自分はCMakeやsh、git 、Pythonに関しては素人なの でAIにいつも助けてもらって いる • 正直、公式のチュートリアル よりわかりやすい面も（日本 語で聞けるというのもあるが ） • ただし常に一発で完動品が出 てくるとは限らないので注意

Break • 大体このぐらいで自分の課題で使ったKokkosのAPIは出尽くし ています • 残りは格子QCDコードBridge++から切り出したベンチマーク コードのKokkos移植についての話です • 言語やフレームワークでは直観にかなったインターフェースが 重要 • 個人的には<>や::を沢山書かせることになるC++のインターフ ェースは何処か失敗しているように感じる

格子QCDコードBridge++ • 日本のコミュニティによる格子QCDコード • 2009年開発開始、2012年バージョン1.0、 2024年バージョン2.0 • （似鳥は素粒子物理学も格子QCDも専門外な のでお手柔らかに） • ドメインウォール・フェルミオンという手法 のミニベンチ • 行列ベクトル積の部分を切り出し • 4次元のリンク変数に3x3の特殊ユニタリ行列 • 格子点上にクォークのスピノール（フェルミ オン） • ただしフェルミオンには5次元目の方向に Ns(〜10)個並んでいる • 高コストだがカイラル対称性をよく保つ

苦労した点（主にソフトウェア工学） Kokkos side Bridge++ side View<TwoSpin****> (nvol, 2, 3, Ns) Field (nin, nvol, nex) KField.hpp libbridgekokkos.so API calls (copy and reordering) • プログラム全体をKokkosで書き換えてしまうわけにはいかない • Bridge++には独自のFieldというデータ構造がありそちらは温存したい • 元のコードはMakefileでビルドされている • Kokkosのヘッダ汚染が伝搬すると全体をCMakeでビルドする羽目になる • 仕方がないのでpImplというイディオムを採用することに • 教科書で見たときにはなんでこんな回りくどいことをするのかと思った • Kokkosで書かれた部分は動的ライブラリ形式に • CMakeでKokkosのランタイムとまとめてlibbridgekokkos.soに • 副作用としてGPU版とOpenMP版を実行時に選べるようになった

データ構造 • GPUやWide SIMDでは AoSよりSoAが連続アクセ スになってよいとされる • Fortranでいうと(3,N)より は(N,3) • しかし実数で24ストリー ムになるのはちょっと多 すぎるかと思った • GPUだとfloat4までまとめ てload/storeできる • 複素2成分のspinor単位に するとピッタリ • Kokkos::complex<float> を2つパックして利用 • これでスピンが6ストリー ム、ゲージが最小で3スト リームに スピン： • 3色、4複素数 (Dirac) • 5次元目にNs View<TwoSpin****> spin(“spin”, nvol/2, 2, 3, Ns); ゲージ： • 3×3複素数, 4方向 • 3行目は再構築も可能 View<TwoSpin****> gauge(“gauge”, nvol/2, 4, 5, 2); Nvol = Nx × Ny × Nz × Nt nvol/2はeven-odd法のため

カーネル構造 Lexical Even Odd 3,0 3,1 3,2 3,3 3,4 3,5 3,6 3,7 3,1 3,3 3,5 3,7 3,0 3,2 3,4 3,6 2,0 2,1 2,2 2,3 2,4 2,5 2,6 2,7 2,0 2,2 2,4 2,6 2,1 2,3 2,5 2,7 1,0 1,1 1,2 1,3 1,4 1,5 1,6 1,7 1,1 1,3 1,5 1,7 1,0 1,2 1,4 1,6 0,0 0,1 0,2 0,3 0,4 0,5 0,6 0,7 0,0 0,2 0,4 0,6 0,1 0,3 0,5 0,7 • Even-odd法でD†Dを計算す ると、2種類のカーネルが交 互に出てくる • Stencil Kernel • Even Odd要素のホッピング • キャッシュがちょっと効く • MPIの袖通信が必要 成分を偶奇で分割す ると結局このDを解 けばいいことになる • Stream Kernel • 5次元方向の三重対角行列と その逆行列 • Element-by-element • 再利用性がないので実効バン ド幅のベンチマークになる （私が勝手に命名したものです）

性能（NVIDIA GH200 / AMD MI250) 324×8 GH200 MI250 before after Stencil timing (sec) 0.21 2.01 1.71 Stream timing (sec) 0.23 1.70 0.87 Overall timing (sec) 0.44 3.70 2.56 Overal TFLOPS (peak) 6.7 (67) 0.8 (23) 1.2 Stream TB/s (peak) 2.4 (4.0) 0.3 (1.6) 0.6 (※)他をスキップするダミーカーネルでやっているので総和は完全には一致しません • AMDの方、Tilingを手動で変えることでStreamカーネルは及第点に • Stencilカーネルはレジスタスピルが大量に出ている • ずっとNVIDIAで開発していてAMDでもいきなり動いたのは素晴らし いが、Performance portabilityは達成できている？ • ただしKokkosマターよりはAMDマターかも（ハードウェアとコンパイラ）

AMD向け変更点
#ifdef __HIPCC__
#define TILING ,{256, 1}
#else
#define TILING
#endif
Kokkos::parallel_for(
Kokkos::MDRangePolicy< Kokkos::Rank<2> > ({0,0}, {4*nvol, 3} TILING),
KOKKOS_LAMBDA(int index, int icolor)
{
Kernel kw, kv;
kw.load (w, index, icolor);
kv.mult (kw, params);
kv.store (v, index, icolor);
});

• Streamカーネルでは色とスピンが混ざらないから格子点あた
り12スレッド使っている
• 明示的に色が別のタイルになるようにすることで性能向上した
• AMDはTilingやバンクコンフリクト回避のpaddinに敏感

おまけ（プリセットテンプレート）
• Nsがコンパイル時定数だと
全ての処理がレジスタ内で完
結する
• ただパラメータをコンパイル
時定数にしていいかと聞くと
コミュニティの人はいい顔を
しない
• F77時代とかならシミュレー
ション前に決め打ちしてから
コンパイルとかもあった？

• 仕方ないので<Ns>はテンプレ
ート変数としてpresetsにあ
る数全部インスタンス化して
実行時に選べるようにした
• 基底クラス継承と再帰テンプ
レートの共演
• KokkosというよりCUDAで
もHIPでもOpenACCでも採
用可能な部品

struct Fops5D_Base;

template <int Ns>
struct Fops5D_dirac : public Fops5D_Base {
struct Kernel{
float upper[Ns];
float lower[Ns];
};
};
constexpr int Ns_presets[] = {8, 10, 12};
constexpr int Len_presets = sizeof(Ns_presets)
/ sizeof(Ns_presets[0]);
// Recursive template
template <int I>
Fops5D_Base *new_Fops5D_dirac(int Ns){
if(Ns == Ns_presets[I]){
fprintf(stderr, "Ns=%d found in presets ", Ns);
return new Fops5D_dirac<Ns_presets[I]>;
}else{
return new_Fops5D_dirac<I+1>(Ns);
}
}
// Terminate the recursion
Template <>
Fops5D_Base *new_Fops5D_dirac<Len_presets>(int Ns){
fprintf(stderr, "Ns=%d NOT found in presets ", Ns);
return new Fops5D_Base;
}

補遺：ガンマ行列の表現と形式 Dirac（エネルギー）形式： Spin射影： 実数（float4）で実装するとき： Weyl（カイラル）形式：

感想（苦労した点） • CMAKE • ビルドシステムなんてソフトを配布するときに使うと思ってた • 裏でコソコソやるのがUNIX的でないというか、、、（OSSビルド時のト ラブルでCMAKEに言うことを聞かせることができず関連するテキストフ ァイルを編集して逃げ切るようなことも） • 公式ドキュメントが良くなったので皆さんもHello, Kokkosしましょう • Plug-in的な開発 • 既存のコードを壊さずに追加の部品だけをKokkosで作るという作業 • ソフトウェア工学の問題だけど、数値計算だけやってきても経験がなかっ た • 格子QCD • 専門外のことをやるにはDomainwall FermionかつEven-oddはちょっと重 かった • リファレンス実装と数式を行き来してにらめっこ • ただしここを速くできると次世代の実計算に繋がるというモチベーション にはなる

AIとの付き合いかた（まとめに代えて） • Chat GPTなしで今回の開発はできたかすら怪しい • 実際に書いてもらったものを採用しているわけではないがCMAKEの 部分などは色々と聞いた • そもそも調べごとの初手がGoogleからChat GPTに移動してしまいど こからどこまでがAIの寄与か判別が難しい • Kokkosのようなマイナー（少なくとも日本では）なツールの内容も知 っていたことの驚いた（しかも日本語で聞ける） • Python、C++、MPI、CUDAは比較的よく知ってる • ハルシネーションはまだあるとはいえ • そもそも調べごとには取っ掛かりが必要、公式のマニュアルを読んで いるといのは最終段階にまで進んでいる状態 • Googleだって100%確かな何かを返してくるわけではないし、それを 言い出せば人間だって（以下自粛） • 新しいものが出てきたときに「そんなものから本物の知と教養は手に 入らない」と怒る人は見かけるが、若い人は自然と順応している