コグノスケ


link 未来から過去へ表示(*)  link 過去から未来へ表示

link もっと前
2021年5月29日 >>> 2021年5月16日
link もっと後

2021年5月29日

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編)

編集者:すずき(2023/09/24 08:54)

コメント一覧

  • コメントはありません。
open/close この記事にコメントする



2021年5月28日

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のときと全く同じでした。海外からものを買うことがほとんどなくて、消費税の存在をすぐ忘れちゃうんですよね……。

編集者:すずき(2021/06/28 15:29)

コメント一覧

  • コメントはありません。
open/close この記事にコメントする



2021年5月23日

TRONが世界標準??

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

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

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

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

編集者:すずき(2022/04/05 12:00)

コメント一覧

  • コメントはありません。
open/close この記事にコメントする



2021年5月22日

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

目次: ベンチマーク

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

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

Nクイーン問題の自作。

編集者:すずき(2024/02/27 01:56)

コメント一覧

  • コメントはありません。
open/close この記事にコメントする



2021年5月16日

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を無効にしてください。でないと初期化時にエラーが発生して動かないです。

poclのビルドとインストール
$ 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を一部変えている程度です。

OpenCLのテスト用ソースコード

#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;
}
poclの実行
$ 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)が参考になります。

poclのデバッグ出力(CUDA版)
$ 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'
...

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

編集者:すずき(2023/09/24 11:57)

コメント一覧

  • コメントはありません。
open/close この記事にコメントする



link もっと前
2021年5月29日 >>> 2021年5月16日
link もっと後

管理用メニュー

link 記事を新規作成

<2021>
<<<05>>>
------1
2345678
9101112131415
16171819202122
23242526272829
3031-----

最近のコメント5件

  • link 20年6月19日
    すずきさん (04/06 22:54)
    「ディレクトリを予め作成しておけば良いです...」
  • link 20年6月19日
    斎藤さん (04/06 16:25)
    「「Preferencesというメニューか...」
  • link 21年3月13日
    すずきさん (03/05 15:13)
    「あー、このプログラムがまずいんですね。ご...」
  • link 21年3月13日
    emkさん (03/05 12:44)
    「キャストでvolatileを外してアクセ...」
  • link 24年1月24日
    すずきさん (02/19 18:37)
    「簡単にできる方法はPowerShellの...」

最近の記事3件

  • link 24年4月17日
    すずき (04/18 22:44)
    「[VSCodeとMarkdownとPlantUMLのローカルサーバー] 目次: LinuxVSCodeのPlantUML Ex...」
  • link 23年4月10日
    すずき (04/18 22:30)
    「[Linux - まとめリンク] 目次: Linuxカーネル、ドライバ関連。Linuxのstruct pageって何?Linu...」
  • link 20年2月22日
    すずき (04/17 02:22)
    「[Zephyr - まとめリンク] 目次: Zephyr導入、ブート周りHello! Zephyr OS!!Hello! Ze...」
link もっとみる

こんてんつ

open/close wiki
open/close Linux JM
open/close Java API

過去の日記

open/close 2002年
open/close 2003年
open/close 2004年
open/close 2005年
open/close 2006年
open/close 2007年
open/close 2008年
open/close 2009年
open/close 2010年
open/close 2011年
open/close 2012年
open/close 2013年
open/close 2014年
open/close 2015年
open/close 2016年
open/close 2017年
open/close 2018年
open/close 2019年
open/close 2020年
open/close 2021年
open/close 2022年
open/close 2023年
open/close 2024年
open/close 過去日記について

その他の情報

open/close アクセス統計
open/close サーバ一覧
open/close サイトの情報

合計:  counter total
本日:  counter today

link About www.katsuster.net
RDFファイル RSS 1.0

最終更新: 04/18 22:44