コグノスケ

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

過去から未来へ表示

もっと前

2021年7月6日 >>> 2021年6月23日

もっと後

2021年7月6日

permalink

編集する

OpenOCDとHiFive UnleashedのSPI Flashその2 - OpenOCDのSPI Flashドライバの仕組み（書き込み編、通常版）

目次: OpenOCD

SPI Flashにどのように書き込みするか見ていきます。OpenOCDのSPI Flash書き込みは大きく2つの方法に分けられます。ヘルパープログラムを使う方法と、SPIを直接制御する方法です。順に紹介します。

ヘルパープログラム

後述しますがJTAGからSPIインタフェースを制御すればSPI Flashの書き換えは可能です。しかし速度が出ません。そのためOpenOCDはヘルパープログラムを使った高速な書き込みを実装しています。

SiFive SoC向けのSPI Flashインタフェースのコードを見ると、下記のような謎の #includeが出てきます。参照先のファイルにはバイナリが羅列されています。

ヘルパープログラムのバイナリを保持する配列の定義


// openocd/src/flash/nor/fespi.c

static const uint8_t algorithm_bin[] = {
#include "../../../contrib/loaders/flash/fespi/fespi.inc"
};


// openocd/contrib/loaders/flash/fespi/fespi.inc

/* Autogenerated with ../../../../src/helper/bin2char.sh */
0x6f,0x00,0xc0,0x01,0x73,0x00,0x10,0x00,0x6f,0x00,0xc0,0x02,0x6f,0x00,0x00,0x05,
0x6f,0x00,0xc0,0x05,0x6f,0x00,0x00,0x07,0x6f,0x00,0x00,0x0a,0x83,0xc2,0x05,0x00,

...

このファイルの元となるコード（アセンブラです）は下記のディレクトリにあります。RISC-V向けのクロスコンパイラを用意して、このディレクトリでmakeを実行するとfespi.incを自分で作成することも可能です。

ヘルパープログラムのソースコードの位置

$ ls openocd/contrib/loaders/flash/fespi/

Makefile  fespi.S  fespi.inc

Makefileではfespi.Sをコンパイルしてバイナリをfespi.incに変換します。バイナリからfespi.incへの変換にはsrc/helper/bin2char.shというヘルパースクリプトを使います。

SPI Flashへの書き込み

前置きが長くなりましたが、フラッシュの書き込みルーチンは下記のとおりです。

OpenOCDのSPI Flash書き込みドライバ（SiFive SPI用）


// openocd/src/flash/nor/fespi.c

static int fespi_write(struct flash_bank *bank, const uint8_t *buffer,
		uint32_t offset, uint32_t count)
{

...

	struct working_area *algorithm_wa;
	if (target_alloc_working_area(target, sizeof(algorithm_bin),
				&algorithm_wa) != ERROR_OK) {
		LOG_WARNING("Couldn't allocate %zd-byte working area.",
				sizeof(algorithm_bin));
		algorithm_wa = NULL;
	} else {
		retval = target_write_buffer(target, algorithm_wa->address,
				sizeof(algorithm_bin), algorithm_bin);    //★ヘルパープログラム本体algorithm_binをターゲットのメモリに書く
		if (retval != ERROR_OK) {
			LOG_ERROR("Failed to write code to " TARGET_ADDR_FMT ": %d",
					algorithm_wa->address, retval);
			target_free_working_area(target, algorithm_wa);
			algorithm_wa = NULL;
		}
	}

...

	page_offset = offset % page_size;
	/* central part, aligned words */
	while (count > 0) {
		/* clip block at page boundary */
		if (page_offset + count > page_size)
			cur_count = page_size - page_offset;
		else
			cur_count = count;

		if (algorithm_wa)
			retval = steps_add_buffer_write(as, buffer, offset, cur_count);    //★バッファに書き込むデータを追加する（通常はこちら）
		else
			retval = slow_fespi_write_buffer(bank, buffer, offset, cur_count);    //★JTAGからSPIを直接操作して書き込むモード（遅い）
		if (retval != ERROR_OK)
			goto err;

		page_offset = 0;
		buffer += cur_count;
		offset += cur_count;
		count -= cur_count;
	}

	//★ヘルパープログラムに書き込むデータを全部渡し、フラッシュに書き込みしてもらう
	if (algorithm_wa)
		retval = steps_execute(as, bank, algorithm_wa, data_wa);

...


static int steps_execute(struct algorithm_steps *as,
		struct flash_bank *bank, struct working_area *algorithm_wa,
		struct working_area *data_wa)
{

...

	while (!as_empty(as)) {
		keep_alive();
		uint8_t *data_buf = malloc(data_wa->size);
		unsigned bytes = as_compile(as, data_buf, data_wa->size);
		retval = target_write_buffer(target, data_wa->address, bytes,
				data_buf);
		free(data_buf);
		if (retval != ERROR_OK) {
			LOG_ERROR("Failed to write data to " TARGET_ADDR_FMT ": %d",
					data_wa->address, retval);
			goto exit;
		}

		retval = target_run_algorithm(target, 0, NULL, 2, reg_params,
				algorithm_wa->address, algorithm_wa->address + 4,
				10000, NULL);    //★ヘルパープログラムでflash write

...


// openocd/src/target/target.c

int target_run_algorithm(struct target *target,
		int num_mem_params, struct mem_param *mem_params,
		int num_reg_params, struct reg_param *reg_param,
		uint32_t entry_point, uint32_t exit_point,
		int timeout_ms, void *arch_info)
{
	int retval = ERROR_FAIL;

	if (!target_was_examined(target)) {
		LOG_ERROR("Target not examined yet");
		goto done;
	}
	if (!target->type->run_algorithm) {
		LOG_ERROR("Target type '%s' does not support %s",
				target_type_name(target), __func__);
		goto done;
	}

	target->running_alg = true;
	retval = target->type->run_algorithm(target,
			num_mem_params, mem_params,
			num_reg_params, reg_param,
			entry_point, exit_point, timeout_ms, arch_info);    //★riscv_run_algorithmへ


// openocd/src/target/riscv/riscv.c

/* Algorithm must end with a software breakpoint instruction. */
static int riscv_run_algorithm(struct target *target, int num_mem_params,
		struct mem_param *mem_params, int num_reg_params,
		struct reg_param *reg_params, target_addr_t entry_point,
		target_addr_t exit_point, int timeout_ms, void *arch_info)
{

...

	/* Save registers */
	struct reg *reg_pc = register_get_by_name(target->reg_cache, "pc", true);
	if (!reg_pc || reg_pc->type->get(reg_pc) != ERROR_OK)
		return ERROR_FAIL;
	uint64_t saved_pc = buf_get_u64(reg_pc->value, 0, reg_pc->size);    //★復帰するときに使う
	LOG_DEBUG("saved_pc=0x%" PRIx64, saved_pc);

...

	/* Run algorithm */
	LOG_DEBUG("resume at 0x%" TARGET_PRIxADDR, entry_point);
	if (riscv_resume(target, 0, entry_point, 0, 0, true) != ERROR_OK)    //★ヘルパープログラムを起動する
		return ERROR_FAIL;

...

かなり複雑で全部説明するのは難しいですが、処理の大まかな流れとしては下記のようになっています。

ヘルパープログラムをターゲットのメモリに書く
ヘルパープログラムに渡す制御データを作成する（SPIの制御コマンド＋書き込むデータ列で構成されています）
ヘルパープログラムの後ろに制御データを書く
ヘルパープログラムを実行する

あとはデータが尽きるまで3つ目と4つ目の処理を実行し続けます。関数でいうとsteps_execute() です。1回で書き込む量は「ページ」のサイズに依存します。ページサイズはSPI Flashデバイスによって違いますが、128バイトか256バイトが多いです。

OpenOCDが対応しているSPI Flashは下記のファイルにまとまっていて、

SPI Flashのパラメータを保持している配列の定義


// openocd/src/flash/nor/spi.c

const struct flash_device flash_devices[] = {
	/* name, read_cmd, qread_cmd, pprog_cmd, erase_cmd, chip_erase_cmd, device_id,
	 * pagesize, sectorsize, size_in_bytes */
	FLASH_ID("st m25p05",           0x03, 0x00, 0x02, 0xd8, 0xc7, 0x00102020, 0x80,  0x8000,  0x10000),
	FLASH_ID("st m25p10",           0x03, 0x00, 0x02, 0xd8, 0xc7, 0x00112020, 0x80,  0x8000,  0x20000),
	FLASH_ID("st m25p20",           0x03, 0x00, 0x02, 0xd8, 0xc7, 0x00122020, 0x100, 0x10000, 0x40000),

...

	FLASH_ID("issi is25lp128d",     0x03, 0xeb, 0x02, 0xd8, 0xc7, 0x0018609d, 0x100, 0x10000, 0x1000000),
	FLASH_ID("issi is25wp128d",     0x03, 0xeb, 0x02, 0xd8, 0xc7, 0x0018709d, 0x100, 0x10000, 0x1000000),
	FLASH_ID("issi is25lp256d",     0x13, 0xec, 0x12, 0xdc, 0xc7, 0x0019609d, 0x100, 0x10000, 0x2000000),
	FLASH_ID("issi is25wp256d",     0x13, 0xec, 0x12, 0xdc, 0xc7, 0x0019709d, 0x100, 0x10000, 0x2000000),

...

例えばHiFive Unleashedに搭載されているSPI FlashはISSI is25wp256dですから、ページサイズは0x100 = 256バイトとわかります。

この情報だけでは不安ならデータシートの8.10 PAGE PROGRAM OPERATION (PP, 02h or 4PP, 12h) を見ていただくと確実です。"The Page Program (PP/4PP) instruction allows up to 256 bytes data ..." とあります。

SPI Flashの書き込みにはもっと簡単な（ただし遅い）方法もあります。今回の処理は実用的な反面、書き込みとは直接関係ない機構（ヘルパープログラム）が登場してやや複雑でした。じっくり見るのは面倒だなと思った方は、次回の簡易版を見ていただくほうがわかりやすいかもしれません。

続きはまた今度。

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

コメント一覧

コメントはありません。

この記事にコメントする

2021年7月5日

permalink

編集する

OpenOCDとHiFive UnleashedのSPI Flashその1 - UnleashedとOpenOCD

目次: OpenOCD

OpenOCDではSiFive HiFive1のSPI Flashへの書き込み＆消去ができます。しかしSiFive HiFive UnleashedのSPI Flashに対しては正常に動作しません。書き込みは成功しますが、消去はできません。一見すると成功しているのに一切データが消去できない謎の現象が発生します。

HiFive UnleashedのSPI FlashにOpenOCDで書き込み＆消去を行った例を示します。OpenOCDの制御方法は何通りかあります。直接制御したい場合はポート4444にtelnetすると良いです。

OpenOCDにtelnetで接続


$ telnet localhost 4444

Trying ::1...
Trying 127.0.0.1...
Connected to localhost.
Escape character is '^]'.
Open On-Chip Debugger

> reset halt
JTAG tap: riscv.cpu tap/device found: 0x20000913 (mfg: 0x489 (SiFive Inc), part: 0x0000, ver: 0x2)

私はGDB経由で制御するほうが慣れているので、GDB経由で制御します。GDBの場合はポート3333に接続します。この例ではZephyr用のツールチェーンに含まれるGDBを使っていますが、RV64に対応していれば何でも良いです（RV32専用ではダメです、UnleashedのCPUはRV64なので）。

OpenOCDにGDBで接続

$ riscv64-zephyr-elf-gdb

(gdb) set arch riscv:rv64

The target architecture is assumed to be riscv:rv64

(gdb) target remote :3333

Remote debugging using :3333

(gdb) monitor reset halt

JTAG tap: riscv.cpu tap/device found: 0x20000913 (mfg: 0x489 (SiFive Inc), part: 0x0000, ver: 0x2)

★★monitor xxxxはremote monitor（この場合OpenOCD）へのコマンドと解釈される

(gdb) monitor flash probe 0

Found flash device 'issi is25wp256d' (ID 0x0019709d)
device needs paging or 4-byte addresses - not implemented    ★★ん？なんだこれ
flash 'fespi' found at 0x20000000

(gdb) x/32x 0x20000000

0x20000000:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000010:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000020:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000030:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000040:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000050:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000060:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000070:     0xffffffff      0xffffffff      0xffffffff      0xffffffff

(gdb) monitor flash fillw 0x20000000 0xddccbbaa 0x10

Disabling abstract command writes to CSRs.
wrote 64 bytes to 0x20000000 in 0.203977s (0.306 KiB/s)

★★書き込みに成功


(gdb) x/32x 0x20000000

0x20000000:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000010:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000020:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000030:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000040:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000050:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000060:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000070:     0xffffffff      0xffffffff      0xffffffff      0xffffffff

★★書き込めた


(gdb) monitor flash erase_address 0x20000000 0x10000

erased address 0x20000000 (length 65536) in 0.204166s (313.470 KiB/s)

★★消去も成功したように見えるものの


(gdb) x/32x 0x20000000

0x20000000:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000010:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000020:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000030:     0xddccbbaa      0xddccbbaa      0xddccbbaa      0xddccbbaa
0x20000040:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000050:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000060:     0xffffffff      0xffffffff      0xffffffff      0xffffffff
0x20000070:     0xffffffff      0xffffffff      0xffffffff      0xffffffff

★★消えないぞ……？

書き込めるのに消せないとは？これいかに？この謎を追います。最初にOpenOCDがどうやってSPI Flashを書き換えているのか？次にどうやって消去しているのか？を追います。妙な警告メッセージが出ている点も気になります。

SPI Flashとは？

その名の通りSPIで接続されたFlashメモリデバイスのことです。SPIはSerial Peripheral Interfaceの略で、CE#（Chip EnableもしくはChip Select, CS# とも）、SCLK、SDI、SDOの4線で構成される非常にシンプルなI/Oバスです。

個人的には「入力と出力が常に同時に行われるのがSPIの大きな特徴」に思います。

SPIの入出力波形の例（SPI FlashのRead Product Identificationコマンド）

例えばこの波形図ではSoCからSPI FlashにInstruction(0xAB) と3バイトのダミーデータを出力（SI側の信号）しているとき、SPI FlashからSoCに意味のないデータが返ってきます（SO側の信号）。

入出力が別々に行われるバス（Ethernetなど）に慣れていると、無意味なデータを受け取る＝無駄なリクエストを相手に送っている、ことを示しますから、無駄な入力はやめてくれ〜としばらく混乱しました。しかしまあ、わかってしまえば非常に単純な話でして、常に入出力が同時に行われる仕様なので、無駄は発生していません。また、無意味な入力データは受け取った後に無視すれば良いだけです。

続きはまた次回。

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

コメント一覧

コメントはありません。

この記事にコメントする

2021年7月3日

permalink

編集する

OpenCLのOSS実装poclを調べるその8 - poclとClang/LLVMの引数

目次: OpenCL

前回に続いてpoclとClang/LLVMの関係について説明します。Clangは引数の体系が2つあります。1つは通常使用するGCCと互換があるオプション、もう1つはclang -cc1オプションを指定したときに使う内部用オプションです。正式な名前がわからないので適当に呼んでいます。正式な名前をご存知の方は教えていただけると嬉しいです。

前者と後者は -Iや -oのように同じ名前の場合もありますし、顕著に異なる場合もあります。RISC-V向けのmarchやmabiは差が顕著なので、例として紹介します。

GCC互換オプション

$ riscv32-unknown-elf-gcc b.c -c -march=rv32gc -mabi=ilp32f

$ clang --target=riscv32 b.c -c -march=rv32gc -mabi=ilp32f

GCCとClangのオプションがほぼ同じですね。次にClang内部用のオプションを確認します。内部用のオプションは -vを指定すると表示できます。

Clang内部用オプションを表示

$ clang --target=riscv32 b.c -c -march=rv32gc -mabi=ilp32f -v

Debian clang version 11.0.1-2
Target: riscv32--
Thread model: posix
InstalledDir: /usr/bin
 (in-process)
 "/usr/lib/llvm-11/bin/clang" -cc1 -triple riscv32-- -emit-obj -mrelax-all -disable-free -disable-llvm-verifier -discard-value-names -main-file-name b.c -mrelocation-model static -mframe-pointer=all -fmath-errno -fno-rounding-math -mconstructor-aliases -nostdsysteminc -target-feature +m -target-feature +a -target-feature +f -target-feature +d -target-feature +c -target-feature +relax -target-feature -save-restore -target-abi ilp32f -msmall-data-limit 8 -fno-split-dwarf-inlining -debugger-tuning=gdb -v -resource-dir /usr/lib/llvm-11/lib/clang/11.0.1 -internal-isystem include -fdebug-compilation-dir /home/katsuhiro/share/projects/c/clang_test -ferror-limit 19 -fno-signed-char -fgnuc-version=4.2.1 -fcolor-diagnostics -faddrsig -o b.o -x c b.c
clang -cc1 version 11.0.1 based upon LLVM 11.0.1 default target x86_64-pc-linux-gnu
ignoring nonexistent directory "include"
#include "..." search starts here:
#include <...> search starts here:
 /usr/lib/llvm-11/lib/clang/11.0.1/include
End of search list.

先ほどと全く違うことがわかると思います。このままだと色々ごちゃごちゃ表示されていてわかりにくいので、対応するオプションだけ抜粋します。

Clang内部用オプション（対応部分のみ抜粋）

--target=riscv32: -triple riscv32--
-march=rv32gc   : -target-feature +m -target-feature +a -target-feature +f -target-feature +d -target-feature +c
-mabi=ilp32f    : -target-abi ilp32f

なお -target-featureはカンマ区切りで複数指定することもできるようです。上記で言えば -target-feature +m,+a,+f,+d,+cにしても良いです。

poclと何の関係が？

なぜこの話をしたのかというとpocl内でLLVMを呼び出す際は、Clang内部用オプションを使わなければならない箇所があるからです。一番わかりやすい（文字列の形でオプションを指定している）のはpocl_llvm_build_program() です。

poclがClang内部用オプションを使っている箇所（文字列で設定するタイプ）


// pocl/lib/CL/pocl_llvm_build.cc

int pocl_llvm_build_program(cl_program program,
                            unsigned device_i,
                            const char *user_options_cstr,
                            char *program_bc_path,
                            cl_uint num_input_headers,
                            const cl_program *input_headers,
                            const char **header_include_names,
                            int linking_program)

{

...

  CompilerInvocation &pocl_build = CI.getInvocation();

...

  // This is required otherwise the initialization fails with
  // unknown triple ''
  ss << "-triple=" << device->llvm_target_triplet << " ";    //★これ
  if (device->llvm_cpu != NULL)
    ss << "-target-cpu " << device->llvm_cpu << " ";    //★これも同様

...

  std::istream_iterator<std::string> begin(ss);
  std::istream_iterator<std::string> end;
  std::istream_iterator<std::string> i = begin;
  std::vector<const char*> itemcstrs;
  std::vector<std::string> itemstrs;
  while (i != end) {
    itemstrs.push_back(*i);    //★std::vectorにオプションのstd::stringを並べる
    ++i;
  }

  for (unsigned idx = 0; idx < itemstrs.size(); idx++) {
    // note: if itemstrs is modified after this, itemcstrs will be full
    // of invalid pointers! Could make copies, but would have to clean up then...
    itemcstrs.push_back(itemstrs[idx].c_str());    //★std::vectorにオプションの文字列のポインタを並べる
  }

...

  //★コンパイラに上記のオプションを指定する（この時点ではまだコンパイラは起動しない）
  if (!CompilerInvocation::CreateFromArgs(
          pocl_build,    //★CompilerInvocation
#ifndef LLVM_OLDER_THAN_10_0
          ArrayRef<const char *>(itemcstrs.data(),
                                 itemcstrs.data() + itemcstrs.size()),
#else
          itemcstrs.data(), itemcstrs.data() + itemcstrs.size(),
#endif
          diags)) {

あとはpocl_llvm_codegen() も同様ですが、指定の方法が違います。-target-feature +mのような文字列ではなく、TargetMachineのメソッドを呼び出して指定します。コードで見たほうがわかりやすいでしょう。

poclがClang内部用オプションを使っている箇所（関数で設定するタイプ）


// pocl/lib/CL/pocl_llvm_wg.cc

int pocl_llvm_codegen(cl_device_id Device, void *Modp, char **Output,
                      uint64_t *OutputSize) {

...

  llvm::Triple Triple(Device->llvm_target_triplet);
  llvm::TargetMachine *Target = GetTargetMachine(Device, Triple);

  // First try direct object code generation from LLVM, if supported by the
  // LLVM backend for the target.
  bool LLVMGeneratesObjectFiles = true;

  SmallVector<char, 4096> Data;
  llvm::raw_svector_ostream SOS(Data);
  bool cannotEmitFile;

  cannotEmitFile = Target->addPassesToEmitFile(PMObj, SOS,
#ifndef LLVM_OLDER_THAN_7_0
                                  nullptr,
#endif
                                  CODEGEN_FILE_TYPE_NS::CGFT_ObjectFile);

  Target->setTargetFeatureString("+m,+a,+f");    //★-target-feature +m,+a,+fに相当する

独自アクセラレータ向けの実装を行う際にClang/LLVMのオプションを変えたくなることは多々ありますから、今回のオプションの違いは今後の説明でも登場するはずです。たぶん。

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

コメント一覧

コメントはありません。

この記事にコメントする

2021年7月2日

permalink

編集する

OpenCLのOSS実装poclを調べるその7 - poclとLLVMとOpenCLカーネルビルド

目次: OpenCL

今回は少し話題を変えてpoclとClang/LLVMの関係について説明します。poclはOpenCL C言語をビルドするため内部的にClang/LLVMを呼び出します。ビルドはおおまかに2段階に分かれています。設定次第で多少変わりますが、ざっくり各段階で何をしているのか説明します。まだ話題にしていない関数も出てきますが、その話は追々やっていこうと思います。

第1段階OpenCL C -> LLVM IR

最初はclBuildProgram() です。やっていることはOpenCL C -> LLVM IR Bitcode（コンパイル＋OpenCLランタイムとのリンク）です。処理に関連する関数と呼び出し関係は下記のとおりです。

clBuildProgram() とLLVM IR Bitcode生成関数の呼び出し関係

clBuildProgram()
  compile_and_link_program()
    pocl_llvm_build_program(): プリプロセス、コンパイル（出力LLVM IR Bitcode）
      link(): OpenCLランタイムとリンク（未定義シンボルの検出）
      pocl_write_module(): 出力LLVM IR Bitcode

この関数の実行中に3つ一時ファイルが出力されます。本来は消されてしまって残らないファイルもありますが、コードを変更し一時ファイルをあえて残すと、

tempfile-xx-xx-xx-xx-xx.cl: オリジナルのOpenCL Cコード（xx-xx-... の部分はハッシュ値）
tempfile-xx-xx-xx-xx-xx.preproc.cl: プリプロセス後のOpenCL Cコード（xx-xx-... の部分はハッシュ値）
~/.cache/pocl/kcache/XX/XXXXXXXX/program.bc: リンク後のLLVM IR Bitcode（XXはハッシュ値なので名前は様々）

が作成されます。link() は全てのシンボルを問答無用で追加するのではなくて、OpenCLカーネルコードから参照されているシンボル（＝未定義のシンボル）のみを追加します。

第2段階LLVM IR -> バイナリ

次はclEnqueueNDRangeKernel() です。やることはLLVM IR Bitcode -> ターゲットバイナリです。

API名はコンパイルやビルドと関係なさそうに見えますが、poclのホストCPU向け実装を見る限り、NDRangeKernelから起因してLLVM IR Bitcodeからターゲットデバイスのバイナリに変換しています（あと何度も変換しなくて良いように生成したバイナリをキャッシュする）。

処理に関連する関数と呼び出し関係は下記のとおりです。

ホストCPU向け実装の呼び出し関係

clEnqueueNDRangeKernel()
  pocl_command_enqueue()
    pocl_pthread_submit()
      pthread_scheduler_push_command(): キューにコマンドを追加する、ワーカースレッドが起床してコマンドを得る

start_thread(): ワーカースレッド
  pocl_pthread_driver_thread()
    pthread_scheduler_get_work()
    check_cmd_queue_for_device(): キューからコマンドを得る、コマンドがNDRangeKernelの実行要求だったら、下記を呼ぶ
      pocl_pthread_prepare_kernel()
        pocl_check_kernel_dlhandle_cache(): dlopen()
          pocl_check_kernel_disk_cache(): キャッシュ
            llvm_codegen(): ターゲットデバイスバイナリ生成
              pocl_llvm_codegen(): .so.oバイナリ作成
              pocl_invoke_clang(): .soバイナリ作成

独自アクセラレータ向け実装もホストCPU向け実装に習いclEnqueueNDRangeKernel() を起点にバイナリに変換を行います。現状はキャッシュ機構はスキップし、毎回ターゲットバイナリを生成する実装です。このままだと非常に遅いので今後、改善する必要があります。

独自アクセラレータ向け実装（仮）の呼び出し関係

clEnqueueNDRangeKernel()
  pocl_command_enqueue()
    pocl_accel_submit(): ここはターゲットデバイスごとに実装が異なる
      scheduleCommands()
        pocl_exec_command()
          pocl_accel_run(): ここから先はターゲットデバイスごとに実装が異なる
            llvm_codegen(): ターゲットデバイスバイナリ生成
              pocl_llvm_codegen(): .so.oバイナリ作成
              pocl_invoke_clang(): .soバイナリ作成

この関数の実行中に2つ一時ファイルが出力されます。本来は消されてしまって残らないファイルもありますが、コードを変更し一時ファイルをあえて残すと、

tempfile-xx-xx-xx-xx-xx.so.o: 中間バイナリ
XX/XXXXXXXX/kernelName/xx-x-x/kernelName.so: 最終的なOpenCLカーネルバイナリ

が作成されます。中間バイナリと最終バイナリの差は、C言語のコンパイルのときに生成する *.oと実行ファイルとの違いと同様に、前者はアドレスが解決されておらず、後者はアドレス解決済みという差がありました。他にも何か違いがあるかもしれません。

第1段階で紹介した通りclBuildProgram() のlink() にてOpenCLカーネルに必要なシンボルが集められます。そのため中間バイナリと最終バイナリを比較しても、含まれる関数やシンボルはほぼ変わりません（例外はclang_rt.builtinsに依存した関数、例えば __adddf3() など）。

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

コメント一覧

コメントはありません。

この記事にコメントする

2021年7月1日

permalink

編集する

OpenCLのOSS実装poclを調べるその6 - OpenCLカーネルのビルド

目次: OpenCL

引き続き、独自アクセラレータのテンプレート実装pocl/lib/CL/devices/accelの細かな問題を調べます。初期化を突破するとカーネルのビルドclBuildProgram() でコケます。

clBuildProgram() でエラー

POCL: in fn compile_and_link_program at line 664:
  |     ERROR | CL_COMPILER_NOT_AVAILABLE Cannot build a program from sources with pocl that does not have online compiler support

原因はdev->compiler_availableがCL_TRUEになっていないことです。初期化が足りないようです。

compiler_availableがfalseのときの動作


// pocl/lib/CL/pocl_build.c

cl_int
compile_and_link_program(int compile_program,
                         int link_program,
                         cl_program program,
                         cl_uint num_devices,
                         const cl_device_id *device_list,
                         const char *options,
                         cl_uint num_input_headers,
                         const cl_program *input_headers,
                         const char **header_include_names,
                         cl_uint num_input_programs,
                         const cl_program *input_programs,
                         void (CL_CALLBACK *pfn_notify) (cl_program program,
                                                         void *user_data),
                         void *user_data)
{

...

      /* clCreateProgramWithBuiltinKernels */
      /* No build step supported at the moment for built-in kernels. */
      if (program->builtin_kernel_names)
        continue;

      /* clCreateProgramWithSource */
      else if (program->source)
        {
#ifdef OCS_AVAILABLE
          if (device->compiler_available == CL_TRUE)    //★このif文が成立せず
            {
              POCL_MSG_PRINT_INFO ("building from sources for device %d\n",
                                   device_i);
              error = pocl_llvm_build_program (
                  program, device_i, program->compiler_options,
                  program_bc_path, num_input_headers, input_headers,
                  header_include_names, (create_library ? 0 : link_program));
              POCL_GOTO_ERROR_ON ((error != 0), build_error_code,
                                  "pocl_llvm_build_program() failed\n");
            }
          else
#endif
            {    //★こちらにきてエラーになってしまう
              APPEND_TO_MAIN_BUILD_LOG (
                  "Cannot build a program from sources with pocl "
                  "that does not have online compiler support\n");
              POCL_GOTO_ERROR_ON (1, CL_COMPILER_NOT_AVAILABLE, "%s",
                                  program->main_build_log);
            }
        }

他の実装を眺めるとpocl_init_default_device_infos() を呼んで解決しているようなので、先達に習いpocl_accel_init() でpocl_init_default_device_infos() を呼び出すように書き換えましょう。

その他のビルド関連の初期化

もともとあったdev->version, dev->available, dev->profileの初期化はpocl_init_default_device_infos() が行いますから、削除しても良いかもしれません。残っていても特に害はないと思いますけど。

現状では拡張機能dev->extensions = "cl_khr_fp64" を設定しないとdouble型を使ったときにビルドエラーになります。最終的に必要な拡張機能がはっきりするまでは必要になったものを順次追加する形で、次に進みましょう。

拡張機能はただ書けば動くわけではなくOpenCL Cコンパイラが対応している必要があります。当たり前ですね。poclが使っているコンパイラは大御所LLVMですから、RISC-V向けだけ機能が欠けていることはないでしょう。たぶん。

ビルド対象はRISC-Vにしたいので、

llvm_traget_tripletとllvm_cpuの設定値


dev->llvm_target_triplet = "riscv32";
dev->llvm_cpu = "generic-rv32";

としました。llvm_target_tripletはclang -cc1の -tripleオプションに渡されます。またllvm_cpuは -target-cpuオプションに渡されます。有効な値を調べる方法は、

llvm_target_tripletの調べ方

$ clang -print-targets

  Registered Targets:
    aarch64    - AArch64 (little endian)
    aarch64_32 - AArch64 (little endian ILP32)
    aarch64_be - AArch64 (big endian)
    amdgcn     - AMD GCN GPUs
...
    riscv32    - 32-bit RISC-V
    riscv64    - 64-bit RISC-V
    sparc      - Sparc
    sparcel    - Sparc LE
    sparcv9    - Sparc V9
    systemz    - SystemZ
...

llvm_cpuの調べ方

$ clang -target=riscv32 -print-supported-cpus

Debian clang version 11.0.1-2
Target: riscv32
Thread model: posix
InstalledDir: /usr/bin
Available CPUs for this target:

        generic-rv32
        generic-rv64
        rocket-rv32
        rocket-rv64
        sifive-e31
        sifive-u54

Use -mcpu or -mtune to specify the target's processor.
For example, clang --target=aarch64-unknown-linux-gui -mcpu=cortex-a35

GCCはコマンド名がアーキテクチャ名そのものですし、LLVMはヘルプで対応アーキテクチャがわかります。どちらも親切で良いですね。

仮実装

ここまでの変更を反映すると、

初期化関数（仮）


// pocl/lib/CL/device/accel/accel.cc

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

  AccelData *D = new AccelData;
  dev->data = (void *)D;

  pocl_init_default_device_infos (dev);

  //SETUP_DEVICE_CL_VERSION(1, 2);    //★pocl_init_default_device_infosが初期化するのでいらない
  dev->type = CL_DEVICE_TYPE_CUSTOM;
  dev->long_name = (char *)"memory mapped custom device";
  dev->vendor = "pocl";
  //dev->version = "1.2";        //★pocl_init_default_device_infosが初期化するのでいらない
  //dev->available = CL_TRUE;    //★pocl_init_default_device_infosが初期化するのでいらない

  dev->extensions = "cl_khr_fp64";    //★cl_khr_fp64を追加
  //dev->profile = "FULL_PROFILE";     //★pocl_init_default_device_infosが初期化するのでいらない
  dev->max_mem_alloc_size = 100 * 1024 * 1024;

  dev->llvm_target_triplet = "riscv32";    //★追加
  dev->llvm_cpu = "generic-rv32";          //★追加

  dev->final_linkage_flags = final_ld_flags;

  if (!parameters) {
    POCL_ABORT("accel: parameters were not given\n");
  }

まだまだ変更が必要ですが、こだわるのは後にして次に進みます。

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

コメント一覧

コメントはありません。

この記事にコメントする

2021年6月24日

permalink

編集する

プログラムからLLVMを実行するその2 - プリプロセス編

目次: LLVM

準備が終わりましたらClang/LLVMをプログラムから呼びましょう。

LLVMでプリプロセスだけを実行するプログラム


int main(int argc, char *argv[])
{
	bool success;

	clang::CompilerInstance CI;
	clang::CompilerInvocation &build = CI.getInvocation();

	// 引数の配列を作成する
	std::vector<const char*> vec_args;
	vec_args.push_back("-I/usr/include/c++/10");
	vec_args.push_back("-I/usr/include/x86_64-linux-gnu/c++/10");
	vec_args.push_back("-I/usr/include/c++/10/backward");
	vec_args.push_back("-I/usr/lib/llvm-11/lib/clang/11.0.1/include");
	vec_args.push_back("-I/usr/include/x86_64-linux-gnu");
	vec_args.push_back("-I/usr/include");
	vec_args.push_back("-I/path/to/llvm-project/_install/include");

	// エラーメッセージを出力するために使われるクラス
	llvm::IntrusiveRefCntPtr<clang::DiagnosticIDs> diagID = new clang::DiagnosticIDs();
	llvm::IntrusiveRefCntPtr<clang::DiagnosticOptions> diagOpts = new clang::DiagnosticOptions();
	clang::TextDiagnosticBuffer *diagBuffer = new clang::TextDiagnosticBuffer();
	clang::DiagnosticsEngine diags(diagID, diagOpts, diagBuffer);
	CI.createDiagnostics(diagBuffer, false);

	// コンパイラ呼び出し用のインスタンスを作成する
	llvm::ArrayRef<const char*> ref_args(vec_args.data(), vec_args.data() + vec_args.size());
	success = clang::CompilerInvocation::CreateFromArgs(build, ref_args, diags);

	// コンパイラフロントエンドのオプション設定
	//   入力ソースコード: test.cpp
	//   出力ソースコード: test.preproc.cpp
	const char *source_file = "test.cpp";
	const char *preproc_file = "test.preproc.cpp";
	clang::FrontendOptions &fe   = build.getFrontendOpts();
	clang::InputKind ik          = clang::InputKind(clang::Language::CXX);
	clang::FrontendInputFile fif = clang::FrontendInputFile(source_file, ik);

	fe.Inputs.clear();
	fe.Inputs.push_back(fif);
	fe.OutputFile.assign(preproc_file);

	// プリプロセスのオプション設定
	//   言語: C++11
	clang::PreprocessorOptions &po = build.getPreprocessorOpts();
	clang::LangOptions *la         = build.getLangOpts();
	llvm::Triple triple            = llvm::Triple();
	build.setLangDefaults(*la, ik, triple, po.Includes, clang::LangStandard::lang_cxx11);

	// 下記のようにオプションの一部だけ変えることもできる
	//la->CPlusPlus = true;
	//la->CPlusPlus11 = true;

	// プリプロセスのオプション
	//   コメント、定義済みマクロなどは出力しない
	clang::PreprocessorOutputOptions &poo = build.getPreprocessorOutputOpts();
	poo.ShowCPP = true;
	poo.ShowComments = false;
	poo.ShowLineMarkers = false;
	poo.ShowMacros = false;
	poo.ShowMacroComments = false;
	poo.RewriteIncludes = false;

	// プリプロセス実行（失敗したらエラーログを出力する）
	clang::PrintPreprocessedAction Preprocess;
	success = CI.ExecuteAction(Preprocess);
	if (!success) {
		get_build_log(diagBuffer, (CI.hasSourceManager()) ? &CI.getSourceManager() : nullptr);
	}
}

残念ながらこの呼び出し方が正解とは断言できません。探した限りではどう呼び出すべきか書かれたドキュメントも見当たりませんでした。上記の例はpoclを参考にしており、大きな間違いはないはずですが……。何かやらかしていたら教えていただけると嬉しいです。

動作確認はLLVM 12で行いました。他のバージョンだとAPIの引数などが変わっているので、ビルドすら通らないと思います。LLVMの困ったところですね……。

インクルードパスの調べ方

上記のサンプルでは引数で -Iオプションを使ってインクルードパスを指定します。インクルードパスは頑張ってヘッダファイルがある場所を調べても良いですが、おそらく同じ名前のヘッダが複数の場所にあって混乱すると思いますから、PCで動作しているClang++ から拝借するのが簡単です。

clangのインクルードパスを調べる

$ clang++ test.cpp -v

Debian clang version 11.0.1-2
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin
Found candidate GCC installation: /usr/bin/../lib/gcc/x86_64-linux-gnu/10

...

#include "..." search starts here:
#include <...> search starts here:
 /usr/bin/../lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10
 /usr/bin/../lib/gcc/x86_64-linux-gnu/10/../../../../include/x86_64-linux-gnu/c++/10
 /usr/bin/../lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/backward
 /usr/lib/llvm-11/lib/clang/11.0.1/include
 /usr/include/x86_64-linux-gnu
 /usr/include
End of search list.

...

いろいろなメッセージが出力されますが、インクルードパスは "search starts here:" の辺りに書かれています。出力は特に捻りはなくディレクトリ名そのものですので、頭に -Iを足せばオプションの出来上がりです。

実行

プリプロセスを実行します。テスト用のプログラムは下記のとおりです。

テスト用のプログラム


#include <iostream>

int main(int argc, char *argv[])
{
	// This is comment
	std::cout << "Hello, world!!" << std::endl;
}

プリプロセス実行

$ make

$ ./clang_test

ファイル名などは完全に決め打ちのため引数は必要ありません。実行に成功するとプリプロセス後のソースコードtest.preproc.cppが作成されているはずです。

プリプロセス後のソースコード


namespace std
{
  typedef long unsigned int size_t;
  typedef long int ptrdiff_t;


  typedef decltype(nullptr) nullptr_t;

}

...

  static ios_base::Init __ioinit;


}

int main(int argc, char *argv[])
{

 std::cout << "Hello, world!!" << std::endl;
}

私の環境で実行したところ27,000行くらいあるファイルになりました。たった1つしかヘッダをincludeしてないのに凄まじい行数に展開されます。コメントは消えていますが、オプションを変更すれば残すこともできます。PreprocessorOutputOptionsのShowComments = trueにすると残ります。

プリプロセス後のソースコードをビルド＆実行

$ g++ test.preproc.cpp

$ ./a.out
Hello, world!!

プリプロセス後のソースコードをg++ などに渡すとコンパイル可能なので、おそらく変な出力にはなっていないでしょう。

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

コメント一覧

コメントはありません。

この記事にコメントする

2021年6月23日

permalink

編集する

プログラムからLLVMを実行するその1 - 準備編

目次: LLVM

LLVMやClangは実行する方法が2つあります。1つ目はみなさまお馴染みのコマンドラインから実行する方法で、2つ目はプログラムからClangのライブラリを通して実行する方法です。

特に後者のプログラムから実行する方法はGCCでは真似できませんから、LLVMならではの機能と言えるでしょう。ただ、ちょっとインタフェースが不安定というか、バージョンによってちょいちょい変わって動かなくなるようで、そこは玉に瑕ですね。

LLVMビルド＆インストール

Clang/LLVMをプログラムから実行するにはいくつか準備が必要です。大まかに分けるとLLVMのビルド＆インストールと、ヘッダおよびライブラリパスの指定です。

ビルドは以前もチャレンジしました（2019年3月26日の日記参照）。基本的にはcmakeとmake（またはninja）です。それは変わりませんが、いくつか追加したいオプションがあるので再掲します。

LLVMのビルドオプション

$ cmake \
  -G Ninja \
  ../llvm \
  -DCMAKE_INSTALL_PREFIX=`pwd`/../_install \
  -DCMAKE_C_COMPILER=clang \
  -DCMAKE_CXX_COMPILER=clang++ \
  -DCMAKE_BUILD_TYPE=RelWithDebInfo \
  -DBUILD_SHARED_LIBS=ON \
  -DLLVM_ENABLE_ASSERTIONS=ON \
  -DLLVM_TARGETS_TO_BUILD="X86;RISCV;NVPTX" \
  -DLLVM_USE_LINKER=lld \
  -DLLVM_BUILD_LLVM_DYLIB=OFF \
  -DLLVM_LINK_LLVM_DYLIB=OFF \
  -DLLVM_ENABLE_PROJECTS="clang;clang-tools-extra;compiler-rt;debuginfo-tests;libc;libclc;libcxx;libcxxabi;libunwind;lld;lldb"

ざっくり意図を説明すると下記のとおりです。オプションの正確な意味についてはLLVM公式ドキュメント（Build LLVM with CMake - LLVM 12 documentation 参照）を見てください。

CMAKE_INSTALL_PREFIX: インストール先を指定します。システムに既にインストールされているLLVMを破壊しないよう、ビルドディレクトリの隣の _installディレクトリにインストールする指定です。
LLVM_TARGETS_TO_BUILD: 以前（2019年3月27日の日記参照）も使いましたが、特定ターゲットのみをビルドするオプションで、ビルド時間の短縮に繋がります。2つ以上指定する場合はセミコロンで繋ぎましょう。例ではx86とRISC-V向けにしていますが、お好きなアーキテクチャを足してください。
LLVM_BUILD_LLVM_DYLIB: 全てのライブラリを1つのライブラリlibLLVM.soに集約するオプションです。興味があればこのオプションのON/OFFにより後述するllvm-configの出力がどう変化するか確認すると面白いかもしれません。
LLVM_ENABLE_PROJECTS: LLVMはLLVM以外にも多彩なツールを持っています。どのツールをビルドするか選択するオプションです。全ては必要ないですが少なくともclangは後で必要になります。例では全部入りにしています。

CMakeの実行が成功したら、ninja installを呼びましょう。インストールまで進むはずです。

Makefileの作成

ヘッダインクルードパスの指定、ライブラリパスの指定のためにMakefileを書きます。パスの細かい値について心配する必要はありません。llvm-configというツールが用意されており、ほぼ全て自動的に用意してくれます。Makefileの一例を示すと、

テスト用のMakefile


LLVM_CONFIG_PATH  = /path/to/llvm-project/_install/bin
LLVM_CONFIG       = $(LLVM_CONFIG_PATH)/llvm-config --link-shared

CPPFLAGS = $(shell $(LLVM_CONFIG) --cppflags)
CFLAGS   = $(shell $(LLVM_CONFIG) --cflags) -g
CXXFLAGS = $(shell $(LLVM_CONFIG) --cxxflags) -g
LDFLAGS  = $(shell $(LLVM_CONFIG) --ldflags) 
LIBS     = -lclang-cpp $(LLVM_CONFIG) --libs --system-libs engine)

clang_test: main.o
	$(CXX) $(CXXFLAGS) $(LDFLAGS) -o $(APP) $< $(LIBS)

基本的にはllvm-config --xxxflagsとするとオプションに指定すべき文字列が出力されますから、素直に各種FLAGSに渡すだけです。もちろん何かオプションを追加するのも自由です。例では -gを足しています。

LIBSのところがちょっと格好悪いのは、llvm-configでlibclang-cppにリンクするような方法が見当たらなかったからです。良い方法をご存知の方は教えていただけると嬉しいです。

これで準備完了です。続きは次回に。

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