コグノスケ

Raspberry Pi 3のオーディオ 番外編2 - 仕様がわかりにくい

目次: Raspberry Pi

Raspberry Piは何も苦労せず動くので最高に便利ですけど、今回みたいにハードを直叩きしたいケースや、もう少し発展させてカーネルやドライバ開発用に使うことを考えると最高とは言えないですね。ちょっとボードとSoCの仕様がわかりにくいです。

ボードの仕様についてはSchematicsは公開されているものの、コネクタ近辺しか記述がなくGPIOピンと信号名の関係がわかりません。SoC（BCM283x）の仕様は歯抜け（Broadcomが一般人に見せたくない部分を削っている？）で意味不明な点がチラホラあります。

ARM系SoCの仕様公開

ちょっとイマイチ感はあるもののBroadcom並の情報公開をしてくれるのは、ありがたいことです。他のベンダーですとAllWinner, Amlogic, Rockchip辺りはセットトップボックス（STB）向けSoCの仕様を公開しています。非常にありがたいです。

仕様がオープンなだけあって、シングルボードコンピュータ（SBC）ではほぼこの3社のSoCが採用されます。あとはSamsungが採用されるくらいかな？

我らが日本ベンダーはSoCの情報をほぼ公開しません。SBCに日本のSoCが採用されることもほとんどありません。一般開発者が日本のSoCに触れる機会はあまりないです。仕様書がないのか、仕様公開するメリットを感じないのか、SBCのような小規模顧客を相手する余裕がないのか……、真相は知りませんけどちょっと残念ですね。

Raspberry Pi 3のオーディオ その7 - 高い音が出ない謎

目次: Raspberry Pi

その1 - 波形観察（2021年5月2日の日記参照）でRaspberry Pi 3では高い音が出ない現象が発生していました。この原因について調べたメモです。結論だけ言えば、PWMのハード的には出せるはず、でも根本原因はわかりませんでした。

PWMの素性を調べる

RasPi 3のHW PWMは出力Duty比を約20usごとにしか変更できないようです。確認方法は下記のようにしました。

• FIFOを有効にする: CTL（アドレス3f20c000）のUSEFx（ビット5と13）をセット
• DMAを無効にする: DMAC（アドレス3f20c008）のENAB（ビット31）をクリア
• FIFO空きをチェックする: STA（アドレス3f20c004）のFULL1（ビット0）がクリアされるまで待つ
• FIF1（アドレス3f20c018）に左右両チャネルのデータを書く、つまり2回Writeする
• FIFO空きチェックに戻る

これにより最速でPWM Duty比を変更することができます。下記は設定例です。

3f20c000  00002121 00000032  00000303 70776d30
3f20c010  00000800 00000000  70776d30 70776d30
3f20c020  00000800 00000800  70776d30 70776d30

Duty比は何でも良いですけど、0％100％が一番波形が見やすいと思います。この設定例で言えば0x000, 0x000と0x800, 0x800を交互に書きます。

最高速度でPWMのDuty比を切り替えたときの波形（黄色1chオーディオ出力、水色2ch PWM出力）

上記操作をしたときのPWMの波形を見ると約20usに一度しか波形が変わっていません。従ってRasPi 3のオーディオ出力は25kHz程度が上限です。

クロックとの関係

なぜ20usに一度しか変更できないか気になりますが、BCM2835の仕様書には一切記載がありません。試行錯誤してみた結果、clk_pwm_domainのクロック周波数を早くすると変更間隔も速くなるみたいです。

PWMのクロック設定はPWMCTL（アドレス3f1010a0）とPWMDIV（3f1010a4）にあるようです。これも仕様書には載っていませんので、Linuxのクロックドライバdrivers/clk/bcm/clk-bcm2835.cから読み取るしかありません。イマイチな仕様書ですね……。

3f1010a0  00000096 00005000  00000200 00000000
3f1010b0  00000000 00000000  0000636d 0000636d

クロックドライバの実装をみるに、CTLレジスタのビット0〜3がクロック源です。6なのでPLLD_COREというクロック源が選択されているようです。ビット9が分周するかどうかの設定です。上記の設定だと有効でも無効でもクロックの速度が変わりませんでした。デフォルトが5分周なのか？どこかでデフォルトの分周比を設定しているのか？どちらなのかは良くわかりません。

DIVレジスタは分周比を表すようです。仕様がイマイチわかりませんが固定小数点のようで、上位ビットが整数部分、下位12ビットが小数点以下を表しているようです。上記の設定だと5分周を意味しているはず。

最高速度でPWMのDuty比を切り替えたときの波形（黄色1chオーディオ出力、水色2ch PWM出力）分周比設定0x4000

分周比を0x4000に変更するとDuty比の変更速度も5/4倍になって、16us間隔になるので、上位ビットが整数部分であることはほぼ間違いなさそうです。ちなみにクロック制御部のレジスタを変更する際は、上位16ビットに0x5a00をORしてwrite しないといけません。

クロック源の周波数をdebugfsから確認すると500MHzみたいです。5分周だから100MHz駆動でしょうか？

# cat /sys/kernel/debug/clk/plld_core/clk_rate
500000006

チャネル数は2ch、FIFO幅は32bitですから、100MHz / 2 / 32 = 1.5625MHzつまり0.64us間隔ならまだわかるんですが、どうしてその30倍もある20us間隔になるんでしょう。さっぱりです。

残った謎

PWMの能力的にはギリギリ24kHzも再生できるはずですが、以前示したように通常のオーディオ出力経路（CPU → GPU → DMA → PWMの経路と思われる）を使うと、24kHzが再生できません。24kHzどころかfs = 32kHzでの16kHzでさえ再生不可能です。

PWM以外に、例えばGPU側で何かPCMデータを弄っているなど、高音を再生する際の制約があるのかもしれません。

Google先生と擬人化画像

艦これが流行ったあとくらいでしたか、Googleで「赤城」を検索すると空母やプラモデルの写真ではなく、艦これの絵がたくさん出てきて空母が出てこなくなりました。

ところが最近（？）Google検索結果の傾向が変わったとTwitterで見たので、試してみました。少なくとも「赤城」は空母が出てくるようになっています。

検索結果が擬人化キャラで埋め尽くされようと、世の中のトレンドの一部ですし間違いとはいえません。Googleはそんなこと気にしないのかと思っていましたが、あえて変えたってことは、内心ダメだこりゃって思ってたんですかね……？

その他の空母

赤城だけ特別扱いか？全部擬人化を排除したか？どちらか気になったので、試しに歴代空母をGoogle画像検索で探して、擬人化率を調べました。結構違っていて面白いです。

• 瑞鶴: 86％（19/22）
• 鳳翔: 92％（23/25）
• 赤城: 5.9％（1/17）
• 加賀: 67％（14/21）
• 龍驤: 88％（21/24）
• 蒼龍: 86％（19/22）
• 飛龍: 22％（4/18）
• 翔鶴: 91％（21/23）

以下はカウントに使った検索結果のキャプチャ画像です。

瑞鶴

鳳翔

赤城

加賀

龍驤

蒼龍

飛龍

翔鶴

赤城だけ特異的に擬人化の絵の結果が少ないです。何か限られた対象だけ特殊処理が働いているとかですかね？

分類

検索結果を見たら一目瞭然なんですが、擬人化絵と空母というか艦船の画像のアスペクト比が明らかに違います。

• 艦船: 横長（船は横長）
• 艦これ: 縦長（人間は縦長）

このような傾向があって1画面の表示結果が多い＝艦これ優勢、とわかります。その他にも、

• 艦船: 白黒（戦時中の古い船の写真が多い）
• 艦これ: 肌色（カラーの絵が多い）

ヒストグラムも割と特徴的です。この件に限ってはあまり難しいことを考えなくても、アスペクト比1:1を境界にして、色の特徴を見たら擬人化画像を分離できるなあなんてことを思いました。手元で分類するくらいなら十分ではないでしょうか。

ハックしようと思えばいくらでもできちゃうんで、検索エンジンの結果としてはダメですけど。

OpenCLのOSS実装poclを調べる その1 - 動かしてみよう

目次: OpenCL

最近OpenCLのオープンソース実装poclについて調べています。わかったことのメモです。

OpenCLはclGetDeviceIDs() を呼ぶときにデバイスの種類を指定します。KhronosのAPIドキュメント（clGetDeviceIDs(3) Manual Page）を見ると、デバイスの種類は4つ定義されています（DEFAULTとALLはデバイスの種類ではないので除外）。

CL_DEVICE_TYPE_CPU

OpenCLデバイスはホストCPUです。ホストCPUはシングルもしくはマルチコアCPUでOpenCLプログラムを実行します。OpenCLでマルチコアCPUの並列プログラムを書く人は居るのかな。あまり聞いたことがないですね……。

CL_DEVICE_TYPE_GPU

OpenCLデバイスはGPUです。GPUはOpenCLプログラムだけではなくて、OpenGLやDirectXのような3DグラフィクスAPIも高速に実行できます。GeForceやRadeonはこちらですね。

CL_DEVICE_TYPE_ACCELERATOR

OpenCLデバイスはアクセラレータ（例えばIBM CELL Blade）です。これらのデバイスはホストCPUとペリフェラル接続（PCI Expressなど）を使用して通信します。

CL_DEVICE_TYPE_CUSTOM

OpenCLデバイスはアクセラレータですが、OpenCL言語でカーネルを書くことができないタイプです。FPGAなんかが該当するんですかね？

これらのうちpoclがサポートしているのはCPUとGPUです。GPUはNVIDIAのCUDAとAMDのHSAに対応しているようです。全部LLVMがビルドしてくれるわけで、すごいよLLVMさん。ACCELERATORはテンプレート実装のみで、そのままでは動作しないので、注意が必要です。

ビルド、インストール

CPU（pthread版）、GPU（CUDA版）、ACCELERATORを有効にしたpoclのビルド、インストール方法は下記のとおりです。実際に動かすときはACCELERATORを無効にしてください。でないと初期化時にエラーが発生して動かないです。

$ cmake -G Ninja \
  -DCMAKE_INSTALL_PREFIX=`pwd`/_install \
  -DENABLE_CUDA=ON \
  -DENABLE_ACCEL_DEVICE=ON \
  ../

$ ninja
$ ninja install

基本的には必要なオプションがあればONにするだけですから、ビルドとインストールはそんなに難しくないはずです。

実行

実行方法はややクセがあります。OpenCLは実装がたくさんあるので、直接OpenCLライブラリをリンクするのではなく、ICD Loaderと呼ばれるライブラリ（2020年7月14日の日記参照）を間に噛ませることが多いです。

ソースコードは Oak Ridge大学のサイトとほぼ同じです。行数を減らしたのと、OpenCL 2.2に合わせて使うAPIを一部変えている程度です。

#define CL_TARGET_OPENCL_VERSION 220

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <CL/opencl.h>
 
// OpenCL kernel. Each work item takes care of one element of c
const char *kernelSource =                                        "" \
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable                    \n" \
"__kernel void vecAdd(__global double *a,                         \n" \
"                     __global double *b,                         \n" \
"                     __global double *c,                         \n" \
"                     const unsigned int n)                      \n" \
"{                                                               \n" \
"    // Get our global thread ID                                 \n" \
"    int id = get_global_id(0);                                  \n" \
"                                                                \n" \
"    // Make sure we do not go out of bounds                     \n" \
"    if (id < n)                                                 \n" \
"        c[id] = a[id] + b[id];                                  \n" \
"}                                                               \n" \
                                                                "\n" ;
 
int main(int argc, char *argv[])
{
    // Length of vectors
    int n = 100000;
    size_t bytes = n * sizeof(double);

    cl_platform_id cpPlatform;
    cl_device_id device_id;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;

    // Device input/output buffers
    cl_mem d_a, d_b;
    cl_mem d_c;

    // Allocate memory for each vector on host
    double *h_a = (double *)malloc(bytes);
    double *h_b = (double *)malloc(bytes);
    double *h_c = (double *)malloc(bytes);
 
    // Initialize vectors on host
    for (int i = 0; i < n; i++) {
        h_a[i] = sinf(i) * sinf(i);
        h_b[i] = cosf(i) * cosf(i);
    }

    // Number of work items in each local work group
    size_t localSize = 64;

    // Number of total work items - localSize must be devisor
    size_t globalSize = ceil(n / (float)localSize) * localSize;

    cl_int err;

    // Bind to platform
    err = clGetPlatformIDs(1, &cpPlatform, NULL);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Get ID for the device
    err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Create a context 
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    queue = clCreateCommandQueueWithProperties(context, device_id, 0, &err);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Create the compute program from the source buffer
    program = clCreateProgramWithSource(context, 1,
                                        (const char **)&kernelSource, NULL, &err);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Build the program executable
    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Create the compute kernel in the program we wish to run
    kernel = clCreateKernel(program, "vecAdd", &err);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Create the input and output arrays in device memory for our calculation
    d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
    d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
    d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);

    // Write our data set into the input array in device memory
    err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,
                               bytes, h_a, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }
    err = clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,
                               bytes, h_b, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Set the arguments to our compute kernel
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }
    err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }
    err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }
    err = clSetKernelArg(kernel, 3, sizeof(unsigned int), &n);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Execute the kernel over the entire range of the data set 
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
                                 &globalSize, &localSize, 0, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("%s:%d err:%d\n", __func__, __LINE__, err);
        return 0;
    }

    // Wait for the command queue to get serviced before reading back results
    clFinish(queue);

    // Read the results from the device
    clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,
                        bytes, h_c, 0, NULL, NULL);

    //Sum up vector c and print result divided by n, this should equal 1 within error
    double sum = 0;
    for (int i = 0; i < n; i++) {
        sum += h_c[i];
    }
    printf("final result: %f\n", sum / n);

    clReleaseMemObject(d_a);
    clReleaseMemObject(d_b);
    clReleaseMemObject(d_c);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);

    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

$ gcc a.c -g -O0 -Wall -lOpenCL -lm

★ICD Loaderにpoclのライブラリ名を教える必要がある

$ cat /etc/OpenCL/vendors/pocl_test.icd
libpocl.so.2.7.0

$ LD_LIBRARY_PATH=/path/to/pocl/build/_install/lib \
  ./a.out

final result: 1.000000

動作しました。良かった。

OpenCLのAPIはICD Loaderが提供し、OpenCLの実装は各ICDが提供します。App → ICD Loader → ICDという呼び出し関係です（詳しくは 2020年7月14日の日記参照）。一見すると煩雑ですが、アプリケーションはICDのことは知らなくても良いのが利点です。今回の例でいえばアプリケーションはlibpocl.soをリンクせずとも、poclの実装を使うことができます。

デバッグ

これだけだと面白くないし、何が動いているかすらわからないので、デバッグ出力を全開にして観察します。POCL_DEBUGという環境変数を使います。その他のデバッグ方法はPoCLのドキュメント（Debugging OpenCL applications with PoCL）が参考になります。

$ LD_LIBRARY_PATH=/path/to/pocl/build/_install/lib \
  POCL_DEBUG=all \
  ./a.out

** Final POCL_DEBUG flags: FFFFFFFFFFFFFFFF
[2021-03-13 07:05:14.074169726]POCL: in fn pocl_init_devices at line 571:
  |   GENERAL |  Installing SIGFPE handler...
[2021-03-13 07:05:14.128006157]POCL: in fn pocl_cuda_init at line 287:
  |   GENERAL |  [CUDA] GPU architecture = sm_61
[2021-03-13 07:05:14.128050269]POCL: in fn findLibDevice at line 560:
  |      CUDA | looking for libdevice at '/usr/lib/nvvm/libdevice/libdevice.10.bc'
...

動作しました。よきかな。タイムスタンプの日付がかなりズレますね……？ま、実害はないし良いか。

ベンチマーク - まとめリンク

目次: ベンチマーク

一覧が欲しくなったので作りました。

• 最速のyes
• 不安定なyes
• yesの高速化（パイプ限定）
• C言語でSEGV
• ELFバイナリでSEGVその1 - 92バイト
• ELFバイナリでSEGVその2 - 64バイト
• デノーマルフラッシュ
• デノーマルフラッシュの速度改善効果
• マイクロアーキとシングルスレッド性能
• FizzBuzzを速くする1（自作アルゴリズム）
• FizzBuzzを速くする2（高速アルゴリズムの紹介）
• FizzBuzzを速くする3（CPUを変えてみよう）
• FizzBuzzを速くする4（うまくいかないこともある）
• FizzBuzzを速くする5（SIMD最適化の紹介）
• FizzBuzzを速くする6（1桁落とし）
• FizzBuzzを速くする7（コンパイラによる違い）
• SIMDを使ったお手軽最適化 - その1
• SIMDを使ったお手軽最適化 - その2
• 100万回のHello, World!
• 100万回のHello, World! - バイナリサイズを削って遊ぼう
• 100万回のHello, World! - バイナリサイズを削って遊ぼう、究極編その1
• 100万回のHello, World! - バイナリサイズを削って遊ぼう、究極編その2

メモリクリアでおなじみmemset()関数の自作。

• memsetのベンチマーク（x86_64, Ryzen 7 2700編）
• memsetのベンチマーク（AArch64, Cortex-A72編）
• memsetのベンチマーク（AArch64, Cortex-A53編）
• memsetのベンチマーク（RISC-V 64, U74-MC編）
• muslのmemset関数
• memsetに一番効く最適化
• ぼくの考えた最強のmemset
• バイトをコピーするSIMD命令
• glibcのmemsetのクセ
• glibcのmemsetは強かった

Nクイーン問題の自作。

• C言語でNクイーン問題
• C言語でNクイーン問題、一意解
• C言語でNクイーン問題、高速化
• JavaScriptでNクイーン問題
• スクリプト言語始めました（PythonとRubyでNクイーン問題）
• インタプリタon Java VM（PythonとRubyでNクイーン問題）
• JavaScriptでNクイーン問題、その2

TRONが世界標準？？

Quoraのとある項目なぜTRON OSが「非常に優れていたが外圧で潰された」とか「組み込みで世界標準OSだ」とかいう誇張された伝説をいまだに信じている人が大勢いるのですか？ - Quora が話題になっていました。そんな話を信じている人が居るんですね。TRONが世界標準……私の知らない世界線でTRONが覇権を獲ったのでしょうか……。

松下電器（おそらく日本一のTRON推しの会社でした）に居た自分すら、そんなこと思ったことありませんでした。

その松下電器でさえBTRONはもちろんiTRONすらギブアップです。いまやレコーダーやテレビのOSはLinux/BSDカーネルを採用しています。iTRONアプリも残ってはいますが、過去資産の作り直しは面倒＆旨味がないのが理由だったと思います。

メモ: 技術系の話はFacebookから転記しておくことにした。

RISC-V 64 CPU第2号が我が家に来た

目次: RISC-V

SiFiveのHiFive Unmatchedを購入しました。現状、世界最速のLinuxが動作するRISC-V 64bit SoC とのことです。

ボードにはSDカードが付属しておりFreedom USDKという独自の環境がインストールされています。ボード上にはUSB接続のシリアル端子があり、電源を入れればLinuxが起動し、ユーザroot、パスワードsifiveでログインできるようになっています。

ぱっと見はPCと同じmini-ITXマザーボードですけど、バックパネルを見るとSDカードの差し込み口、USBシリアル用のmicroB端子が出ていて、どちらかというとSBC（シングルボードコンピュータ）です。PCっぽさがありません。

HiFive Unmatchedのバックパネル

本当はグラフィックカードを装着してGUIを使うべきですが、昨今のグラフィックカード品薄＆異常な値上がりのおかげで全く買う気が起きないので、しばらくシリアルコンソールで使おうと思います。

インストールされているカーネルは、
Linux unmatched 5.11.10 #1 SMP Wed Apr 7 17:37:34 UTC 2021 riscv64 riscv64 riscv64 GNU/Linux
でした。5.11はStableカーネルではあるものの、既にEOLです。まあ、開発用ボードだしこんなもんか。

購入時の同じ罠

Crowd Supplyから購入しました。本体 $679, 消費税が7,100円、合計で7万円くらいでした。HiFive Unleashedほどではないにせよ、SBCにしては良いお値段です。

UPSが米国→日本まで持ってきて、国内はクロネコヤマトが運びます。受け取りの際に、消費税を着払いでクロネコに払う必要があります。私は消費税のことを忘れていて、何だこの金は？？と混乱しました。Unleashedのときと全く同じでした。海外からものを買うことがほとんどなくて、消費税の存在をすぐ忘れちゃうんですよね……。

memsetのベンチマーク（RISC-V 64, U74-MC編）

目次: ベンチマーク

（参考）コード一式はGitHubに置きました（GitHubへのリンク）

Linuxが動くRISC-Vボードを買ったので、RISC-V 64でもmemsetをやってみました。環境はボードがSiFive HiFive Unmatchedで、SoCがSiFive Freedom U740で、コアがU74-MCという名前です。動作周波数は書いてないですね。OSはFreedom USDKというSiFive独自？の環境です。メモリはDDR4-2400のようです（Schematics hifive-unmatched-schematics-v3.pdfより）。

特徴的な点は、

• glibc C実装が最速
• アセンブラ実装がない（O2のglibc C実装と同じ性能）

あと個人的に残念だった点としては、U74コアの速度です。前世代のHiFive Unleashedに搭載されていたU54コアはCortex-A53の足下にも及びませんでした（2019年5月27日の日記参照）。

U74はCortex-A72レベルとまでは言いませんが、Cortex-A53は超えてくると期待していましたが、少なくともmemsetに関しては負けています。半分くらいの速度しか出ていません……。

gcc -O3 -fno-builtinの測定結果（SiFive U74-MC編）

gcc -O2 -fno-builtinの測定結果（SiFive U74-MC編）

OpenCLのOSS実装poclを調べる その2 - 独自アクセラレータのテンプレート実装を眺める

目次: OpenCL

CPUでもGPUでもないデバイスでOpenCLを動かすとしたらどうしたら良いでしょうか？答えとしては、その1で紹介したとおり、CL_DEVICE_TYPE_ACCELERATORを実装すれば良いです。が、イチから作るのはとっても大変です。

poclのテンプレート実装

素晴らしいことにpoclにはテンプレートらしき実装がpocl/lib/CL/devices/accelに用意されています。やりたいこととは微妙に違うことが後々わかりますが、イチから作るよりははるかにマシです。このテンプレートを改造しましょう。

テンプレートの名前はaccelでいかにもアクセラレータに見えますが、デバイスタイプはCL_DEVICE_TYPE_ACCELERATORではなくCL_DEVICE_TYPE_CUSTOMです。CUSTOMは「コンパイルが可能なデバイス」ではなく、ビルトインカーネルのみを実行するデバイスです。ユーザー定義のカーネルを実行することは考えられていません。

ユーザー定義カーネルが実行できることが独自アクセラレータの売りですから、何とかしてCUSTOMではなくACCELERATORになるように実装を改造する必要があります。これはなんとも先が長そうです……。

OpenCLのOSS実装poclを調べる その3 - デバイス数の取得処理

目次: OpenCL

独自アクセラレータのテンプレート実装pocl/lib/CL/devices/accelはデバイスタイプがCUSTOMになっているのが最大の難関ですが、その他にも色々問題があります。

最初に遭遇する問題はデバイス数を取得する処理のエラー処理が間違っていることです。現状のコードだとちょっと特殊な環境変数を渡さないと動きません。

// pocl/lib/CL/devices/devices.c

static unsigned device_count[POCL_NUM_DEVICE_TYPES];

...

cl_int
pocl_init_devices ()
{

...


  /* Init operations */
  for (i = 0; i < POCL_NUM_DEVICE_TYPES; ++i)
    {

...

      /* Probe and add the result to the number of probed devices */
      assert(pocl_device_ops[i].probe);
      device_count[i] = pocl_device_ops[i].probe(&pocl_device_ops[i]);    //★デバイス数を取得する★
      pocl_num_devices += device_count[i];
    }

...

  dev_index = 0;
  /* Init infos for each probed devices */
  for (i = 0; i < POCL_NUM_DEVICE_TYPES; ++i)
    {
      if (pocl_devices_init_ops[i] == NULL)
        continue;
      str_toupper (dev_name, pocl_device_ops[i].device_name);
      assert(pocl_device_ops[i].init);
      for (j = 0; j < device_count[i]; ++j)    //★デバイス数42億と誤解したまま処理しようとしてクラッシュする★
        {


// pocl/lib/CL/devices/accel/accel.cc

unsigned int pocl_accel_probe(struct pocl_device_ops *ops) {
  //★POCL_DEVICESという環境変数が見つからないとき、-1というエラー値を返す★
  //★本来エラー値である -1だが、デバイス数として解釈され42億になってしまう★
  int env_count = pocl_device_get_env_count(ops->device_name);
  return env_count;
}


// pocl/lib/CL/devices/devices.c

/**
 * Get the number of specified devices from environment
 */
int pocl_device_get_env_count(const char *dev_type)
{
  const char *dev_env = getenv(POCL_DEVICES_ENV);
  char *ptr, *saveptr = NULL, *tofree, *token;
  unsigned int dev_count = 0;
  if (dev_env == NULL)
    {
      return -1;    //★ここにくる★
    }
  ptr = tofree = strdup(dev_env);
  while ((token = strtok_r (ptr, " ", &saveptr)) != NULL)
    {
      if(strcmp(token, dev_type) == 0)
        dev_count++;
      ptr = NULL;
    }
  POCL_MEM_FREE(tofree);

  return dev_count;
}

このような実装になっておりaccelのデバイス数が42億（！）と解釈されてしまい、42億回デバイスを列挙しようとしてクラッシュします。バグのような気がしますけど、サンプル実装ですのであまり文句を言っても仕方ありません。

環境変数POCL_DEVICES="pthread -1 CUDA -1 accel 1" のようにデバイス数を明示的に渡せば回避可能です。最終的にはpocl_accel_probe() が正しくデバイス数を返すような実装を追加する必要があるでしょうが、この場は環境変数で切り抜けます。

OpenCLのOSS実装poclを調べる その4 - デバイスのパラメータを渡す環境変数

目次: OpenCL

引き続き、独自アクセラレータのテンプレート実装pocl/lib/CL/devices/accelの細かな問題を調べます。デバイス数の取得の問題を回避すると、次はデバイスのパラメータを渡す問題に遭遇します。

// pocl/lib/CL/devices/accel/accel.cc

void pocl_accel_init_device_ops(struct pocl_device_ops *ops) {

  ops->device_name = "accel";    //★★デバイス名はaccel
  ops->init = pocl_accel_init;

...


// pocl/lib/CL/devices/devices.c

cl_int
pocl_init_devices ()
{

...

  dev_index = 0;
  /* Init infos for each probed devices */
  for (i = 0; i < POCL_NUM_DEVICE_TYPES; ++i)
    {
      if (pocl_devices_init_ops[i] == NULL)
        continue;
      str_toupper (dev_name, pocl_device_ops[i].device_name);    //★★dev_nameはデバイス名を大文字に変換したACCELになる
      assert(pocl_device_ops[i].init);
      for (j = 0; j < device_count[i]; ++j)
        {

...

          /* Check if there are device-specific parameters set in the
             POCL_DEVICEn_PARAMETERS env. */
          POCL_GOTO_ERROR_ON (
              (snprintf (env_name, 1024, "POCL_%s%d_PARAMETERS", dev_name, j)    //★★環境変数名を生成する箇所
               < 0),
              CL_OUT_OF_HOST_MEMORY, "Unable to generate the env string.");
          errcode = pocl_devices[dev_index].ops->init (
              j, &pocl_devices[dev_index], getenv (env_name));

...

実装ではpocl_accel_init() にて環境変数の値をパースしてデバイスのパラメータを取得します。環境変数名はデバイス番号によって変化しますが、0番目のデバイスであればPOCL_ACCEL0_PARAMETERSという名前になります。環境変数名は上記にあるとおりpocl_init_devices() で決めています。

困ったことに環境変数が見つからないとabort() してしまうので、環境変数には最低でも何か1つ数値を渡す必要があります。なお1つ目の値はレジスタ領域のベースアドレスだと解釈されるようです。

他の実装（pthreadとcuda）は環境変数を使わないので、同様の問題は存在しません。最終的にはaccelも環境変数に頼らない実装に変えていく必要がありますが、今はそのままにしておきます。

今更HiFive UnleashedのJTAGに気づいた

目次: RISC-V

昔買ったHiFive Unleashedというボード、JTAGはちょっと特殊なコネクタが必要だと思っていて接続を諦めていました。ところが今日、ファンを掃除したあと動作確認をした際に、USB Serialと一緒にUSB JTAGも用意されていることに気づきました。

HiFive Unleashedはディスコンでもう手に入りませんし、後継機種のHiFive Unmatchedも買った今となっては、かなり今更感ありますが……。

ちょっと試したところ、簡単にOpenOCDが繋がりSMPモードにすると5コア（rv64imacなE51とrv64gcなU54 x 4）が見えました。

adapter speed 10000

adapter driver ftdi
ftdi_device_desc "Dual RS232-HS"
ftdi_vid_pid 0x0403 0x6010

ftdi_layout_init 0x0008 0x001b
ftdi_layout_signal nSRST -oe 0x0020 -data 0x0020

set _CHIPNAME riscv
jtag newtap $_CHIPNAME cpu -irlen 5 -expected-id 0x20000913

set _TARGETNAME $_CHIPNAME.cpu
target create $_TARGETNAME.0 riscv -chain-position $_TARGETNAME -rtos hwthread
target create $_TARGETNAME.1 riscv -chain-position $_TARGETNAME -coreid 1
target create $_TARGETNAME.2 riscv -chain-position $_TARGETNAME -coreid 2
target create $_TARGETNAME.3 riscv -chain-position $_TARGETNAME -coreid 3
target create $_TARGETNAME.4 riscv -chain-position $_TARGETNAME -coreid 4
target smp $_TARGETNAME.0 $_TARGETNAME.1 $_TARGETNAME.2 $_TARGETNAME.3 $_TARGETNAME.4

init
halt

SiFiveのSoCはFU540も後継のFU740も命令セットの違うコアを混載するのが好きですね。混載は良いとしてせめて同じ命令セットにしてほしかったです。単純なSMPしか持ってないOSだと制御できないじゃん……。

メモ: 技術系の話はFacebookから転記しておくことにした。加筆。

Zephyr on HiFive Unleashed

目次: Zephyr

JTAGが繋がった（2021年6月5日の日記参照）記念にZephyrをHiFive Unleashedで動かしてみました。使うハードウェアはUARTだけ、コードもJTAGでITMにロードして動作させるだけなら楽勝だろと思いきや、全然UARTから文字が出力されず1日掛かってしまいました。

*** Booting Zephyr OS build zephyr-v2.6.0-39-g2bf63134e8f0  ***
thread_a: Hello World from cpu 0 on hifive_unleashed!
thread_b: Hello World from cpu 0 on hifive_unleashed!
thread_a: Hello World from cpu 0 on hifive_unleashed!
thread_b: Hello World from cpu 0 on hifive_unleashed!
thread_a: Hello World from cpu 0 on hifive_unleashed!
thread_b: Hello World from cpu 0 on hifive_unleashed!

原因はコアのPLL設定が間違っていて、コアクロックから分周して作るUARTのボーレートもおかしな周波数になっていたからです。

FU540のリセット直後は外部クロック源（33.33MHz）をそのまま使って動きます。起動後PLLを1GHz（33.33 / 1 * 120 / 4 = 999.9MHz）に設定し、コアクロックをPLL側に切り替えなければなりません。

PLLの設定はFSBLという2段目のブートローダーが行うので、通常はOSが気にする必要はありません。しかしRTOSは大抵ブートローダーを経由しませんから、ブートローダーに隠れた設定も拾ってきて実装しないと動かないことがあります。

ブートローダーがあって当たり前のリッチ系SoCでRTOSを動かそうとすると、大抵この「暗黙のうちに設定されている何か」が抜け落ちて動かなかったりおかしくなったり、何かと面倒が起きます。

幸いなことにSiFiveのブートローダーはコードが公開されており、完全ブラックボックスのSoCと比べれば、難易度は低い部類です。ありがたいですね。

メモ: 技術系の話はFacebookから転記しておくことにした。加筆。

GTKのウインドウサイズ変更がシビアすぎる

目次: Linux

TwitterでGUIの話をしている人を見かけて思い出した話です。私はGNU/Debian LinuxのGUIをLightDM + GTKにして使っています。

GUIにこだわりはなく一時期デフォルトだったLightDMを使い続けているだけです。デフォルトになっただけあって、良くできていると思うし特に不満はないです……が、1点だけ言わせてもらえば、ウインドウサイズ変更の判定が厳しすぎませんか？

LightDM + GTKのウインドウサイズ変更判定の位置

特に左辺、上辺、右辺が1pxしか反応してくれないので、マウス操作が非常にシビアです。手が震えてきます。

操作が良くない

Twitterで上記の話をしていたところ、Alt + 右マウスボタンでウインドウサイズを変えられるよ、と教えてもらいました。今度からそうします。1pxに合わせるのは手が疲れる……。

OpenCLのOSS実装poclを調べる その5 - デバイスの初期化

目次: OpenCL

引き続き、独自アクセラレータのテンプレート実装pocl/lib/CL/devices/accelの細かな問題を調べます。次の問題はOpenCLの初期化です。clGetPlatformIDs() から初期化関数pocl_accel_init() に辿り着いたところでabort() が呼ばれクラッシュします。

  |   GENERAL |  accel: accelerator at 0x1000 with 0 builtin kernels
Could not open /dev/mem

テンプレート実装の意図としては /dev/memをopen() してメモリマップされたハードウェアのレジスタを読み書きしたいようです。今回は実際のハードウェア相手ではないので、レジスタの読み書きではなく /dev/memの代わりにバイナリファイルを開いてもらうように書き換えます。

// pocl/lib/CL/devices/accel/accel.cc

cl_int pocl_accel_init(unsigned j, cl_device_id dev, const char *parameters) {

...

  POCL_MSG_PRINT_INFO("accel: accelerator at 0x%zx with %zu builtin kernels\n",
                      D->BaseAddress, D->SupportedKernels.size());

  int mem_fd = open("/dev/mem", O_RDWR | O_SYNC);
  if (mem_fd == -1) {
    POCL_ABORT("Could not open /dev/mem\n");  //★★このabortでクラッシュ
  }

...

ファイル名を書き換えて突破するとデバイスが持っているメモリのサイズを取得し、メモリマップしようとする部分で怒られます。

  |   GENERAL |  accel: accelerator at 0x1000 with 0 builtin kernels
a.out: ../lib/CL/devices/accel/accel.cc:196: void MMAPRegion::Map(size_t, size_t, int): Assertion `Data != MAP_FAILED && "MMAPRegion mapping failed"' failed.

サイズを取得している箇所は下記のとおりです。

// pocl/lib/CL/devices/accel/accel.cc

cl_int pocl_accel_init(unsigned j, cl_device_id dev, const char *parameters) {

...

  uint32_t ctrl_size = D->ControlMemory.Read32(ACCEL_INFO_CTRL_SIZE);
  uint32_t imem_size = D->ControlMemory.Read32(ACCEL_INFO_IMEM_SIZE);
  uint32_t dmem_size = D->ControlMemory.Read32(ACCEL_INFO_DMEM_SIZE);
  uint32_t pmem_size = D->ControlMemory.Read32(ACCEL_INFO_PMEM_SIZE);

  uint32_t max_region =
      std::max(std::max(ctrl_size, imem_size), std::max(dmem_size, pmem_size));

  D->InstructionMemory.Map(D->BaseAddress + max_region, imem_size, mem_fd);

... 

バイナリファイルを書き換えて何か適当な値が読めるようにしてやりすごすか、面倒ならばD->ControlMemory.Write32(ACCEL_INFO_CTRL_SIZE, 0x2000); のように固定値を書いておくと次に進みます。今は実際のデバイスが相手ではないので、とりあえず先に進めて後で考えましょう。