Ubuntu24.04のCUDAでC++23を使う設定

Ubuntu24.04環境でC++23のCUDAコードを書こうとしたら環境構築が結構面倒だったのでメモ。

nvidia/cuda:13.3.1-cudnn-devel-ubuntu24.04 のDockerイメージで確認した。

※C++23をフルサポートしているのはCUDA 13.3 以降らしい

パッケージの追加

2026/10時点でUbuntu24.04にはg++ 13が入っている。std::ranges::toがヘッダに入ってなかったりするのでg++ 14を追加する。

sudo apt install g++-14

nvccのコマンドライン引数

nvccでコンパイルするときには -std=c++23 の指定に加えて -ccbin で上記のg++ 14を指定する(-ccbinを指定しないとg++ 13の方が使われてしまう)。

nvcc -std=c++23 -o cpp23.exe -ccbin /usr/bin/g++-14 cpp23.cu

お試し用のソースコード cpp23.cu は以下のコードを使った(内容に意味はない)。

#include <vector>
#include <ranges>
#include <print>

int main(void)
{
    constexpr int length = 32 * 4;
    auto values = std::views::iota(1, 1 + length) | std::ranges::to<std::vector<std::uint8_t>>();

    for (auto v : values) {
        std::println("{:3d}", v);
    }

    return 0;
}

CMakeLists.txt

CMakeLists.txtは CMAKE_CUDA_HOST_COMPILER と target_compile_options の指定が必要。

cmake_minimum_required(VERSION 3.24)    # CMAKE_CUDA_ARCHITECTURESのnative

set(CMAKE_CUDA_ARCHITECTURES native)
set(CMAKE_CUDA_HOST_COMPILER /usr/bin/g++-14)  # nvccが担当しない箇所のコンパイラにg++ 14を指定する(-ccbinとして渡される)
project(cpp23_cuda LANGUAGES CXX CUDA)

find_package(CUDAToolkit REQUIRED)

add_executable(${PROJECT_NAME})
target_sources(${PROJECT_NAME} PRIVATE
    cpp23.cu
)
target_compile_options(${PROJECT_NAME} PRIVATE
    $<$<COMPILE_LANGUAGE:CUDA>:-std=c++23>  # コマンドライン引数の -std=c++23 を直接指定する
)

CMAKE_CUDA_HOST_COMPILERは-ccbin相当の指定になる。

問題は -std=c++23 の方で、以下のように set_target_properties を使おうとしても今回の環境では非対応のようでエラーになってしまう。

# こっちの書き方だとダメ
set_target_properties(${PROJECT_NAME} PROPERTIES
    CUDA_STANDARD 23
    CUDA_STANDARD_REQUIRED ON
)

cmake実行時に以下のエラーになる。

CMake Error in CMakeLists.txt: Target "cpp23_cuda" requires the language dialect "CUDA23" (with compiler extensions). But the current compiler "NVIDIA" does not support this, or CMake does not know the flags to enable it.

CUDAの__shfl_xor_syncとfloat版atomicMaxのメモ

久しぶりにCUDAコードを書こうとしたら色々変わってたり忘れていたので自分用にメモ

スレッド間通信で__shfl_xor_sync()を使う場合のメモ

float y = __shfl_xor_sync(0xffffffff, x, mask, 8);

第一引数は基本的にall Fでよさそう。AIに質問したら分岐でシャッフル処理に参加しないスレッドがいる場合に該当するビットを0にしてデッドロックなどを防ぐための引数とのこと。

第二引数を渡した値が相手先のスレッドで戻り値として返ってくる。

通信相手のスレッドは第三引数を使って threadIdx.x ^ mask で決まるらしい。

第四引数は2の累乗を指定可能で上記のように8を指定するとスレッドが8個ずつのサブグループになってそれぞれで通信することになる。

例えば mask=0x02 だとthreadIdx.xが以下の組み合わせ同士で値の交換が起きる。(スレッドIDは2進数で表現)

# mask は2進数で 010b

スレッド 000b ┓交換
スレッド 010b ┛

スレッド 001b ┓交換
スレッド 011b ┛

スレッド 100b ┓交換
スレッド 110b ┛

スレッド 101b ┓交換
スレッド 111b ┛

ワープ内に限られるけどsharedメモリとかを使わずにスレッド間でデータ交換できるのは便利。

float版atomicMax

atomicMax()は整数型しか用意されてなさそう。AIに聞いたらatomicCAS()を使って整数値の交換で対処しろ、とのこと。

※最小比較を逆転させるとatomicMin()も実装できる

__device__ __forceinline__ float atomicMax_float(float* addr, float value)
{
    float old = *addr;  // 最初に今の値を取り出す
    float assumed;

    if(old >= value) {
        return old;  // すでに最大値が入っていれば何もせず終了
    }

    // 値の交換が成功するまでループする
    do {
        assumed = old;
        if (assumed >= value) {
            break;  // ほかのスレッドが自分よりmaxな値で更新していれば終了(ループ2周目以降に意味のある処理)
        }
        // 値の交換を試みる(float値をint値に偽装して処理する)
        old = __int_as_float(atomicCAS((unsigned int*)addr, __float_as_int(assumed), __float_as_int(value)));
    } while (old != assumed);  // ほかのスレッドが値を更新した場合は old != assumed になるのでリトライする

    return old;
}

__float_as_int()や__int_as_float()はビット列を維持したままデータ型解釈のみを読み替えるらしい。

VSCodeで2段階リモート接続時の設定JSONの場所

毎回忘れて困るのでメモ。

  • Windows環境にVSCodeが入っていて
  • 以下どちらかのルートで2段階でremote接続している状況(VSCode経由ではコンテナを起動していない)
    • WSL > 既存のDockerコンテナにアタッチ
    • SSH > 既存のDockerコンテナにアタッチ

リモート接続の設定JSONは以下の場所に入っている。

C:\Users\<ユーザー名>\AppData\Roaming\Code\User\globalStorage\ms-vscode-remote.remote-containers\imageConfigs

Dev Containersでosrf/ros:humble-desktopと表示されているコンテナの設定ファイル名は「osrf%2fros%3ahumble-desktop.json」

どうしてもWSLg環境でAWSIMを動かしたい時に見る記事

注意事項

今回のやり方は不具合が発生する可能性が高いため自力で問題切り分け・解決が可能な方に向けた情報共有です。広くおススメできる方法ではありません。

達成したこと

Windows環境でAWSIMを起動し自動運転AIチャレンジのサンプルがそれらしく動くことを確認した。

WSLg上でAWSIMを起動する様子

※公式サイトの推奨環境ページに記載されている通り自動運転AIチャレンジの環境はUbuntu22.04を前提としている。今回は無理やりWindowsに環境作ったという報告

環境の構成

自動運転AIチャレンジの環境はシミュレータであるAWSIMと制御系のAutowareからなっている。(と思う。まだ始めたばかりでよく分かっていない)

Autowareの方はWindowsのWSL2環境でUbuntu22.04を用意して大会公式が用意してくれている環境構築手順そのままで構築すればOK。

問題はAWSIMの方でVulkanを使っているのでWSLg環境でDockerコンテナ内から実行することができない(エラーになる)。
※大会公式の環境ではUbuntu22.04上にDockerコンテナを作る形になっている

対策としてWSLgのIssueにあるようにkisak-mesaのドライバを入れてVulkanをD3D12に読み替えてバイパスしてもらうようにすれば、とりあえずたいていのVulkanアプリは動く。
ただし、自分が以前試した限りUbuntu24.04しかこの方法は使えなかった。参考

ただし、AWSIMはUbuntu22.04, ROS2 Humbleを前提にビルドされているので無理やり何とかする必要がある。

以上をまとめると下図のようにWindows上で22.04と24.04の2つのWSL環境を動かしAutowareはDockerコンテナ内、AWSIMはコンテナを使わず24.04上で直接実行する。

構成図

AWSIMを無理やり動かすための対策内容

WSL2上のUbuntu24.04でVulkanアプリが動くようにする

上述したWSLgのIssueの通りでOK。

sudo add-apt-repository ppa:kisak/kisak-mesa
sudo apt update
sudo apt upgrade

依存ライブラリを無理やりインストール

AWSIMの依存ライブラリをUbuntu22.04、ROS2 humbleからUbuntu24.04環境に持ってくる。

※バージョンが一致しないものを無理やり持ってくるので動く保証はない点に注意。一見動いているように見えても正しく動いていない可能性もある

Ubuntu22.04 から

足りないのは以下のファイル。/usr/lib/x86_64-linux-gnu ディレクトリに入っているのでWSL2のUbuntu22.04環境からコピーしてくる。

  • libconsole_bridge.so.1.0
  • libfmt.so.8.1.1
  • liborocos-kdl.so.1.5.1
  • libspdlog.so.1.9.2

24.04環境にコピーしたら以下のシンボリックリンクを張る。

libfmt.so.8 -> libfmt.so.8.1.1
liborocos-kdl.so.1.5 -> liborocos-kdl.so.1.5.1
libspdlog.so.1 -> libspdlog.so.1.9.2

ROS2 humble から

公式のリリースページからros2-humble-日付-linux-jammy-amd64.tar.bz2をダウンロードしてきてtar.bz2ファイルの中から以下のファイルをUbuntu24.04環境に入れる。

  • libclass_loader.so
  • librcl_action.so
  • librclcpp_action.so

入れる場所は分かりやすいように /opt/ros/humble/lib に入れた。

※Ubuntu24.04に対応するのはROS2 jazzyなので24.04環境に本来humbleディレクトリは存在しない

libpython3.10.so.1.0

Python3.10はPPAがあるので以下の手順ですんなり入る。

sudo add-apt-repository ppa:deadsnakes/ppa
sudo apt install libpython3.10

AWSIMファイル転送

Ubuntu22.04側で大会公式の環境構築手順を実施すればAWSIMのファイルもダウンロードされているので以下のファイルをUbuntu24.04側へコピーする。

  • aichallenge-racingkart/aichallenge/simulator/AWSIM
    • AWSIMのファイル一式。ディレクトリごとコピーする
  • aichallenge-racingkart/vehicle/cyclonedds.xml
    • CycloneDDSの設定ファイル。Autoware側でも同じファイルを使うので22.04側から削除しないように注意

ファイルの置き場所は任意で問題ないはず。今回はユーザのホームディレクトリ直下に置いた。

実行手順

AWSIM起動

WSL2インスタンスを起動するたびに以下を1回実行する。(Autoware側はDockerコンテナで実行してくれている)

# ROS2のディスカバリー通信が届くようにloデバイスをマルチキャストOnに設定する
sudo ip link set multicast on lo
# CycloneDDSでエラーにならないようにバッファサイズを増やす
sudo sysctl -w net.core.rmem_max=2147483647 >/dev/null

make dev手順の代わりに以下のコマンドで同等の手順になる。

CYCLONEDDS_URI="file://$HOME/cyclonedds.xml" \
ROS_LOCALHOST_ONLY=0 \
RMW_IMPLEMENTATION=rmw_cyclonedds_cpp \
ROS_DOMAIN_ID=1 \
LD_LIBRARY_PATH="$HOME/AWSIM/AWSIM_Data/Plugins:/opt/ros/humble/lib" \
$HOME/AWSIM/AWSIM.x86_64 \
  --start-mode off \
  --vehicles 1 \
  --laps 600 \
  --timeout 60000000

LD_LIBRARY_PATHにはAWSIMのPluginsディレクトリと無理やり持ってきたHumbleのライブラリ置き場を指定する。

AWSIMのコマンドライン引数に何を指定すれば良いのかは aichallenge-racingkart/aichallenge/run_simulator.bash を見れば分かる。

Autoware起動

make dev手順の代わりに以下のコマンドで同等の手順になる。

make autoware-simulator

Appendix デバッグ・問題切り分け作業内容

AWSIM エラー内容の確認

AWSIM実行中のエラーは以下のログファイルに出力されている。

/$HOME/.config/unity3d/TIERIV/AWSIM/Player.log

以下のように問題のあるPluginがログに残っている。

Unable to preload the following plugins:
    libfastrtps.so
    libfastrtps.so.2.6.3
    librmw_fastrtps_cpp.so
    librmw_fastrtps_shared_cpp.so

※自動運転AIチャレンジではCycloneDDSが使われているので問題ないが、今回の対策を実行しても上記のFastRTPS関連のエラーは残ったままなので注意

ROS2通信が上手くいかない場合は先にログを見た方がよさそう。

足りないライブラリの探し方

今回AWSIMのPluginがロードできない問題は依存ライブラリが足りないケースのみだったためlddスクリプトでロードできていない.soを1つずつ調べて足りないライブラリを探した。

# libRobotecGPULidar.soの依存関係を調べる例
LD_LIBRARY_PATH="$HOME/AWSIM/AWSIM_Data/Plugins:/opt/ros/humble/lib" \
ldd AWSIM/AWSIM_Data/Plugins/libRobotecGPULidar.so

結論

コンテナを使わないUbuntu24.04環境に無理やり依存ライブラリを持ってくることでWindows環境でも自動運転AIチャレンジの環境を作ることができた。
ただし、ライブラリの互換性問題で一見動いているように見えても正しく動かない可能性があるため広くおススメできる状態にはならなかった。
この課題は自動運転AIチャレンジの環境がUbuntu24.04に移行すれば解消しそうに思える(※)ので現状は別途Ubuntu22.04環境を用意するのが良いと思わる。

※AWSIMがUbuntu24.04前提ならmesaドライバのVulkan対応のみで動くため

以上

専用ライブラリを使わずに点群をglTFファイルに書き出してみた

CARLAシミュレータで出したLiDAR点群を可視化したいなぁと思ってnumpy.ndarrayの3次元座標をglTFファイルに出力するPythonスクリプトを書いてみた。

glTFファイルはテキスト形式だとJSONフォーマットらしくて、ざっくりした説明は公式のAPIリファレンスガイドPDFに書いてある。

今回は点群1つ1つに色を付けるのではなく、種類ごとに色を分けるだけにとどめた。 そうするとmeshesが1個で中に複数のprimitives[]を持てばよいらしい。

meshes[]
  ┣ primitives[0]
  ┃    ┣ attributes.POSITION  # accessors[]のインデックス番号(0始まり)を入れる
  ┃    ┗ material  # materials[]のインデックス番号(0始まり)を入れる
  ┗ primitives[1]

座標データは accessors > bufferViews > buffers と参照関係になるらしい。そしてbuffers[]のuriにはbase64エンコードした座標値を入れる。テキストのリストを書けないのは残念。テキストエディタで座標値変更できないね。手作業で座標値を更新できないならライブラリでglTF出力すれば良いのでは

ソースコード

from __future__ import annotations
import copy
import json
import base64
import numpy as np


GLTF_BASE_DICT = {
    'asset': {
        'version': '2.0',
        'generator': 'glTF Point Cloud from python script'
    },
    'scene': 0,
    'scenes': [{'nodes': [0]}],
    'nodes': [{'mesh': 0}],
    'meshes': [{'primitives': []}],  # ここに複数入る
    'materials': [],                 # ここに複数入る
    'accessors': [],                 # ここに複数入る
    'bufferViews': [],               # ここに複数入る
    'buffers': [],                   # ここに複数入る
}


def _get_material(color: tuple[int,int,int]) -> dict:
    """
    {
      "pbrMetallicRoughness": {
        "baseColorFactor": [color[0]/255, color[1]/255, color[2]/255, 1.0],
        "metallicFactor": 0.0,
        "roughnessFactor": 1.0
      }
    }
    """
    rgba = [value / 255 for value in color] + [1.0]
    return dict(pbrMetallicRoughness=dict(baseColorFactor=rgba, metallicFactor=0.0, roughnessFactor=1.0))


def _get_buffer(points: np.ndarray) -> dict:
    """
    {
      "byteLength": ...,
      "uri": "data:application/octet-stream;base64,..."
    }
    """
    binary_data = points.astype(np.float32).tobytes()
    base64_str = base64.b64encode(binary_data).decode('utf-8')
    uri = 'data:application/gltf-buffer;base64,' + base64_str
    return dict(byteLength=len(binary_data), uri=uri)


def _get_view(points: np.ndarray, buffer_index) -> dict:
    """
    {
      "buffer": buffer_index,
      "byteOffset": 0,
      "byteLength": ...,
      "target": 34962  # OpenGL buffer target?
    }
    """
    return dict(buffer=buffer_index, byteOffset=0, byteLength=len(points.astype(np.float32).tobytes()), target=34962)


def _get_accessor(points: np.ndarray, view_index) -> dict:
    """
    {
      "bufferView": view_index,
      "byteOffset": 0,
      "componentType": 5126,  # GL_FLOAT
      "count": 3,
      "type": "VEC3",
      "max": [1.0, 1.0, 0.0],
      "min": [0.0, 0.0, 0.0]
    }
    """
    count = points.shape[0]  # pointsのshapeは(N, 3)
    max_value = np.max(points, axis=0)
    min_value = np.min(points, axis=0)
    return dict(bufferView=view_index, byteOffset=0, componentType=5126, count=count, type='VEC3', max=max_value.tolist(), min=min_value.tolist())


def _get_primitive(attr_index, material_index) -> dict:
    """
    {
        "attributes": {
        "POSITION": attr_index
        },
        "mode": 0,  # POINTS
        "material": material_index
    }
    """
    return dict(attributes=dict(POSITION=attr_index), mode=0, material=material_index)


def add_points(gltf_dict: dict, points: np.ndarray, color: tuple[int,int,int]):
    materials: list = gltf_dict['materials']
    material_index = len(materials)
    materials.append(_get_material(color))

    buffers: list = gltf_dict['buffers']
    buffer_index = len(buffers)
    buffers.append(_get_buffer(points))

    bufferViews: list = gltf_dict['bufferViews']
    view_index = len(bufferViews)
    bufferViews.append(_get_view(points, buffer_index=buffer_index))

    accessors: list = gltf_dict['accessors']
    accessor_index = len(accessors)
    accessors.append(_get_accessor(points, view_index=view_index))

    primitives: list = gltf_dict['meshes'][0]['primitives']
    primitives.append(_get_primitive(attr_index=accessor_index, material_index=material_index))


def write_gltf(point_list, color_list, output_filepath):
    gltf_dict = copy.deepcopy(GLTF_BASE_DICT)
    for point, color in zip(point_list, color_list):
        add_points(gltf_dict, point, color)

    with open(output_filepath, mode='w', encoding='utf-8', newline='') as f:
        json.dump(gltf_dict, f, indent=4)


if __name__ == '__main__':
    # お試し
    r = np.array([[0, 0, 0], [1, 0, 0], [0, 1, 0], [1, 1, 0]])
    g = np.array([[0.5, 0, 0.125], [0, 0.5, 0.125], [0.5, 0.5, 0.125], [0, 0, 0.125]])
    b = np.array([[0.25, 0, 0.25], [0, 0.25, 0.25], [0.25, 0.25, 0.25], [0, 0, 0.25]])
    write_gltf([r, g, b], [(255, 0, 0), (0, 255, 0), (0, 0, 255)], 'test.gltf')

お試し実行結果

glTFビューアで見たらそれっぽい位置関係で表示できたのでたぶん大丈夫?

ビューアで閲覧した画像

点が小さくて視力検査みたいになっちゃった。

PyAVで動画切り出し

PyAVでmp4ファイルの一部を切り出した(例えば開始40秒から20秒間の動画を切り出す)新しい動画を作ろうとしたけど、なかなかうまくいかなかったのでソースコードを残しておく。

使用したバージョン

  • av 16.1.0

ポイント

VideoFrameの以下のattributeを再設定しないと狙った通りの動画が生成されなかった。

attribute 設定値
dts None
pts None
time_base Fraction(1, rate)

rateは動画のフレームレート。30fpsならrate=30

ソースコード

import av
from av.container import InputContainer
from av.video import VideoFrame
from fractions import Fraction


def clip(in_file, out_file, start, end):  # in_fileのstart秒からend秒までをout_fileへ切り出す
    rate = 30  # 本来はin_stream.codec_context.framerate か in_stream.codec_context.rate から自動設定したかったが今回試した環境だとどちらもNoneが返るためハードコードした
    in_container: InputContainer
    with av.open(in_file) as in_container, av.open(out_file, mode='w') as out_conteiner:
        in_stream = in_container.streams.video[0]
        # out_stream = out_conteiner.add_stream_from_template(in_stream)  # 自分の環境だと UnknownCodecError になるのでadd_stream_from_template()メソッドではダメだった
        out_stream = out_conteiner.add_stream('h264', rate, width=in_stream.width, height=in_stream.height)
        in_container.seek(int(start / in_stream.time_base), stream=in_stream)
        for pkt in in_container.demux(in_stream):
            frame: VideoFrame
            for frame in pkt.decode():
                if frame.time < start:
                    continue
                if frame.time > end:
                    break
                frame.dts = None
                frame.pts = None
                frame.time_base = Fraction(1, rate)
                for packet in out_stream.encode(frame):
                    out_conteiner.mux(packet)
        # Flush stream
        for packet in out_stream.encode():
            out_conteiner.mux(packet)

VSCodeからAWSのSSM接続経由でSSHする時にインスタンスID指定しない方法

たぶんできそうな案が出来たのでめも。動作確認できたらタイトルから「案」を削除予定。 動作確認できた。プロファイル指定漏れがあった個所を修正済み。

結論

以下のモチベーションで色々試したりGoogle Geminiにアドバイスを求めたりしてたどり着いた方法。

以降のawsコマンドに--profile my_profile引数を渡している箇所は自分のプロファイル名に読み替えること。

# C:\Users\ユーザ名\.ssh\config
Host my_auto_ec2
    # EC2インスタンスのNameタグに付けた名前
    HostName my-instance-name
    # EC2インスタンスへログインするユーザ名
    User my-ec2-user
    ProxyCommand C:\\Windows\\System32\\WindowsPowerShell\\v1.0\\powershell.exe -Command "$id = (aws ec2 describe-instances --profile my_profile | jq -r --arg name Name --arg value %h '.Reservations[].Instances[] | select(.Tags[].Key==$name) | select(.Tags[].Value==$value) | .InstanceId') ; aws ssm start-session --profile my_profile --target $id --document-name AWS-StartSSHSession --parameters portNumber=%p"

この設定でssh my_auto_ec2で接続できる。

ProxyCommand に書いたPowerShellコマンドの簡単な解説

  • powershell.exe -Command "コマンド1 ; コマンド2"形式でPowerShellにコマンドを2つ続けて実行させる
  • "$id = (コマンド)"でコマンド実行結果を変数idに入れる
  • jqlang/jqは--arg 変数名 変数の内容形式のコマンドライン引数でクエリ文字列の内部で変数を使える(複雑なエスケープ指定を避けられる)

おまけ

表題と関係ないがjqlang/jqの変数を活用すると以下のような複雑な処理も書けた。

# Nameタグが「my-instance」から始まるインスタンスの名前, インスタンスタイプ, 状態をそれぞれカンマ区切りで表示する
aws ec2 describe-instances --profile my_profile | jq -r --arg name Name --arg value 'my-instance' '.Reservations[].Instances[] | select(.Tags[].Key==$name) | select(.Tags[].Value | startswith($value)) | [(.Tags[] | select(.Key==$name).Value), .InstanceType, .State.Name]' | jq -r --arg sep ', ' 'join($sep)'

出力例

my-instance-micro, t3.nano, running
my-instance-large, t3.large, stopped
my-instance-xlarge, t3a.xlarge, stopping