HOMA

向井さんのターンにできた HOMA というプロトコルが最近また LWN に登場していた・・・というか、この記事がきっかけで話題になったのかもしれない。

そのへんからリンクされている記事を冷やかすと Google の Snap というユーザスペースの独自プロトコルが紹介されていた。上の podcast episode で「大企業こういうのやってないんですかねー」「やってたら論文くらいありそうなもんですけどねー」とかいう話をしていたわけだが、論文あったわ。Amazon に至ってはハードウェアまで作ってるし

ユーザランドで実装できるならわざわざカーネルでやらなくても良いのでは・・・という意見は論文中で反論されているが、しょうじき説得力はいまいち。とはいえ TCL/TK を発明した大学の先生が自らばりばり Linux カーネルモジュールを書いているという圧倒的つよつよ感の前に、外野は正座するしかない。経歴から年齢を推定してしまう。

Annapurna Labs

エゴサーチをしていたら:

ということで関連記事を探すとこんなのがあった:

EC2 の特定インスタンスでフラグをたてるとハードウェアアクセラレーションの力で TCP と UDP が自動的に速くなるという新機能。この機能は SRD (Scalable Realiable Datagram) というプロプリエタリなプロトコルの上に作られているらしく、2020 年の論文がリンクされている。

著者を見るとみな Annapurna Labs という子会社の人が書いている。2015 年に Amazon が買収したイスラエルの会社らしい。”Nitro” SoC をはじめ AWS のカスタムシリコン事業(?)を牽引しているもよう。

買収されたチップデザイン会社の活躍といわれて思い出すのが Apple が 2008 年に買った PA Semi. Apple A series の SoC デザインを牽引し、スマホ性能で他社をぶっちぎったことがよく知られている。

Annapurna Labs は Amazon にとっての PA-Semi としてサーバ性能で他社をぶっちぎるのに一役買っている、と言える。

去年 Google Cloud が Intel から VP を引き抜いて SoC をがんばると宣言したニュースがあり、なんなのだろうと思っていた。Annapurna Labs を踏まえるとごく自然な動きなのかもしれない。(注記: わたくし Google 勤務ですが特になんのインサイダー知識もございません。)

こういうハードウェアのスタートアップが活躍なんて「シリコン」バレーだなあ・・・と思いかけたが、P.A Semi (“Palo Alto Semiconductor”) はともかく他はイスラエルなのであっちの方がよっぽど Silicon Valley だね. そして Amazon はシアトルなのだった。ただし Annapurna Lab の創業者はベイエリアにいるらしく、これは Lab 126 などの Amazon ハードウェア部門がそのへんにあるからだろうと思われる。

EV

Bill Gates の本を読んだ副作用で EV について考え直す。

自分はアパート住まいなので自宅での充電ができず、それが deal breaker なので次もガス車を買うだろうと思っていた。あと田舎への遠出に伴う充電の難しさも障壁だった。どのくらいがんばればこうした問題を克服できるのか。

田舎への遠出については割と簡単に対処ができる。今使っているガス車を売らずに残しておけば良い(壊れない限りは)。あと、田舎にもそこそこあるね charging station. 計画的に給電すればなんとかなりそうにも見える。

普段の給電。アパート勢の基本戦略は職場での充電らしいが、自分の家は妻による子の送迎や買い物が主な車の用途で職場充電はできない。そして残念ながら買い物・送迎ルートには fast charger が存在しない。確実に fast charge できそうな最寄りの station は片道 5-10 分くらいの場所にある。わざわざそこに出向く必要がある。充電の頻度は、バッテリーの容量にもよるが週に一回くらいだろうか。週末の外出 (2-3h 程度のドライブを伴う)でも乗るなら週に二回。

つまり週に 1-2 回、充電のためだけの外出 (1 時間弱) が必要になる。これが EV を買った場合の時間的コスト。たとえば自分が早朝に外出し充電しながら読書する 、というような形で支払うことになるだろう。かったるいが、可能といえば可能。

車種。週末はガス車を使い近所でのみ EV という選択肢を考えたが、小型の車はバッテリー容量が小さいので不便()。子が乗る後部座席も狭い。それなりに capacity のある SUV の方が現実的に見える。小型車を現実的に使えるのは持ち家勢のみ。車本体の値段は、ガス車の 5 割増くらいだろうか。高いのみならず、まだ選択肢が少ない。在庫もない。なおこれまで乗っていた日本車メーカーは残念ながらだいぶ EV に出遅れており、他を探さないといけない。


というわけで、出費と時間を投じれば EV への乗り換えはいちおう可能な雰囲気。ただレイオフの飛び交う不景気の最中に大きな出費は控えたいので、当面は指をくわえて EV リテラシーを高めておくのがせいぜい。ロシアの戦争が終わり、中国の疫病鎖国が終わり、市場からそれらの傷が癒え、そのとき勤務先がまだ息をしていたら、我が家に EV がやってくることもある、かもしれない。温暖化阻止に協力できる日は遠いが、そういうもんです。それまで今の車が故障しないことを祈るばかり。

仕事について考える時間

このごろ仕事のことを考える時間がないなと思う。もちろん勤務時間中は仕事のことを考えてるんだけど、そういうのじゃなくて、一歩さがってよく考えるような時間。

仕事が順当でなくなることはよくある。はっきりと行き詰まることもあるし、なんとなく調子が下がっていって気がつくと前に進んでいないこともある。それでも目先には表面的な仕事があるので、気をつけないと自分の不調に気づけず、延々と目先のどうでもいい仕事に手を動かして肝心の大仕事が進まない。それどころか状況が悪化していくこともある。たとえば締め切りが近づくにつれ選択肢は狭まっていく。

だから定期的に一歩さがって自分に問う必要がある:いま自分の仕事は順調なのだろうか・順調でないとしたら何が原因で、どう対処すべきなのだろうか。仕事を個人的にレビューすると言っても良い。

独身だった遠い昔、これは当たり前の所作だった。仕事しかしてない身分だったので、たとえば食事をしながら仕事のことをぼんやり考えているのは普通だった。「仕事のことを考える時間がない」という事態を想像するのは難しかったと思う。しかし家事子守という second shift を抱えハードコア九時五時の毎日を生きる今、業務時間外で仕事のことを考えるのは難しい。仕事が終わった瞬間から、たとえば晩飯調理の手順について考え始めている。そのあと食事をしたらもう子を寝かせる時間。自分は一緒に寝てしまうのでそこで一日が終わる。次の日は早くから起きるけれど、そんな朝から仕事のこととか考えたくない。朝くらい課外活動させてくれ。Podcast 準備で論文読むとかさ。あるでしょ。


自分は仕事のレビューを勤務中にやるのが得意でない。埋め込まれた環境から物理的に距離を置かないと、目先の仕事から注意を切り離せない。たとえば自分は毎週金曜日の仕事に weekly review の時間をとっているけれど、普段の仕事と連続していているとなんとなくタスクのリストを shuffle するだけで終わってしまい、批評性が生まれない。

家はどうかというと、家事をしながら仕事のことを考えることはゼロではないが、仕事の空気は押し出されがち。家には家の気圧がある。

なんとか時間を確保しようと、以前は夜や週末に時間を使っていた。たとえば子供が寝たあとに散歩をしながら考え事をしたり、週末の朝コーヒー屋にでかけたりしていた。あとは通勤中に考え事をすることもあった。でもここ数年の大きな生活の変化で、こういう時間は全部どこかに行ってしまった。パンデミックの最中は仕事どころじゃなかったので気がづかなかったけれど、いざ日常が帰ってくると失ったものに気づく・・・ことすらなく、ただうっすらと息苦しい。最近ようやく思い出した。そういえば仕事のこと考えてないじゃん、みたいな。

仕事のことを考えないと、仕事から関心が薄れる。関心が薄れると考えない。そんなループで仕事のどうでも良さが増していく。家族がいる身で仕事のことばかり考えているのが良いとも限らないのである程度の detachment は必要だけれども、あまりに関心が下がると仕事の成果がでないので、それはそれでよくない。バランスが必要。今はあまりに関心が足りていない。関心が低いので深刻さが沸かないところ、長い Thanksgiving 休暇にあわせがんばって気にかけようとしている今。なお、今は朝 5 時です。


仕事について考える時間を仕事の外の日常に取り戻すべきなのだろうか。あまり良いアデアには思えない。家には家の優先度があって、そこに仕事を持ち込むと摩擦が起こる。それよりは、なんとか業務時間中にレビューできるよう工夫する方が良いのではないか。たとえば在宅勤務の金曜日に半日くらいコーヒー屋に行ってみるとか。これに限らず WFH を活かすのは良い気がするな。出勤のある日も、ランチはいつものオフィスから離れた場所へ、ノートなりラップトップなりを片手に行ってみる。あるいはランチはスキップして、そのぶんコーヒーブレイクに散歩する。

目先の仕事に使う時間が減ってしまうが、それは仕方ない。場合によってはなんとかしてもう少し勤務時間を捻出できないか工夫してみても良いかもしれない。朝、子の遅い食事の終わりを待たずにさっさと出勤してしまうとか。仕事と家を混ぜるよりは、仕事の家の境界を交渉する方が健全というか、チート感が少ない気がする。ハードコア九時五時だと現状の給料を正当化できないんですよ、とかいって。

日常でも、たとえばジョギング中に podcast を聞くのは諦めて仕事のことに頭を使う日があっても良い。あと交通事故で自転車が壊れて以来出勤意欲および手段を失いほとんど WFH してるんだけど、いいかげん自転車を買って所定の日数は会社に行ったほうが良さそう。通勤という時間は、無駄といえば無駄なんだけど、考え事バッファとしてあってもいい気がする。


仕事に限らず、自分からは「考える時間」が失われているなと思う。前のブログの目次を見ると、2020 年の 2 月と 3 月の間には明らかな頻度の断然があり、それは回復していない。今となってはどうやってそんなにたくさん書いていたのか想像すらできない。そういえば会社で昼休みに書いていたこともあった気がするな。家だとそういう切り替えができない。自分は会社オフィスのある生活に最適化された旧世代なのだろうね。まあ家は狭いので仕方なし。

Social Media 復帰 (に失敗)

ここ二か月くらい、またソーシャルメディアで管でも巻いてみるかと思い、Twitter に新しくアカウントを作りログインしていた。動機としては、近所・・・つまりベイエリアとか・・・にすんでいる日本人の知り合いが増えたらいいなというもの。

が、この目論見はうまくいかなかった。全然。

まず、近所に住んでいるとおぼしき人々の Twitter 上での発言を自分が楽しめなかった。海外在住日本人のソーシャルメディアはどうしても expat 的な自意識が働きがち。必ずしも出羽の守とまではいかないけれど、日本から見た海外在住者としての自分をコンテンツにするメディアの圧力がある。そういうの、地元民として接したい身からすると鬱陶しい。東京に住んでたらその手のことわざわざ書かないでしょ、みたいな。あとベイエリア固有の現象として、つよつよエンジニアとしての自分をコンテンツにしている人も多い。意識低く生きている自分はそういうのも眩しすぎて苦手。これらのマジョリティ(?)を排除していくとちょっと個性のある主婦っぽい人のアカウントしかフォローに残らず、こうした人々の内容は割と好きなんだけど、主婦だけフォローしてる中年男性の自分どうなの・・・という居心地の悪さが拭えない。

Expat 及びつよつよエンジニアの皆さんは、実際にオフラインで会って話すと必ずしも悪い人々ではないのだが、Twitter というメディアとしての圧力が人々の中にある自分の苦手な側面をあぶりだし、強調してしまう。それが厳しい。

新しい知り合いが増えそうにないなら、せめて旧知の関係を温故しようと地元の知り合いをフォローしてみた。これは悪くなかった。知り合いたちに expatty な面がないとはいわないけれど、相手の素性を知っているので近況報告としての側面がネガティブな気持ちを上書きしてくれる(場合もある)。ただ、我ながら知り合いの少なさを痛感してしまった。知り合いが少ない以上に、Twitter 上で発言している知り合いが少ない。もうちょっと発言してくれればいいのだが。

ただいざ久々に Twitter をしてみると、人々の発言が少なくなるのもわかる。というのは、面白いことを書くのは難しいのだよね。たとえば自分だと会社に行って・・・というか最近は出勤すらせずもっぱら家で働いて、家事子守して寝る、みたいな暮らしをしており、書くことがない。自分は最近は炊事を熱心にしているのでメシ写真とかを投稿してみたものの、別に特段素敵なものを作っているわけでもない。週末のハイキング先の写真とかは面白くないなりに記録としてはアリだけど、だいたい現地にアンテナがなく投稿できない。

結果として自分の発言にコンテンツとしての価値がうまれない。仕事とかテクノロジー話とかを書けば一定程度の面白さはあるんだろうけれど、地元の人々と接する文脈では役に立たないどころか、非地元の日本語圏アカウントを呼び寄せてしまい関係形成の観点からうれしくない。つまらないことでも書けばいいといえばいいけど、それだとだれも自分をフォローしないよね。つまり Twitter で知り合いを増やしたいなら自分に「面白さ」すなわちコンテンツとしての価値が必要で、しかしそんなもの、普通はないのですよ。プログラマという職業的な側面や、インターネット中毒者としての側面をフィルタすると特に。

個人のコンテンツ化という視点でみると、この地の人々が expat やつよつよエンジニアになるのもわかる。それなら比較的簡単(というと怒られそうだけど)にコンテンツとしての価値を提供できるのである。自分がローカルで価値のあるコンテンツになろうと思ったら地元民として熱心に活動する必要があり、つまり買い物・外食・イベント・地元ニュースなどに気を配らなければならず、じっさいその手のお役立ちアカウントをフォローするとありがたいわけだが、自分がそんな活動をする気にはなれない。つまり自分は Twitter 上に存在する価値がない。Twitter は「発信」することがある人たちのプラットホームなのだと思い知らされた。退屈な庶民に出る幕なし。

なお自分の妻氏は地元民としてかなりアクティブに活動しており面白いコンテンツをもっているのだが、それらはもっぱらプライベートなグループチャットに配信されているのだそうな。もったいない(余計なお世話)・・・。

この「コンテンツの価値フィルタ」によってアクティブなユーザを絞り込んでいる以上、Twitter 上に新しい人間関係を期待するのは筋が悪い。潜在的な近隣住人が大量にフィルタアウトされてしまうし、自分自身もフィルタされ他の人の目につかないから。そういうのは Twitter 以外の場所でやれということなのでしょう。


それでも Twitter 上で発言を続けている逞しい知人たちをフォローするためぼちぼち眺めたりつまらないことを書いたりしていたら、先の大富豪買収騒ぎが勃発した。個人的に大富豪の動向はどうでもいいのだが、人々がその話ばかりするせいで場の空気が悪くなってしまった。平和を求めてソーシャルメディアしてるんだから、キナ臭い話はしないでほしい。加えて、何人かの知り合いたちが Twitter から Mastodon に引っ越してしまった。もともとアクティブな知人の少ない timeline なので、数人いなくなるだけでもう 2-3 人しかアクティブな知り合いが残っていない。

といったところ嫌気が差して観念し、サインアウトしアプリは消した。ソーシャルメディア・・・というか Twitter ・・・に復帰しようとする試みは二か月くらいで頓挫したのだった。まあ、間が悪すぎたね。なお Mastodon は標準で RSS をサポートしているので、彼の地に旅立った人々の発言は feed reader で拝見してます。


問いが二つある。1. Twitter は、もし使うのであればどのように使うべきものなのか。2. 知り合いと交流したいならどこでやるのがよいのか。

まず Twitter の使い方について。ここは、自分を含む多くの人にとっては何かを発言する場所ではなくなっていた。自分は従来のアカウントで podcast の宣伝とかをする必要はあるが、なにか個人的なことを書く動機はない。何かを発言する場でもないし、従来の意味でのソーシャル活動、つまり管巻きとか社交、をする場でもないと思う。不可能ではないが、プラットフォームのインセンティブと合ってない。

ではどう使うのが正しいかというと、面白いこと、興味のあることを投稿している人・・・すなわち Twitter 有名人・・・をフォローして眺めるのに使うのだろうね。個人的な人間関係は混ぜず、発言もせずに、情報収集、ニュースメディアとして使う。自分だったら、興味深いプログラミング・ソフトウェア開発に関係するアカウントを沢山フォローして眺めるのがよいのでしょう。タイムラインもアルゴリズムに任せて、流行っているものを織り込んでいい。これはそのうちまた別のアカウントでも作って試してみたい。ただ子ができてから情報中毒者を引退し咀嚼スループットが失われているので、今は無理だな。意識が高まったときにでも挑戦します。そういえば昔 Nuzzel というツールでそんな読み方をしていたのを思い出した。Twitter に買収されてなくなっちゃったのだよな。

2. 知り合いと交流したいならどこでやるのがよいのか。たぶん都合のいいメディアというのは存在しないのでなんらかの努力が望まれる。ただここに書くには長すぎる話題なので割愛いたします。とりあえず Twitter でないことはわかった。そういう時代は 10 年くらい前でおわってました。

Book: How To Avoid A Climate Disaster

How To Avoid A Climate Disaster: Gates, Bill: 9780241448304: Amazon.com: Books

Bill Gates が書いた Zero Carbon Emission 指南書。自分は環境リテラシーが半世紀前で止まっており、なんとかしたいと思っていたので聴いた。その目的は完全に満たせた。Amazon のレビューが一万件くらいついててすごい。電力(how we plug), マテリアル(how we make things), 農業畜産(how we grow things), 交通 (how we get around), 冷暖房(how we keep cool and stay warm) と CO2 発生源を順番に紹介し、その問題と潜在的な解決を議論している。知らない話が多かった。なおこの項目はだいたい CO2 のシェア順になっている。

そのあと、テクノロジーのイノベーションがないとまったく不可能なので起業家とサイエンティストがんばれ、政府はちゃんとインセンティブを出せ、いらん法律は直せ、民間人は購買行動を通じてシグナルを伝えろ、と主張する。

しょうじき全然できる気がせず、人類滅びちゃうかもな・・・という気持ちにならないでもない。Climate Tech, まったく風評を聞かないがちゃんと儲かってるのだろうか。自分の好奇心外だっただけならいいんだけど。Biden も国民にばらまいてインフレを呼び込んだ税金を climate tech に回していたらよかったのに・・・。

子供を持たないアメリカ人のうち数パーセントは気候変動を意思決定の理由に挙げているらしいと以前読んだ。この話を知ったときはずいぶん想像力旺盛だなと思ったけれど、真面目に調べるほどそういう気持ちになりそうな気はする。

このあとも気が向いたらぼちぼち類書をこなしていきたい。

Year Of Desktop Linux

注文していた Desktop PC が届いた。HP Z2 Small Form Factor G9 Workstation をそれなりに盛ったスペックにしたもの。セールでそこそこ割り引かれていた。速い。

自分は以前 Laptop で Linux を使っていたのでかなりバイアスされているが、Year of Desktop Linux はもう来てるなと思う。これは別に Desktop Linux が良くなったからではなく、Desktop Computer というものの地位が低下したから。だいたいみんな PC 買うなら Laptop なわけです。Desktop PC を買うのって、ゲーマーとか「クリエイター」とか、一部の物好きな人たち。そういう人たちは laptop も持っている。desktop は二台目。

Desktop Linux は Desktop PC ユーザーの多くが必要とするものは持ってないのでニッチだが、そのニッチにはよく適合していると思う。つまり、プログラマとか、科学技術計算したいアカデミックの人とか。そういう人々の需要は満たしている。

そしてお店・・・というかオンラインで普通に変える。HP, Dell, Lenovo – 各社普通に Linux モデルを売っている。そういうのを買えばインストールのトラブルは(基本的には)ないのでセットアップはわりかし一瞬である。ライセンスとかドライバとか別途インストールする面倒が少ないので Windows よりよっぽどラク。

そして Linux が苦手なカジュアルかつモダンな計算機の用途は他のデバイスが満たしてくれる。音楽はスマホで聴くし、ビデオ会議は laptop でやるし、など。あと言うまでもないけど色々なものがウェブベースになったのは当然 Linux Desktop の実用性に寄与している。

こうなると別に手元ではなくクラウドとかに置けばいいのではという気にもなり、それは一理ある。クラウド各社は Cloud Workstation 的なものを提供しており、それらのデフォルト OS は Linux である。自分も仕事ではこういうのを付与され、使っている。手元にあったでかいデスクトップは処分した。ただ、個人で持つには高いのだよね Cloud VM.

まあとにかく Desktop PC を買った自慢です。画面がでかい。CPU が速い。ファンの音は許容範囲。満足。

読書活動 PPMP Ch.9, 10

CHAPTER 9 Parallel histogram

  • 面白くなってまいりました。
  • Baseline approach では、atomicAdd() を使います。Kernel から使える AtomicOps あるの?まじで?速いわけがないのはわかるだけに存在する事実にびっくり。しかも modern GPU の atomic ops は last level cache 上で完結できるらしい(hit した場合)。”Because the access time to the last-level cache is in tens of cycles rather than hundreds of cycles, the throughput of atomic operations is improved by at least an order of magnitude compared to early generations of GPU. This is an important reason why most modern GPUs support atomic operations in the last-level cache.” すごいじゃん!
  • Privatization: は thread block 間で独立に histogram を計算しておき、最後に合算する。contention が減るので速くなる。ただし相変わらず atmicAdd は必要。
  • 発展形で、per-thread block の histogram を shared memory に置く。もっと速い。atomicAdd は shared memory に対しても使える。
  • Shared memory って要するにプログラマがコントロールできる L1 cache なのだね。consistency とかないので write back や inter-core の invalidation とかが必要ないぶん便利ではあるが、それをプログラマが exploit しないといけないので大変。CPU は基本的にそういうのナシなので諦めがつきやすい。むかしのゲーム機とかだと consistency level いじったりするので似てたといえば似てたかな。まともな OS を載せるにはそういうの負荷だけれど。

CHAPTER 10 Reduction

  • reduce() ね。
  • 並列でやるには reduction tree をつくる。概念的なツリーであってデータ構造ではない。
  • このあと最適化の話があり、control divergence の最小化、memory coalescing の最大化、そして thread coarsening と全部入りの最適化ステップを踏んでく。面白い。
  • 特に control divergence の最適化で同じコードパスを通るスレッドが同じ warp に収まるよう細工するのはお見事。
  • こういうの、レイテンシ以外にどういう指標をみてチューンできるのかプロファイラの機能を知りたいねえ。

読書活動 PPMP Ch.8

CHAPTER 8 Stencil

  • Stencil は偏微分を計算するのに使います。Convolution に似てますが、問題の性質から最適化の仕方も変わります。
  • 問題柄三次元で計算がちですが、三次元のタイルはあまりでかくできませせん。したがって halo の割合が大きく DRAM アクセス帯域の効率を下げがちです。
  • そこで thread coarsening をします!つまりループを回して cube の一つの次元を iterate していく。Tile をシフトする必要があるが、再利用する cell は shared memory 間でのコピーなのでレイテンシの問題がない。
  • これは面白いところだよな。CPU だと ring buffer 的にインデクス側で rotate しそうなものだけれど(というか GPU でも同じことはできるけど、説明を簡略化するためにそういうロジックは省いているのかもしれない。)
  • アクセスパターンによっては shared memory を使うまでもなく、ループをまたいでレジスタ(ローカル変数)を使いまわせる。これを register coarsening という。

この本は地の文で「X 行目から Y 行目では ZZ をしています」とコードを解説するんだけど、あまりにフォローするのが難しいので文を読むのは諦めてコードをにらんで済ますことが多い。Literate programming スタイルで書いてくれればよいのになあ。

読書活動 PMPP Ch.7

CHAPTER 7 Convolution

さて具体的なテクニックの章に入ってまいりました。Convolution とか簡単でしょと思っていたが、ここまでの章を読むといくらでも複雑な高速化があり得ることがわかり、やや恐怖。

  • Tile 化するが、Convolution という計算の性質上 input のタイルを output のタイルより大きくしないといけない。
  • Input のタイルを output のタイルと同じ大きさに揃え、はみ出す分は L2 が hit することを期待するアプローチもある。って L2 あんのね。
  • どのアプローチが一番速い、という話はしない。ベンチマークしろということか・・・。

読書活動: PMPP Ch. 6

昨日は休カフェイン日だったので良く寝られた。すなわち若干寝坊。

CHAPTER 6 Performance considerations

  • Memory Coalescing: DRAM の burst access を生かせるよう隣接スレッドが隣接アドレスのメモリに同時アクセスする (ex. row major matrix の同一 row を一緒に読む)
  • Interleaved data distribution: DRAM アクセスは channel が全部使われるようにしつつ、各 channel が異なる bank を順番にアクセスできるようにする。
  • DRAM アクセスがいい感じに散らばっていないといけないとか、キャッシュがある CPU のメモリアクセスのルールと全然違うよなあ。GPU のキャッシュアクセスに相当する shared memory はプログラマが明示的にロードするから、その(locality が高く速いメモリへの) ロードの仕方は spread していないといけない、ということなわけだが・・・。
  • 帯域の(並列の)広さぴったりにしつつ、かつ帯域のシーケンシャルなスループットを最大化する、というのが interleaved data distribution なわけだが、そんなの aware なコードを書くとか不可能だっつーの。まあ CPU でも同じ水準の awareness を持ってコード書いてるわけじゃないから似たようなもんか・・・。
  • この問題は CPU にはないのかというと、まあ、ないね。理論上個々のコアからの並列アクセスは interleaved な方がいいわけだけど、そんなに足並み揃えられないからねえ。GPU もこういう芸当ができるのは SM 単位、という理解でいいのだろうか。
  • そして GPU にもキャッシュはあるらしい。が、これは十分にタイミングの近いアクセスが coalesce されるようにする目的という。
  • Thread coarsening: 本来なら並列化する処理を直列化して同じカーネルの実行に押し込めることで、カーネル単位のオーバーヘッドを減らす。たとえば shared memory へのデータのロードとか。Thread に分割しても下のハードウェアの並列度が低くて実際には parallel にならない場合に有効。
  • Shared memory をカーネルの実行を跨いで渡すことができないのでこういうのが必要になる。勝手に色々やってくれる CPU cache のありがたさが身に染みるわ。
  • Check list が載っており、個々の項目を復習しつつ「この章で使うよー」と紹介。今更だが、GPU のプログラミングって何かを速くするのが全てだから、全ての章が高速化の話なのだね。なかなか認知負荷の高い本である。
  • プロファイラの使い方も教えてほしいけど、「リファレンスみてね」でおしまい。

Link: Performing Calculations on a GPU | Apple Developer Documentation

Performing Calculations on a GPU | Apple Developer Documentation

Mac には Metal があっていいなあとひやかす。Metal が CUDA より良いとは思えないが、それでも GPGPU の環境が標準でついてきてラップトップでもさくっと動かせるというのは Linux や Android で暮らしている身からすると羨ましい。Windows プログラマも同じであろう。ラップトップの Intel GPU じゃ GPGPU も盛り上がらないだろうし、そもそもコード書く準備がかったるいからね Windows. Linux はコードを書くのは構わないけど GPGPU の API がないのは同じ。

読書活動 PMPP Ch. 5

CHAPTER 5 Memory architecture and data locality

  • Shared memory のサイズは実行時に計算して kernel call の configuration <<..>> の中に指定できる。
  • A100 の shared memory は 162KB / SM. L1 cache より小さい。自分の買った T1000 という GPU はもっと小さいことでしょう・・・。同一世代の SM は値段によらず同じ構成なのかな。たぶんそうなのでしょう。
  • “Shared Memory” は on-die なのに “Local Memory” が off-die という命名はどうなんだ・・・とおもったが、shared memory のサイズの小ささを見ると local memory どう考えても on-die に置けませんね、とある種の納得がある。
  • レジスタがめちゃくそ沢山ある (数万?) のに対し shared memory が妙に小さいの、CPU から考えると不思議。レジスタが沢山必要なのはわかるんだけど。
  • 行列計算の例は memory locality の説明用には素晴らしい。一方でこういうコード書かないと性能がでないなら GPU でプログラミングとか無理ぽ・・・というきもち。機械学習してる人とか GPU つかってるけどカーネル自分で書いてるのは一握りのエリート/底辺レイヤのひとだけなのだろうねえ。そしてこの本はそういう人向けに書かれているのだった。

GPGPU, HPC の難しさを垣間見る章でありました。

ARM GPU Cores

によると、Mali GPU には “Shader Core” というのが 10 個くらい入っている。Mali-G72 は最大で 32 core までいけるらしいが、そういうのはデスクトップとか用なのでしょう。 この Shader Core が CUDA でいう SM 相当で、世代に名前がついている。

上のリンクによると Valhall という世代は “processing unit” が、2 つ, Bifrost 世代は 3 つ入っている。この “processing unit” が CUDA でいう “CUDA Core” 相当だと思えばよさそう。Utgard という世代は vertex unit と fragment unit がわかれており、unified shader architecture になってない。

Bifrost というやつは 16fp が 8 個 SIMD できると書いてある。この計算と3 つ processing unit があるという事実の関係はよくわからない。

そしてこいつらは SM みたいに thread を oversubscribe して latency tolerance をたかめる、みたいなことはどれくらいしてくれるのだろうね。


そういえば Apple の M1 Ultra は 64 core GPU で、てきとうにぐぐった記事によるとベンチマークのスコアは RTX 3090 に対し 2x+ 遅くらいである。RTX 3090 はこの Wikipedia ページによると SM が 82 個で、各 SM に 128 CUDA Core くらいある。

RTX の SM 82 個と M1 Ultra の 64 core は comparable である。とすると Apple GPU の各コアの processing unit みたいなやつは 128 を半分くらいにして 64 個くらいはいってるのかな?(こんな雑な計算していいのかな?)

ここでは M1 (無印) GPU は “the M1 GPU supports 24576 = 1024 * 24 simultaneous threads” と言っているが、ハードウェアが並列実行できる SIMD の幅はよくわからない。それは開示されていない、という理解でよさそう。

一歩さがって CUDA Core per SM の数をプログラマが意識する必要があるのかというと・・・無いな。重要なのは SM が持てる thread の数, thread block の数, register の数とかで、これは M1/Metal でもクエリできるっぽい。

つまり SM あたりの CUDA Core の数というのは基本的にマーケティング用の数字盛り上げツール、マニアのトリビアであって、特に有用ではないのだろうな。

とはいえ上の 64 という processing unit per GPU Core の概算はさすがに一桁くらい違いそうな気がする。だって ARM の GPU は 3? 8? とかなわけです。そして M1 も Axx もそういうコア単位のデザインは共有してると思われるわけです。8-20 倍ということはないだろうさすがに。


まとめると

  • ARM の GPU は随分遅そうである。そしてメモリアクセスの隠蔽とかがどうなってるのかはさっぱりわからない。レジスタの数もなにもかもわからない。
  • Apple M1 の GPU もよくわからない。ただ Metal のおかげでいくらかは性能をクエリできる。レジスタもいっぱいあるので NVIDIA GPU 的な latency tolerance は期待できる。
  • Apple の M1 Ultra がなぜ RTX 3090 と同じ桁くらい速いのかはよくわからない。

NVIDIA はプロプリエタリなりに情報が色々オープンで偉かった。AMD とか Intel とか調べてないけどどれくらいアーキテクチャを開示してるのだろうね。

読書活動 PMPP Ch. 3, 4

ひきつづき Programming Massively Parallel Processors – 4th Edition. だらだらインターネットしてたら出遅れたが読んでいきます。

CHAPTER 3 Multidimensional grids and data.

  • 三次元まである構造に対しどういうハードウェアの支援があるのかはわからなかった・・・・。ただデータの数 (ex. 画像のピクセル数) ぶんだけスレッドがあるという抽象は面白いというか、スレッドプール脳の自分はメンタルモデルを補正しなければいけないと思いました。

CHAPTER 4 Compute architecture and scheduling

  • Stream Multiprocessor (SM) は複数の CUDA Core を持つ.
  • SM には速いメモリ(キャッシュみたいの?) がついている
  • Thread block = SM! こういうのが知りたかった!
  • SM には数十しかコアがないのに 1024 per block のスレッドを許すのだから、SM は 1024 スレッドぶんのローカルなステート(stack, pc, etc.) はコアの外のどこかに保持するのだね。この shared memory に入っているのだろうか。
  • __syncthreads()!
  • WARP はスケジューリングの単位で、つまりほんとに並列に動く。現状は 32 thread = 1 WARP らしい。つまり最低でも各 SM は 32 cuda core. 64 core ある SM は 2 WARP を並列に動かさせるということか。それは誰得なのだろうな・・・。
  • とおもいきや SM の Core が WARP より小さい (ex. 16 core) “processing block” に分割され、これが SIMD の単位らしい。なんでやねん!WARP = SIMD ならわかるが、これが違うなら WARM の存在意義とは…
  • というとこれは歴史的経緯で、昔は 1 SM = 1 WARP (= 32 Core) みたいな構成だったらしい。メンタルモデル的には WARP = SIMD と思っておけばよいもよう。
  • そして Volta (2017) 以降では WARP 内で control divergence が起きても同一 WARP 内のそれぞれのコードパスが同時に動けるらしい。SIMD とはなんだったのか。わけわかんねーそりゃ消費電力も増えるわけだわ。
  • なお、厳密には SM は 1 つ “以上” つまりハイエンドだと 2 つとかの thread block を持てる模様。これも SM = WARP というマッピングを破壊しているが、アーキテクチャの進化の帰結なのだろうね。
  • Latency Tolerance. これすごいよな。OS がスケジューリングに関与せずソフトウェアからスレッドのスケジューリングが見えないとかプロプリエタリにも程があるしプロファイリングとかも困難でクソだなとおもっていたが、DRAM アクセスの瞬間にスレッドを切り替えるんだからハードウェア側でやらざるを得ない。仕方ないですね・・・と納得した。というか CPU もメモリアクセスでストールするとかはプロファイラでは見えない(なんらかのカウンタはあると思うけど)ので、可視性は大差ないとも言える。
  • そして latency を上手に隠したいなら thread block のサイズはなるべくデカくしておいた方が良いということなんですね・・・とおもいきや SM は複数の thread blocks を同時にホストできるので、かならずしも block size を 1024 にしなくてもよい。とはい同時にホストできる block 数にも上限があるので、多いに越したことはないそうな。更にスレッドが使うレジスタ数でも utilization が制限される!大変だなー・・・。あるカーネルが必要とするレジスタ数とかクエリできるのだろうか。デバイスの性能がわかっても、コードの性質がわからないとな。

以上。Chapter 4 は面白かったね。

疑問があるとすれば、モバイルの GPU の “コア” たとえば M1 が 10 GPU Core というときそのコアは CUDA の語彙で何に相当するのだろうか。SM? CUDA Core? 常識的に考えたら SM だが、ならその “Core / SM” の中にはいくつの CUDA Core 相当があるの?まあそれは CUDA の教科書には出てこないよなあ。

読書活動: Programming Massively Parallel Processors: A Hands-on Approach

半額セールに釣られて新しい PC を注文した。しかも調子にのって使い道もない NVIDIA の GPU もつけてしまった。受注生産らしく届くのは一カ月以上先。せっかくなので CUDA の本でも読んで気分を盛り上げていきたい。

Programming Massively Parallel Processors – 4th Edition

前の版を 5 年くらい前にざっと読んだ気がするんだけど全然覚えてない。奇しくも今年改版されている。そして Elsevier は PDF を売っている。偉い。

今日は二章まで。CUDA の kernel って shader とは随分違うよねえ。引数・・・はあるが、戻り値がない。そして引数もグローバルな何かである。Thread Block とか完全に存在価値が不明だけれど、それはのちの章で明らかになるらしい。

Book: ADHD 2.0

Amazon.com: ADHD 2.0: New Science and Essential Strategies for Thriving with Distraction – from Childhood Through Adulthood (Audible Audio Edition): Edward M. Hallowell, John J. Ratey, Fred Sanders, Random House Audio: Audible Books & Originals

子の落ち着きじゃなく、まさか ADHD じゃないよな・・・という気持ちと、学校とかでは落ち着きがないと ADHD と診断されて薬を処方されてしまうホラーストーリーなどへの心配から理解を深めようと読んだ。

それなりに理解は深まった。あと子供の ADHD の改善のために親ができることと子の健やかさのために親ができることはかなりの割合で一致しており、自分たちはまあまあ実施していることがわかり、なんとなく安心した。

あと素人診断するものではないが、自分の子は単に集中力がないだけだな・・・という感想。さっさとメシくって歯を磨け。ついでに自分もなんとなく集中力ない気がしていたけれど、ADHD 的な特性はカスりもしていなかった。人間、ADHD とは無関係に集中力がないこともある。