この記事は、 NTT docomo Business Advent Calendar 2025 19日目の記事です。 こんにちは、イノベーションセンターの鈴ヶ嶺です。普段はAIアクセラレータの検証に関する業務に従事しています。 本記事では、まずTenstorrentのAIアクセラレータアーキテクチャを紹介し、その特徴について説明します。次に、複数の演算を1つのkernelに統合するfused kernelによる最適化に注目し、標準正規乱数(randn)を例にTenstorrentのアクセラレータにおける具体的な実装方法と性能評価を共有します。その結果、従来の演算の組み合わせの標準正規乱数の実装と比較して、fused kernel実装により約4倍の高速化を確認しました。 Tenstorrentとは オンチップ計算を活かしたFlash Attention fused kernelの実装と評価 実装 性能評価 まとめ Tenstorrentとは Tenstorrent Inc. は次世代AIアクセラレータを製造する半導体メーカーです。 オープン戦略を掲げており、アクセラレータにはRISC-Vを採用し、ソフトウェアに関してはOSS ( https://github.com/tenstorrent ) として積極的に公開されています。 2025年12月現在ではDEC、AMD、Apple、Teslaを歴任した半導体業界の著名なJim Keller氏がCEOを務めています。 TenstorrentのAIアクセラレータのアーキテクチャについて紹介します。 引用: https://speakerdeck.com/tenstorrent_japan/tensix-core-akitekutiyajie-shuo?slide=7 アクセラレータはTensix Coreと呼ばれる5つのBaby RISC-V、2つのNetwork-on-Chip(NoC)、SRAMで構成されるものが複数搭載されています。 一般的なハードウェア管理キャッシュを持たない構成となっており、明示的にコア付近のSRAMを操作する分散メモリ型のNear Memory Computing(NMC)な設計です。 5つのRISC-Vコアは独立な動作が可能なMIMD(Multiple Instruction、 Multiple Data)アーキテクチャです。 多くの処理は典型的にはデータ読み出しを行うReader kernel(RISC-V 1)、 計算をするCompute kernel(RISC-V 2、 3、 4)、 データ書き込みを行うWriter kernel(RISC-V 5)に分けて実行されます。 後述する標準正規乱数のfused kernel実装ではデータ読み込みが不要のためCompute、Writer kernelのみの実装となっており、処理に合わせて自由度を高く調整できます。 16x16を基本としてtileベースの演算エンジンを積んでおり、Compute kernelはこのエンジンを呼び出します。 kernel間のデータはCircular Buffer (CB)と呼ばれるSRAM上のFIFOキューでやり取りをします。 ホストとのデータ交換は外側のDRAM(GDDR)を介して行われます。 その他の技術詳細は日本法人のTenstorrent Japanから以下にさまざまな資料が公開されているためご参照ください。 https://speakerdeck.com/tenstorrent_japan オンチップ計算を活かしたFlash Attention アクセラレータの特徴として、低コスト化のためにHBM(High Bandwidth Memory)などの高コストなメモリを使わない設計となっています。 そのためできるだけDRAM往復によるオーバーヘッドを避けるために、オンチップのSRAM上で計算する工夫がされます。 ここではLLMのAttention計算の事例を取り上げて、どのようにTenstorrentのAIアクセラレータで効率化されるのかを説明します。 https://github.com/tenstorrent/tt-metal/blob/main/tech_reports/FlashAttention/FlashAttention.md LLMのAttentionはそのまま計算すると、巨大な中間行列によりHBM、 DRAMへのデータ移動がオーバーヘッドとなることが知られております。 FlashAttention 1 2 は、その課題に対して行列をチャンクに分割し、より高速なSRAM上で計算しデータ移動のオーバーヘッドを削減し、高速化する手法です。 TenstorrentのAIアクセラレータでも、このFlashAttentionを適用可能です。 大容量のSRAMを利用して実装され中間データがDRAMに書き込まれないため高速化されます。 以下の図のようにベースライン実装と比較して平均して20倍高速に動作します。 引用: https://github.com/tenstorrent/tt-metal/blob/main/tech_reports/FlashAttention/images/image3.png fused kernelの実装と評価 AIアクセラレータの実行は複数のkernelの実行による、中間計算結果のメモリアクセスや起動オーバーヘッドが課題となります。 そこで複数の計算処理を1つのkernelに統合するfused kernelにより性能を向上させる処理がよく用いられます。 例えばLLMのAttentionなどは計算を最適化するために1つのfused kernelとして実装されています。 ttnn.transformer.scaled_dot_product_attention(input_tensor_q: ttnn.Tensor, input_tensor_k: ttnn.Tensor, input_tensor_v: ttnn.Tensor, *, attn_mask: ttnn.Tensor = None, is_causal: bool = true, scale: float = None, sliding_window_size: int = None, memory_config: ttnn.MemoryConfig = None, program_config: SDPAProgramConfig = None, compute_kernel_config: ttnn.DeviceComputeKernelConfig = None, attention_sink: ttnn.Tensor = None) → ttnn.Tensor https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.transformer.scaled_dot_product_attention.html#ttnn.transformer.scaled_dot_product_attention ここではttnnに実装されていない標準正規乱数を生成するrandnを実装します。 randnは一般的な PyTorchの torch.randn や Numpyの np.random.randn などではサポートされています。 標準正規乱数には、Box-Muller法 3 を用います。 実装 新規のOperation追加は、次のように手順で行います。 https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/adding_new_ttnn_operation.html まず、ホスト側での処理を抜粋すると以下のように実装します。 ttnn/cpp/ttnn/operations/randn/device/randn_device_operation.[cpp|hpp] ではOperationの引数やバリデーションを実装します。 struct RandnDeviceOperation { struct operation_attributes_t { const ttnn::Shape shape; // テンソルの形状 DataType dtype; Layout layout; const MemoryConfig memory_config; MeshDevice* device; const DeviceComputeKernelConfig compute_kernel_config; uint32_t seed; // 乱数seed }; // ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ void RandnDeviceOperation:: validate_inputs ( const operation_attributes_t& operation_attributes, const tensor_args_t& tensor_args) { TT_FATAL ( operation_attributes.dtype == DataType::FLOAT32 || operation_attributes.dtype == DataType::BFLOAT16, "Randn: Output tensor must be Float32 or Bfloat16" ); // dtypeによるバリデーション TT_FATAL (operation_attributes.layout == Layout::TILE, "Randn: Not currently supporting row major layout" ); // メモリレイアウトのバリデーション } アクセラレータ上のkernel実行の詳細は ttnn/cpp/ttnn/operations/randn/device/randn_program_factory.cpp に記述します。 ユーティリティ関数 tt::tt_metal::split_work_to_cores 4 によるコアごとの処理を均等に分散 CreateCircularBuffer によるCB(FIFOキュー)の作成 CreateKernel によるCompute、 Writer kernelの作成 SetRuntimeArgs kernel実行の引数の設定 // split_work_to_coresにより、それぞれのコアに処理を割り振る auto [num_cores, all_cores, core_group_1, core_group_2, units_per_core_group_1, units_per_core_group_2] = split_work_to_cores (grid, units_to_divide); // CBの作成(2tile分の出力ができるように確保する) constexpr uint32_t dst_cb_id = CBIndex::c_0; CircularBufferConfig cb_output_config = CircularBufferConfig (in_out_num_tiles * dtype_tile_size, {{dst_cb_id, out_data_format}}) . set_page_size (dst_cb_id, dtype_tile_size); tt_metal:: CreateCircularBuffer (program, all_cores, cb_output_config); // Writer kernelの設定 const std :: string kernels_dir_path = "ttnn/cpp/ttnn/operations/randn/device/kernels/" ; std :: vector < uint32_t > writer_compile_time_args{dst_cb_id}; tt::tt_metal:: TensorAccessorArgs (output. buffer ()). append_to (writer_compile_time_args); const std :: string writer_file_path = kernels_dir_path + "writer_standard_normal.cpp" ; KernelHandle writer_kernel_id = tt_metal:: CreateKernel ( program, writer_file_path, all_cores, WriterDataMovementConfig (writer_compile_time_args)); // Compute kernelの設定 const std :: vector < uint32_t > compute_compile_time_args{dst_cb_id}; const std :: string compute_file_path = kernels_dir_path + "compute_standard_normal.cpp" ; auto [math_fidelity, math_approx_mode, fp32_dest_acc_en, packer_l1_acc, dst_full_sync_en] = get_compute_kernel_config_args (device-> arch (), operation_attributes.compute_kernel_config); KernelHandle compute_kernel_id = CreateKernel ( program, compute_file_path, all_cores, ComputeConfig{ .math_fidelity = math_fidelity, // 計算の精度 ref: https://speakerdeck.com/tenstorrent_japan/tensix-core-akitekutiyajie-shuo?slide=26 .fp32_dest_acc_en = true , .dst_full_sync_en = dst_full_sync_en, .math_approx_mode = math_approx_mode, .compile_args = compute_compile_time_args, .defines = compute_defines, }); // foreach in split_work_to_coresによる割り振り // kernel引数(1コアあたりの乱数生成のtile数、出力のアドレス)の設定 std :: vector < uint32_t > compute_runtime_args = {seed, tile_offset, units_per_core}; SetRuntimeArgs (program, compute_kernel_id, core, compute_runtime_args); std :: vector < uint32_t > writer_runtime_args = {output. buffer ()-> address (), tile_offset, units_per_core}; SetRuntimeArgs (program, writer_kernel_id, core, writer_runtime_args); // end ここからはkernelの実装を説明します。kernel内で利用可能なAPIは以下になります。 https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis.html Compute kernel ttnn/cpp/ttnn/operations/randn/device/kernels/compute_standard_normal.cpp の抜粋を記述します。 tileベースの命令を用いて処理します。 ここで実際にBox-Muller法で標準正規乱数が生成されます。 // Box-Muller法で標準正規乱数 (Z1, Z2) を生成 // Z1 = sqrt(ln(U1) * -2) * cos(U2 * 2pi) // Z2 = sqrt(ln(U1) * -2) * sin(U2 * 2pi) // 出力CBの末尾に2tile確保 cb_reserve_back (dst_cb_id, 2 ); // タイルレジスタを確保 tile_regs_acquire (); // U1、 U2の一様乱数(0, 1)をレジスタ0, 1に生成 rand_tile ( 0 , flt_min, one_minus); rand_tile ( 1 , flt_min, one_minus); // sqrt(ln(U1) * -2)を計算し、レジスタ0に格納 log_tile ( 0 ); mul_unary_tile ( 0 , neg_two); sqrt_tile ( 0 ); // レジスタ2に2piを詰める fill_tile_bitcast ( 2 , two_pi); // U2 * 2piを計算し、レジスタ3, 1に格納 mul_binary_tile ( 1 , 2 , 3 ); mul_binary_tile ( 1 , 2 , 1 ); // cos(U2 * 2pi)を計算し、レジスタ3に格納 cos_tile ( 3 ); // sin(U2 * 2pi)を計算し、レジスタ1に格納 sin_tile ( 1 ); // Z1 = sqrt(ln(U1) * -2) * cos(U2 * 2pi)を計算し、レジスタ3に格納 mul_binary_tile ( 0 , 3 , 3 ); // Z2 = sqrt(ln(U1) * -2) * sin(U2 * 2pi)を計算し、レジスタ1に格納 mul_binary_tile ( 0 , 1 , 1 ); // 出力dtypeが BFLOAT16 の場合は型変換 #ifdef OUTPUT_DTYPE_BFLOAT16 typecast_tile< 0 , 5 >( 3 ); typecast_tile< 0 , 5 >( 1 ); #endif // レジスタ計算の確定、完了待ち tile_regs_commit (); tile_regs_wait (); // レジスタ3, 1のZ1、 Z2をCBへ書き込み pack_tile ( 3 , dst_cb_id); pack_tile ( 1 , dst_cb_id); // レジスタ解放 tile_regs_release (); // CBの末尾に2タイル追加したことを通知 cb_push_back (dst_cb_id, 2 ); 次にWriter kernel ttnn/cpp/ttnn/operations/randn/device/kernels/writer_standard_normal.cpp を抜粋します。 基本的にはCompute kernelからデータを受け取り、そのままNOC経由で書き込みます。 // CBの先頭に2tileがCompute kernelからpushされるまで待つ cb_wait_front (dst_cb_id, 2 ); // CBの読み取りポインタ取得 uint32_t dst_cb_read_base = get_read_ptr (dst_cb_id); uint32_t dst_cb_read0_ptr = dst_cb_read_base; uint32_t dst_cb_read1_ptr = dst_cb_read_base + dst_tile_bytes; // NOCでタイル単位に非同期書き込み noc_async_write_tile (i, output_addrg, dst_cb_read0_ptr); noc_async_write_tile (i + 1 , output_addrg, dst_cb_read1_ptr); // 書き込み完了までバリア noc_async_write_barrier (); // CBから2tile pop cb_pop_front (dst_cb_id, 2 ); 最後にC++やPythonから呼び出すための実装を追加します。 ttnn/cpp/ttnn/operations/randn/device/[randn|randn_pybind].[cpp|hpp] Tensor Randn:: invoke ( const ttnn::Shape& shape, MeshDevice& device, const DataType dtype, const Layout layout, const MemoryConfig& memory_config, const std :: optional <DeviceComputeKernelConfig>& compute_kernel_config, uint32_t seed) { auto tensor = ttnn::prim:: randn (shape, dtype, layout, memory_config, device, compute_kernel_config, seed); if (layout != Layout::TILE) { tensor = ttnn:: to_layout (tensor, layout); } return tensor; } // ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ void bind_randn_operation (py:: module & pymodule) { bind_registered_operation ( pymodule, ttnn::randn, doc, ttnn::pybind_overload_t{ []( const OperationType& self, const ttnn::Shape& shape, MeshDevice& device, const DataType dtype, const Layout layout, const MemoryConfig& memory_config, const std :: optional <DeviceComputeKernelConfig>& compute_kernel_config, uint32_t seed) { return self (shape, device, dtype, layout, memory_config, compute_kernel_config, seed); }, py:: arg ( "shape" ), py:: arg ( "device" ), py:: kw_only (), py:: arg ( "dtype" ) = DataType::BFLOAT16, py:: arg ( "layout" ) = Layout::TILE, py:: arg ( "memory_config" ) = ttnn::DRAM_MEMORY_CONFIG, py:: arg ( "compute_kernel_config" ) = std :: nullopt , py:: arg ( "seed" ) = 0 }); } 今回実装したより詳しい全体コードは、以下のPull Requestを参照ください。 https://github.com/tenstorrent/tt-metal/pull/34508 性能評価 次のスクリプトで実装したttnn.randnと従来のopを組み合わせたttnn.rand + Box-Muller変換の実装と比較します。 補足としてCPUによる実装も計測します。 import math, time, ttnn, torch, numpy as np def rand_box_muller (shape, *, device, dtype, layout, mem, seed): half = (*shape[:- 1 ], shape[- 1 ] // 2 ) u1 = ttnn.rand(half, device=device, dtype=dtype, layout=layout, memory_config=mem, seed=seed + 1234 ) u2 = ttnn.rand(half, device=device, dtype=dtype, layout=layout, memory_config=mem, seed=seed + 4321 ) r = ttnn.sqrt(ttnn.multiply(ttnn.log(u1), - 2.0 )) th = ttnn.multiply(u2, 2.0 * math.pi) z0 = ttnn.multiply(r, ttnn.cos(th)) z1 = ttnn.multiply(r, ttnn.sin(th)) return ttnn.concat([z0, z1], dim=- 1 ) def fused (shape, *, device, dtype, layout, mem, seed): return ttnn.randn(shape, device=device, dtype=dtype, layout=layout, memory_config=mem, seed=seed + 1234 ) def torch_randn (shape, *, dtype, seed): torch.manual_seed(seed+ 1234 ) return torch.randn(shape, dtype=dtype) def bench (name, fn, *, iters, warmup): for i in range (warmup): fn(i) t0 = time.perf_counter_ns() for i in range (iters): fn(i) mean_ms = (time.perf_counter_ns() - t0) / 1e6 / iters print (f "{name}: {mean_ms:.6f} ms/iter" ) return mean_ms DEVICE_ID = 0 SHAPE = ( 1 , 1 , 1024 , 1024 ) ITERS, WARMUP = 10000 , 1000 LAYOUT, MEM, DTYPE = ttnn.TILE_LAYOUT, ttnn.DRAM_MEMORY_CONFIG, ttnn.float32 device = ttnn.open_device(device_id=DEVICE_ID) res_rand_box = bench( "ttnn.rand + Box-Muller" , lambda i: rand_box_muller(SHAPE, device=device, dtype=DTYPE, layout=LAYOUT, mem=MEM, seed=i), iters=ITERS, warmup=WARMUP) res_randn = bench( "ttnn.randn" , lambda i: fused(SHAPE, device=device, dtype=DTYPE, layout=LAYOUT, mem=MEM, seed=i), iters=ITERS, warmup=WARMUP) print (f "Speedup: {res_rand_box / res_randn:.3f}x" ) ttnn.close_device(device) print ( " \n appendix" ) res_torch = bench( "torch.randn" , lambda i: torch_randn(SHAPE, dtype=torch.float32, seed=i), iters=ITERS, warmup=WARMUP) 4つの Tenstorrent Wormhole™ n300s カードを搭載したTT-LoudBoxサーバで実行した結果が次のようになります。 従来のop組み合わせ(rand×2 + log/sqrt/sin/cos/mul + concat)の実装に比べて、今回fused kernelを実装して約4倍の高速化が達成しました。 ちなみに、CPU(Intel® Xeon® Silver 4309Y)の torch.randn で実行したものと比べるとアクセラレータによる並列実行の恩恵を感じることができると思います。 ttnn.rand + Box-Muller: 0.344376 ms/iter ttnn.randn: 0.085173 ms/iter Speedup: 4.043x appendix torch.randn: 4.509201 ms/iter また、出力されたサンプルの分布を可視化しても標準正規分布として問題ないことが次のように確認できました。 import ttnn, matplotlib.pyplot as plt, numpy as np device = ttnn.open_device(device_id= 0 ) x = ttnn.randn( ( 1 , 1 , 1024 , 1024 ), device=device, dtype=ttnn.float32, layout=ttnn.TILE_LAYOUT, memory_config=ttnn.DRAM_MEMORY_CONFIG, seed= 1234 , ) x = ttnn.to_layout(x, ttnn.ROW_MAJOR_LAYOUT) x = ttnn.from_device(x) x = ttnn.to_torch(x).cpu().numpy().ravel() mean = np.mean(x) var = np.var(x) plt.figure(figsize=( 6 , 4 )) plt.hist(x, bins= 100 , density= True , alpha= 0.7 ) plt.axvline(mean, linewidth= 2 , label=f "mean = {mean:.6f}" ) plt.axvspan(mean - np.sqrt(var), mean + np.sqrt(var), alpha= 0.2 , label=f "var = {var:.6f}" ) plt.title( "Histogram of ttnn.randn()" ) plt.xlabel( "Value" ) plt.ylabel( "Probability Density" ) plt.grid( True ) plt.legend() plt.tight_layout() plt.savefig( "fig.png" ) ttnn.close_device(device) まとめ 本記事では、TenstorrentのAIアクセラレータアーキテクチャとその特徴を紹介しました。また、fused kernelによる具体的な最適化の実装方法と従来手法と比較して約4倍の高速化を達成する性能評価結果を共有しました。 明日のアドベントカレンダーもお楽しみに。 Dao, Tri. "Flashattention-2: Faster attention with better parallelism and work partitioning." arXiv preprint arXiv:2307.08691 (2023). ↩ Shah, Jay, et al. "Flashattention-3: Fast and accurate attention with asynchrony and low-precision." Advances in Neural Information Processing Systems 37 (2024): 68658-68685. ↩ Box, George E. P. and Mervin E. Muller. “A Note on the Generation of Random Normal Deviates.” Annals of Mathematical Statistics 29 (1958): 610-611. ↩ https://github.com/tenstorrent/tt-metal/blob/main/METALIUM_GUIDE.md#spmd-in-metalium ↩

この記事は Japan AWS Ambassador Advent Calendar 2025 の12/16付記事となります。 こんにちは、SCSKの木澤です。  今年もAmbassadorアドベントカレンダーへ寄稿させていただきます（公開が1日遅れですみません） ありがたいことに、今年もre:Inventに参加させていただきました。 今回は参加したキーノートやワークショップを通じて感じたAWSの戦略についてお話ししたいと思います。 re:Invent 2025の参加体験 私はAWS Ambassadorに認定いただいていることもあり、今年もありがたく参加させていただきました。 なおre:Inventは規模が桁違いに大きく、セッションやアクティビティも多いので、イベントで満足感を得られるかどうかは各参加者の行動計画に委ねられることと感じています。私の過去の体験としてはこんな感じで 2022 ブレイクアウトセッション中心 適当に予約したらバス移動が多くなってしまい、とにかく大変だった（圧倒された） 2023 色々な種類にバランス良く参加し、比較的満足度は高かった 新サービスのセッション！と思って意気揚々と参加したら全然参加者がいなかったこともあった 2024 ワークショップに多くに登録 手は動かして特定サービスの理解は深まったが、自社フィードバックの観点では微妙 今年は過去の体験を踏まえ、キーノートとワークショップ、EXPO巡回を中心に若干余裕を持ったスケジュールで計画しましたが、過去と比較しても満足感は高かったですね。 本イベントに参加する意義や感じ方は人それぞれかと思いますが、私としてはしっかりキーノートでトレンドを抑える、それとワークショップで新サービス中心に手を動かすことに価値があるのではないかというのが結論です。 なお、今年のre:Inventのレポートは、同行者の蛭田さんから多く発信いただきました。 現地の雰囲気を良く伝えられているかと思いますので、ご紹介いたします。 re:Invent 2025 「re:Invent 2025」の記事一覧です。 blog.usize-tech.com   AI革命進行中 さて、re:Invent 2025の話に入る前に現在のトレンドから触れたいと思います。 現在、AI革命がもの凄い速度で進行中ですね。 私がIT業界に入って四半世紀過ぎましたが、これほどの激動を感じたことはありません。 なお約30年前にはインターネット革命がありました。 現在のIT覇権を握る企業には、その頃に起業された方々も多いです。 Amazon　1994創業 Google　1998創業 （国内）楽天　1997創業 AI革命によって、今後数十年の覇権を握る企業が現われるのでは？ と投資家が思ってマネーが渦巻いている…というのが現在の状況と認識しています。 AIにおける業界構造 現在のAI革命におけるステイクホルダーとしては、このように整理するとわかりやすいかと思います。 半導体チップメーカー（nVIDIA社など） クラウド基盤提供ベンダー（AWSなど） AIモデル提供ベンダー（Anthropicなど） それらを用いて開発する開発者、ビルダー 現在のトレンドとして、あいにく 1.チップメーカーに富（利益）が集中する歪な状況  になっています。 データセンターは電力確保の争奪戦 急速なAI利用拡大に伴い、全世界的にデータセンターが不足する状況になっています。 特にAI利用においては多くの電力を消費するGPU等を大量に利用しますので、電力消費が莫大となります。 以下ガートナー社のレポートのリンクを掲載しますが、2030年までに電力需要は2倍に達する予想となっています。 米国では盛んな電力需要に応えるため、原発の新設などが計画されています。 参考：Gartner、データセンターの電力需要は2025年に16％増加し、2030年までに2倍になるとの予測を発表 https://www.gartner.co.jp/ja/newsroom/press-releases/pr-20251119-dc   私が思う印象的なスライド AI基盤のアップデート Matt Garman CTOのキーノートでは、冒頭多めの時間をとって、新しいAI学習用チップであるTrainium3プロセッサと、それを144基搭載したTrn3 UltraServersの発表がありました。 なぜ、主にインフラ担当の Peter DeSantisのキーノートでなく、Matt CEOのキーノートで採りあげるのか、とお感じの方もいるかもしれませんが、やはりAIモデルの開発ベンダー等、大規模利用顧客へのメッセージがあるのかと思います。 推しは電力効率（コスパ） Trainium3プロセッサと、それを144基搭載したTrn3 UltraServersの特徴を挙げますと Trainium3プロセッサ 性能２倍 電力効率４倍 推論にも利用可能 EC2 Trn3 UltraServers 性能4.4倍 メガワットあたりのトークン数5倍 敢えて電力効率を推しているあたりで、AWSとしての推しは「コスパ」なんだろうと推察します。 今後電力が逼迫することが予想される状況においては、より低消費電力で動作するAIサーバの需要が高まるという予測もあるのかと思います。 余談になりますが、1ラックにこれだけ詰めて凄いですね。従来型のIAサーバーでもここまで詰めたら熱で異常が起こりそうです。ASICベースのTrainiumプロセッサだからこそできる芸当なのか、と思うとワクワクします。 AIエージェント関連機能の拡充 Amazon Bedrock AgentCoreは、現在AWSイチオシのサービスと言って差し支えないかと思います。 今年7月の発表、10月のGA以降、急激に利用拡大が進んでいると聞いています。 今回re:Invent中の各種発表の中でも、AWSはAIエージェントを動かす環境として最適な基盤だというメッセージが多く含まれていたのが印象的でした。 （Swamiキーノートから） 主なAIエージェント関連のアップデートは以下となります。 AIエージェントの統制 Policy in Amazon Bedrock AgentCore　　　　AIエージェントの統制 Amazon Bedrock AgentCore – Evaluations　 AIエージェントのオブザーバービリティ Frontier Agent群 Kiro Autonomous Agent　　自律的に課題解決し開発を推進するエージェント AWS Security Agent　　　　セキュリティ専門知識を持ったエージェント。セキュアな開発を促進 AWS DevOps Agent　　　　問題発生時に根本原因の特定し解決を支援するエージェント 今回のre:InventはAIエージェント一色と言っても過言でないプログラムとなっており、キーノートで発表されたアップデート以外にも多くのワークショップでAIエージェントの活用ユースケースについて触れたものがありました。 私が参加したワークショップの概要だけ書きますと… [COP403-R] AIエージェントによるクラウド運用の自動化 障害時の一時切り分け、AIによる障害原因の分析、復旧までAIエージェント経由でやってもらうワークショップ。 DNSホスト名をIPアドレスまたはインスタンスIDに解決するツール VPC Reachability Analyzer を活用してネットワークの問題を特定し修正するツール Amazon CloudWatch のログとメトリクスを取得するツール これらのなツールをAIエージェントに与えることで、障害原因の分析と復旧まで行うことができた。 [PEX315] AWS AI を活用したインシデント対応によるインテリジェントなセキュリティ運用 SecurityHub⇒S3に保存したセキュリティログをAIエージェントに分析させるワークショップ。 自然言語を使ってAIエージェントに「重要な発見を教えて」「どう解決できるか」を指示することで、SecurityHubで発見されたポイントとそれに対応する解決手段を得ることができる。 また、「解決してください」と指示すると、エージェントが問題を解決してくれる。 このようなワークショップ（ハンズオン）がありました。 Amazon Bedrock AgentCoreの周辺機能拡充で「使い勝手の良い」AIエージェント開発環境を提供することや、具体的なAIエージェントのユースケースを伝えることで、より開発者（ビルダー）に、AWSのAIエージェント環境を使ってもらいたい！という意思を明確に感じることができました。   まとめ ここまで、基盤側のアップデートと、AIエージェント環境拡充と両面の話をさせていただきました。 AI基盤　：　コスパの良いインフラ基盤を提供する AIエージェント　：　使い勝手の良い機能を提供する 総合すると、AWSの戦略が見えてくると思います。 歪な構造になっている業界構造を逆転させて、きちんとビルダーが儲かる世界を作りたい。 そして、AWSはエコシステムを提供することで引き続き、業界リーダーでありたい。 そんなメッセージを強く感じた今回の訪問でした。   おまけ 本内容については、2025/12/16に開催した JAWS-UG Sales #4のLTで発表させていただきました。 ご聴講いただいた皆様、ありがとうございました。 Download Now!

はじめに こんにちは！NTTドコモビジネスの2025年夏の現場受け入れ型インターンシップに参加させていただきました、インターン生の竹田です。私は現在高専の専攻科1年生で、普段は船舶におけるサイバーセキュリティに関する研究活動を行っています。 この記事では、私が今回のインターンシップで取り組んだ業務体験内容について紹介します。 はじめに 参加のきっかけ インターンシップ概要 OTセキュリティとOsecTの概要把握 テーマ選定 検討1: 船舶での使用プロトコル調査 NMEA 0183 IEC61162-450 検討2: 現状の船内ネットワーク調査 検討3: 船舶pcapに対する現状のIDS製品の出力検証 Talker IDから資産を出力するZeek・Spicyパーサー作成・検証 Zeek・Spicyとは パーサー検証 まとめ イベントへの参加 SOC見学 ドコモとのコラボ企画 イノベーションセンター内での業務共有 TechLunch インターンシップを通して学んだこと おわりに 参加のきっかけ 私がインターンシップに参加したのは、研究のテーマとして「船舶におけるIDS（侵入検知システム）」を扱っていることが理由です。制御システムのセキュリティと船舶のセキュリティは「可用性」を最重要視するという点で類似しているため、実際に制御システムで使用されているOT-IDS（制御システム向けIDS）がどのように運用されているのかを知りたいと思いました。また、船にIDSを導入することで船舶会社にどのような利点が見込めるのか、導入する場合の技術的な課題は何かについて理解したいという思いもありました。 インターンシップ概要 今回私は、NTTドコモビジネス イノベーションセンターで、OTシステム向けIDS・セキュリティ可視化サービスである、OsecT（オーセクト）の開発に関する業務を行いました。 インターンシップで行った業務体験内容を以下にまとめます。 OTセキュリティとOsecTの概要把握 まず、OTセキュリティの概要を把握するために、イノベーションセンターが社会人向けに提供している教育プログラムを受講させていただきました。 講義では、Stuxnet、Mirai、WannaCryといったマルウェアが制御システムに影響を及ぼした事例に加え、GPSスプーフィングやWebカメラの不正利用など、制御環境にも波及する脅威が紹介されました。また、ランサムウェア被害を受けた企業の約53％が50万ドル以上の身代金を支払ったという 調査結果 を知ることになりました。私はこの数字を聞いて、サイバーリスク保険の存在によって経済的負担が減っているために、莫大な金額を請求されても企業が支払いに踏み切ってしまう現状があるのではないかと思いました。 さらに、ICSサイバーキルチェーンやMITRE ATT&CK for ICSといった、OTセキュリティ特有の分析フレームワークも紹介されました。これまで一般的なITセキュリティのフレームワークしか知らなかったため、制御システム向けに整理された独自の手法が存在することを初めて知り、新鮮な学びとなりました。 OTセキュリティの概要を学ぶ上で特に面白いなと感じたのは、 業界によってセキュリティ要件の枠組みが大きく異なる 点です。 船舶分野 UR E26/27 という国際的に統一された技術要件が存在し、グローバルに共通の基準に基づいて安全性やサイバーセキュリティ対策を整備する流れがある。これは国際航行を前提とするため、世界的に統一されたルールが不可欠となっている。 製造分野 国や地域、さらには分野ごとに異なる規格が並立しており、北米では CSA規格 、EUでは NIS2指令、 半導体分野では SEMI規格 などが例として挙げられる。市場や技術領域の多様性が大きいため、各地域、各分野で独自の要件が策定されている。 この違いは、業界ごとの国際性や供給網の特性を反映していると考えられます。船舶は国際航行を前提とするためグローバルに統一されたルールが不可欠である一方、製造業は市場や技術領域の多様性が大きいため、各地域で独自の要件が策定されやすいのだと理解しました。 単に「OTセキュリティ」と括るのではなく、こうした制度設計や標準化の背景を踏まえて学ぶことの重要性を改めて実感しました。 次に、チームが開発しているサービスであるOsecTについての説明を受けました。 (引用元: https://www.ntt.com/business/services/security/security-management/wideangle/osect.html ) OsecTは2つの基本機能を備えています。 ネットワークの可視化 ネットワークトラフィックを分析し、資産・通信・セキュリティリスク等をwebポータル画面で可視化できます。これにより、ネットワークマップやトラフィック量から重要な端末を視覚的に把握でき、対応の優先順位がつけやすくなります。 脅威・脆弱性の検知 通信状況を学習し、脅威や脆弱性を検知した際にアラートを通知します。検知した情報はwebポータルでも確認可能です。 また、 資産台帳と連携する機能 も持っているため、不審な状況を見つけたときに設置場所や担当者などの情報をもとに、素早く対応することが可能となっています。さらに、資産台帳に登録されていない、いわゆる未把握の機器も表示されるため、台帳の漏れを補完したり、不正に接続された端末の存在を調査したりすることも可能です。これにより、既存の資産管理の精度を高めると同時に、セキュリティリスクを早期に発見する仕組みとしても活用できます。 一方で、現行の仕組みには課題があることも示されました。例えば、IPアドレスを持たない機器については、対応しているOTプロトコル以外は可視化の対象外となるため、 すべての資産を完全に把握できるわけではありません 。 説明を聞いて、これまでIDSというと「脅威検知」をするものだという印象がありましたが、OT-IDSには「ネットワークの可視化」という重要な役割も担っているんだということを知り、とても勉強になりました。 テーマ選定 ここまでの講義・説明を受けて、私はある1つの仮説を立てました。 船舶のプロトコルを分析 すれば船内の資産を把握でき、 資産インベントリの作成・更新補助ができるのではないか。 ※資産インベントリとは、 IACS UR E26/27 で提出が義務付けられている、船舶機器の詳細をまとめた台帳です。 船舶のライフサイクルを通じて船舶資産インベントリは更新・維持されなければならないことを考慮して作成する必要があります。なお、 資産インベントリの管理を手動で行う場合、更新作業が煩雑になる可能性があります。 そのため、資産インベントリの更新を効率化し、精度を高めるために、 システム化が推奨されます。 （Class NK | IACS UR E26 p.31 4-1.2より） この点を踏まえ、各船舶機器が発する固有のシグナルを活用できれば、船舶資産の自動識別が可能となり、資産インベントリの作成・更新作業を大幅に補助できるのではないかと考えました。 しかし、OsecTは船舶特有のプロトコルに対して十分な解析機能を持っていません。解析可能とするためには新たに専用のパーサーを開発・実装する必要があります。 そこで、以下のプロセスで調査・検証を進めることにしました。 検討1: 船舶での使用プロトコル調査 調査の結果、船舶では主に2つのプロトコルが使われていることが分かりました。調査内容を以下にまとめます。 NMEA 0183 NMEAは「National Marine Electronics Association（米国海洋電子機器協会）」の略。 船舶機器間の通信のための仕様で、潮流計、音響測深機、ジャイロコンパス、GPSなどのデータを伝送するために使用される。シリアルポートを使用した片方向通信でASCII形式。 例（GPS）： $GPGGA,052400.00,3539.3146239,N,13945.6411751,E,4,07,0.59,4.987,M,34.035,M,1.0,3403*76 （データ引用元: https://ales-corp.co.jp/technical-information-nmea/ ） Talker ID とよばれる、どの機器がそのデータを送信しているかという識別子がある。上記の例では、 $GPGGA の GP の部分がTalker IDとなる。 Talker ID対応表（一部） 引用元： pynmea2/NMEA0183.pdf at master · Knio/pynmea2 · GitHub | **ID** | **解説** | | --- | --- | | AG | 一般的なオートパイロット | | AP | 磁気コンパス連動のオートパイロット | | CD | DSC（デジタル選択呼出）通信装置 | | CR | 無線ビーコン受信機 | | CS | 衛星通信機 | | CT | MF/HF帯ラジオ電話通信機 | | CV | VHF帯ラジオ電話通信機 | | CX | スキャン機能付き無線受信機 | | DF | 方位探知機 | | EC | ECDIS | | EP | EPIRB | | ER | 機関室モニタリングシステム | | GP | GPS受信機 | | HC | 磁気コンパス | | HE | ジャイロコンパス（真北追従） | | HN | ジャイロコンパス（非真北） | | II | 統合計測機器 | | IN | 統合ナビゲーションシステム | | LC | ロランC | | P | メーカー独自（非標準） | | RA | レーダー/ARPA（自動衝突予防装置） | | SD | 音響測深機 | | SN | その他の電子測位システム | | SS | スキャニング音響測深機 | | TI | ターンレートインジケータ | | VD | ドップラー式速度センサ | | DM | 水中・磁気式速度ログ（船速計） | | VW | 水中・機械式ログ（パドルホイール等） | | WI | 気象センサ | | YX | トランスデューザー | | ZA | 原子時計 | | ZC | クロノメーター | | ZQ | 水晶時計 | | ZV | 電波時計 | IEC61162-450 IECは「International Electrotechnical Commission（国際電気標準会議）」の略で、後ろに標準化された規格の番号を付与して文書や通信規格そのものの名として利用される。 UDPマルチキャストを使って、NMEAセンテンスをEthernetパケットで送信するための規格。 例： UdPbC\0\s:GP0001\$GPGGA,052400.00,3539.3146239,N,13945.6411751,E,4,07,0.59,4.987,M,34.035,M,1.0,3403*76\r\n タグブロック部と呼ばれる、NMEAセンテンスの直前に配置される行の中には sブロック とよばれる送信元識別子がある。 調査の結果、NMEA 0183には Talker ID 、IEC61162-450には sブロック と、どちらのプロトコルにも資産特定に使えそうな識別子がありました。 検討2: 現状の船内ネットワーク調査 次に、現状の船内ネットワークがどのような構成なのか調査しました。 出典： JRC | 船内LANシステム あくまで私の予想ですが、GPSや潮流計などのシリアル出力データ（NMEA 0183）を入出力IF装置がIEC61162-450に変換しているのだと思われます。 検討3: 船舶pcapに対する現状のIDS製品の出力検証 次に、OsecT以外でのOT-IDSについても船舶プロトコルの対応状況を調査するために、以下の環境を想定したpcapファイルを用意し、2つのOT-IDS製品を使ってトラフィック解析しました。 結果として、今回の検証で使用したOT-IDSでは船舶で使われているプロトコルを解析できませんでした。このため、OsecTに適用可能な船舶プロトコルパーサーを新たに作成することにしました。 今回は、IEC61162-450の仕様書が手元にない＋インターネットにもあまり情報がないため、NMEA 0183のTalker IDを足掛かりにしてパーサーの作成を進めました。 Talker IDから資産を出力するZeek・Spicyパーサー作成・検証 OsecTのパーサーはZeek・Spicyで構成されています。 Zeek・Spicyとは Zeek トラフィックを解析して、IPアドレス、MACアドレスやプロトコルなどの情報をログとして出力するOSS。 Spicy Zeekで利用するC++のパーサーをC++で記述することなく簡易に生成するためのツール。 参照： 【日本初紹介】Zeek・Spicyの使い方まとめ - NTT docomo Business Engineers' Blog 今回は、ZeekとSpicyを使って船内ネットワークトラフィックを解析するためのパーサーを作成しました。 パーサー検証 作成したパーサーを使って、実際に船内ネットワークトラフィック（pcapファイル）を解析しました。 今回の検証では、2種類のNMEA 0183メッセージからTalker IDとSentence Typeの2つのフィールドを抽出し、Talker IDからセンサの特定ができているかを確認しました。以下は、ZeekとSpicyを用いて取得したNMEAメッセージの解析結果（logファイル抜粋）です。 GP の場合：GPS HE の場合：Heading （ジャイロコンパス） 上記の通り、Talker IDから発信元のセンサの特定（船舶資産の特定）が行えていることが分かります。 まとめ 今回は、既存のOT-IDSが船舶特有のプロトコルを解析できないという技術課題に対し、NMEA文を解析可能にする独自のパーサーを作成しました。 このパーサーではNMEA 0183内に含まれるTalker IDを足掛かりにして船舶機器を特定できるように設計しており、その情報をOsecTの可視化機能に組み込むことで、船舶特有の機器の可視化の強化や資産インベントリの補助につながる可能性があるというポジティブな成果を得ることができました。 イベントへの参加 IDSに関する業務以外にも複数の社内イベントや見学会に参加させていただきました。 SOC見学 NTTセキュリティのSOC（Security Operation Center）を見学しました。 当初はSOCに対して「脅威を検知したら即座にアラートを出す場所」というイメージを持っていました。しかし、実際にNTTセキュリティが提供するSOCサービスは一定期間トラフィックを収集・分析し、詳細なレポートとしてクライアントに提供するサービス形態であることを知りました。SOCには多様な形があるのだということを知ることができ、とても勉強になりました。 ドコモとのコラボ企画 イノベーションセンターで攻撃インフラの解明や撲滅に向けて活動しているPJがセキュリティ国際会議「Botconf」で発表した内容を、ドコモ本社のインターン生とともに聴講させていただきました。 発表内容は新ロシア派ハクティビストに関する詳細なレポートで、実際の脅威情報分析の手法を学ぶ貴重な機会となりました。また、ハクティビストの情報を安易に拡散することがかえって悪影響を生む可能性があると知り、情報との向き合い方を考え直すきっかけにもなりました。 イノベーションセンター内での業務共有 イノベーションセンター内で活動する5名のインターン生が、それぞれの担当業務や取り組み内容を共有しました。 普段は近くで作業していても、実際に何をしているのか詳しく知る機会は少なかったので、とても新鮮でした。みんな全然違うテーマやプロジェクトに取り組んでいて、「そんなことやってるんだ！」と驚くことも多く、チームの多様性やそれぞれの専門性の広さを実感できて面白かったです。 TechLunch NTTドコモビジネスで不定期開催されているランチ勉強会「TechLunch」に参加しました。 「勉強会」という言葉から堅い雰囲気を想像していましたが、実際は社員の方がそれぞれの興味関心に基づいて自由に発表する形式で、活発かつフランクな学びの場でした。 当日は社員1名とインターン生2名の発表があり、発表中はSlack上でリアルタイムに反応が飛び交い、オープンな雰囲気を感じました。 インターンシップを通して学んだこと インターンに参加する前は、GPSスプーフィングやAISスプーフィングといった船舶に多く見られる攻撃を検知するIDSを作成することを考えていました。しかし、OT-IDSの説明をしていただく中で「資産インベントリ補助」という新たな活用方向を見出すことができ、とても有意義な経験となりました。 また、製造業と船舶業界ではセキュリティ要件の枠組みが大きく異なることも学びました。将来的には製造業のセキュリティ分野に関わりたいと考えているため、今後は制度やプロトコルの違いを意識しながら、より幅広い視点で学びを深めていきたいと思います。 さらに、インターンを通じてセキュリティに関わる仕事の現実的な難しさと奥深さを実感しました。重大なインシデントが起こらないことは、セキュリティに携わる者として理想的な状態です。しかし一方で、何か大きな事件が起こらなければ業界全体の意識や整備が進みにくいという現状もあり、その間にある葛藤のようなものを感じました。 おわりに 今回のインターンはこれまで開発に携わった経験がなかった自分にとって、大きな一歩を踏み出す機会となりました。以前から「何かを作ってみたい」と思いながらも、具体的なアイデアが浮かばなかったり、時間を理由に後回ししてしまったりしていました。そんな中で、実際に手を動かして開発へ取り組めたことは、自分にとって非常に貴重な経験でした。 また、さまざまなイベントにも参加させていただき、自分の所属チームだけでなく他のチームやインターン生、さらにはドコモグループの社員の方々とも関わることができました。部署を越えて多くの方と交流する中で、職場の雰囲気や、人とのつながりの大切さを改めて実感しました。 最後になりましたが、このインターンに関わってくださった皆さまに心より感謝申し上げます。受け入れてくださったチームの皆さま、特にこのインターンで参加したIT/OT統合セキュリティPJおよびOsecT Tech PJの田中さん、前田さん、加島さんには2週間にわたり大変お世話になりました。温かくご指導くださり、本当にありがとうございました。