引き続き、独自アクセラレータのテンプレート実装 pocl/lib/CL/devices/accel の細かな問題を調べます。初期化を突破するとカーネルのビルド 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 になっていないことです。初期化が足りないようです。
// 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 にしたいので、
dev->llvm_target_triplet = "riscv32";
dev->llvm_cpu = "generic-rv32";
としました。llvm_target_triplet は clang -cc1 の -triple オプションに渡されます。また llvm_cpu は -target-cpu オプションに渡されます。有効な値を調べる方法は、
$ 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 ...
$ 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");
}
まだまだ変更が必要ですが、こだわるのは後にして次に進みます。
今回は少し話題を変えて pocl と Clang/LLVM の関係について説明します。pocl は OpenCL C 言語をビルドするため内部的に Clang/LLVM を呼び出します。ビルドはおおまかに 2段階に分かれています。設定次第で多少変わりますが、ざっくり各段階で何をしているのか説明します。まだ話題にしていない関数も出てきますが、その話は追々やっていこうと思います。
最初は clBuildProgram() です。やっていることは OpenCL C -> LLVM IR Bitcode(コンパイル+OpenCL ランタイムとのリンク)です。処理に関連する関数と呼び出し関係は下記のとおりです。
clBuildProgram() compile_and_link_program() pocl_llvm_build_program(): プリプロセス、コンパイル(出力 LLVM IR Bitcode) link(): OpenCL ランタイムとリンク(未定義シンボルの検出) pocl_write_module(): 出力 LLVM IR Bitcode
この関数の実行中に 3つ一時ファイルが出力されます。本来は消されてしまって残らないファイルもありますが、コードを変更し一時ファイルをあえて残すと、
が作成されます。link() は全てのシンボルを問答無用で追加するのではなくて、OpenCL カーネルコードから参照されているシンボル(=未定義のシンボル)のみを追加します。
次は clEnqueueNDRangeKernel() です。やることは LLVM IR Bitcode -> ターゲットバイナリです。
API 名はコンパイルやビルドと関係なさそうに見えますが、pocl のホスト CPU 向け実装を見る限り、NDRangeKernel から起因して LLVM IR Bitcode からターゲットデバイスのバイナリに変換しています(あと何度も変換しなくて良いように生成したバイナリをキャッシュする)。
処理に関連する関数と呼び出し関係は下記のとおりです。
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つ一時ファイルが出力されます。本来は消されてしまって残らないファイルもありますが、コードを変更し一時ファイルをあえて残すと、
が作成されます。中間バイナリと最終バイナリの差は、C 言語のコンパイルのときに生成する *.o と実行ファイルとの違いと同様に、前者はアドレスが解決されておらず、後者はアドレス解決済みという差がありました。他にも何か違いがあるかもしれません。
第 1段階で紹介した通り clBuildProgram() の link() にて OpenCL カーネルに必要なシンボルが集められます。そのため中間バイナリと最終バイナリを比較しても、含まれる関数やシンボルはほぼ変わりません(例外は clang_rt.builtins に依存した関数、例えば __adddf3() など)。
前回に続いて pocl と Clang/LLVM の関係について説明します。Clang は引数の体系が 2つあります。1つは通常使用する GCC と互換があるオプション、もう 1つは clang -cc1 オプションを指定したときに使う内部用オプションです。正式な名前がわからないので適当に呼んでいます。正式な名前をご存知の方は教えていただけると嬉しいです。
前者と後者は -I や -o のように同じ名前の場合もありますし、顕著に異なる場合もあります。RISC-V 向けの march や mabi は差が顕著なので、例として紹介します。
$ 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 --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.
先ほどと全く違うことがわかると思います。このままだと色々ごちゃごちゃ表示されていてわかりにくいので、対応するオプションだけ抜粋します。
--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 内で LLVM を呼び出す際は、Clang 内部用オプションを使わなければならない箇所があるからです。一番わかりやすい(文字列の形でオプションを指定している)のは pocl_llvm_build_program() です。
// 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/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 のオプションを変えたくなることは多々ありますから、今回のオプションの違いは今後の説明でも登場するはずです。たぶん。
OpenOCD では SiFive HiFive1 の SPI Flash への書き込み&消去ができます。しかし SiFive HiFive Unleashed の SPI Flash に対しては正常に動作しません。書き込みは成功しますが、消去はできません。一見すると成功しているのに一切データが消去できない謎の現象が発生します。
HiFive Unleashed の SPI Flash に OpenOCD で書き込み&消去を行った例を示します。OpenOCD の制御方法は何通りかあります。直接制御したい場合はポート 4444 に 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 なので)。
$ 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 は 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 など)に慣れていると、無意味なデータを受け取る=無駄なリクエストを相手に送っている、ことを示しますから、無駄な入力はやめてくれ〜としばらく混乱しました。しかしまあ、わかってしまえば非常に単純な話でして、常に入出力が同時に行われる仕様なので、無駄は発生していません。また、無意味な入力データは受け取った後に無視すれば良いだけです。
続きはまた次回。
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 というヘルパースクリプトを使います。
前置きが長くなりましたが、フラッシュの書き込みルーチンは下記のとおりです。
// 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;
...
かなり複雑で全部説明するのは難しいですが、処理の大まかな流れとしては下記のようになっています。
あとはデータが尽きるまで 3つ目と 4つ目の処理を実行し続けます。関数でいうと steps_execute() です。1回で書き込む量は「ページ」のサイズに依存します。ページサイズは SPI Flash デバイスによって違いますが、128 バイトか 256 バイトが多いです。
OpenOCD が対応している 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 の書き込みにはもっと簡単な(ただし遅い)方法もあります。今回の処理は実用的な反面、書き込みとは直接関係ない機構(ヘルパープログラム)が登場してやや複雑でした。じっくり見るのは面倒だなと思った方は、次回の簡易版を見ていただくほうがわかりやすいかもしれません。
続きはまた今度。
フラッシュの書き込みは、ヘルパープログラムを使うタイプ(上記で追いかけた方)と、SPI を直接制御する方があります。高速に書き込めるのはヘルパープログラムを使うタイプですが、わかりやすいのは 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(bank, 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);
...
ここまでは前回と同じです。前回は steps_add_buffer_write() と steps_execute() が活躍しましたが、今回は slow_fespi_write_buffer() の方を見ます。
static int slow_fespi_write_buffer(struct flash_bank *bank,
const uint8_t *buffer, uint32_t offset, uint32_t len)
{
uint32_t ii;
if (offset & 0xFF000000) {
LOG_ERROR("FESPI interface does not support greater than 3B addressing, can't write to offset 0x%" PRIx32,
offset);
return ERROR_FAIL;
}
/* TODO!!! assert that len < page size */
fespi_tx(bank, SPIFLASH_WRITE_ENABLE);
fespi_txwm_wait(bank);
if (fespi_write_reg(bank, FESPI_REG_CSMODE, FESPI_CSMODE_HOLD) != ERROR_OK) //★CE# 信号を Low にする(負論理、処理の開始を示す)
return ERROR_FAIL;
fespi_tx(bank, SPIFLASH_PAGE_PROGRAM); //★PP コマンド = 0x02
fespi_tx(bank, offset >> 16); //★アドレス
fespi_tx(bank, offset >> 8);
fespi_tx(bank, offset);
for (ii = 0; ii < len; ii++) //★データ
fespi_tx(bank, buffer[ii]);
fespi_txwm_wait(bank); //★送信し終わるまで待つ
if (fespi_write_reg(bank, FESPI_REG_CSMODE, FESPI_CSMODE_AUTO) != ERROR_OK) //★CE# 信号を High にする(負論理、処理の終了を示す)
return ERROR_FAIL;
keep_alive();
return ERROR_OK;
}
// openocd/src/flash/nor/spi.h
/* SPI Flash Commands */
#define SPIFLASH_READ_ID 0x9F /* Read Flash Identification */
#define SPIFLASH_READ_MID 0xAF /* Read Flash Identification, multi-io */
#define SPIFLASH_READ_STATUS 0x05 /* Read Status Register */
#define SPIFLASH_WRITE_ENABLE 0x06 /* Write Enable */
#define SPIFLASH_PAGE_PROGRAM 0x02 /* Page Program */ //★コマンドはこれ
...
コマンド、データ、待機、たったこれだけです。前回はヘルパープログラムの制御などが出てきて複雑でしたが、実は SPI Flash の制御だけ見ると非常にシンプルです。
いよいよ、正常に動かない SPI Flash の消去を見ます。消去は書き込みと違い大量にデータを書く必要がなく、ヘルパープログラムを使った実装はありません。前回紹介した JTAG で SPI を直接制御する方法(簡易版)と似た処理です。
SPI Flash の消去には 3種類あります。ブロック消去に対応していないデバイスもあるようです。ややこしいことに統一された名前がないらしくて、メーカーによっては大きな消去単位をセクタと呼んでいます(Micron がそうだった)。なんじゃそりゃ、訳がわからんから統一してくれ〜。
細かい名前は覚えても意味がない(メーカーによって違うのでどうでも良い)ので、全部消す、大サイズで消す、小サイズで消す、くらいの理解で十分だと思います。
前置きはこれくらいにして、消去処理のコードを見ましょう。
// openocd/src/flash/nor/fespi.c
static int fespi_erase_sector(struct flash_bank *bank, int sector)
{
struct fespi_flash_bank *fespi_info = bank->driver_priv;
int retval;
retval = fespi_tx(bank, SPIFLASH_WRITE_ENABLE);
if (retval != ERROR_OK)
return retval;
retval = fespi_txwm_wait(bank);
if (retval != ERROR_OK)
return retval;
if (fespi_write_reg(bank, FESPI_REG_CSMODE, FESPI_CSMODE_HOLD) != ERROR_OK) //★CE# 信号を Low にする(負論理、処理の開始を示す)
return ERROR_FAIL;
retval = fespi_tx(bank, fespi_info->dev->erase_cmd); //★ブロック 64K 消去コマンド
if (retval != ERROR_OK)
return retval;
sector = bank->sectors[sector].offset;
retval = fespi_tx(bank, sector >> 16); //★アドレス
if (retval != ERROR_OK)
return retval;
retval = fespi_tx(bank, sector >> 8);
if (retval != ERROR_OK)
return retval;
retval = fespi_tx(bank, sector);
if (retval != ERROR_OK)
return retval;
retval = fespi_txwm_wait(bank); //★送信し終わるまで待つ
if (retval != ERROR_OK)
return retval;
if (fespi_write_reg(bank, FESPI_REG_CSMODE, FESPI_CSMODE_AUTO) != ERROR_OK) //★CE# 信号を High にする(負論理、処理の終了を示す)
return ERROR_FAIL;
retval = fespi_wip(bank, FESPI_MAX_TIMEOUT);
if (retval != ERROR_OK)
return retval;
return ERROR_OK;
}
ページ書き込みのときはコマンドが決め打ちでしたが、セクター消去ではコマンドが変数になっています。この変数が持つ値は、前回も紹介した 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),
...
//★HiFive1 が搭載している Flash
FLASH_ID("issi is25lp032", 0x03, 0x00, 0x02, 0xd8, 0xc7, 0x0016609d, 0x100, 0x10000, 0x400000),
FLASH_ID("issi is25lp064", 0x03, 0x00, 0x02, 0xd8, 0xc7, 0x0017609d, 0x100, 0x10000, 0x800000),
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),
//★HiFive Unleashed が搭載している Flash
FLASH_ID("issi is25wp256d", 0x13, 0xec, 0x12, 0xdc, 0xc7, 0x0019709d, 0x100, 0x10000, 0x2000000),
...
HiFive Unleashed が搭載している SPI Flash の品番は ISSI IS25WP256D ですから、erase_cmd は 0xdc です。SPI Flash のデータシートを見ると 0xdc は BLOCK ERASE OPERATION の 4BER64K コマンドです。
コードを見たときにお気づきかもしれませんが、SPI Flash の書き込みや消去で渡しているアドレスは 3バイトしかなく、最大で 16MB までしか扱えません。一方 HiFive Unleashed が搭載している SPI Flash の容量は 256Mbit = 32MB です。
足りないですね?どうやって 16MB 以降の消去や書き込みを行うのでしょうか?方法は 3つあるようです。
OpenOCD は 3番目を期待しているようです。なぜかというと 4BER64K コマンドは 4バイトアドレスを期待するコマンドだからです。
データシートから抜粋したコマンドの仕様は上記の通りです。アドレスを 4バイト送ってくることを期待しています。しかし実装は 3バイトしかアドレスを送っていません。うーん、何かおかしいですね?
これまでに書き込み、消去の仕組みを調べました。その過程でどうして HiFive Unleashed で SPI Flash の消去が正常に動作しないのかわかりました。アドレスの指定が間違っているからです。
// openocd/src/flash/nor/fespi.c
static int fespi_erase_sector(struct flash_bank *bank, int sector)
{
...
retval = fespi_tx(bank, fespi_info->dev->erase_cmd); //★ブロック 64K 消去コマンド(4バイトアドレスを期待)
if (retval != ERROR_OK)
return retval;
sector = bank->sectors[sector].offset;
retval = fespi_tx(bank, sector >> 16); //★アドレス(3バイトしか送っていない)
if (retval != ERROR_OK)
return retval;
retval = fespi_tx(bank, sector >> 8);
if (retval != ERROR_OK)
return retval;
retval = fespi_tx(bank, sector);
if (retval != ERROR_OK)
return retval;
消去コマンドは 0xdc(4BER64K コマンド)で 4バイトアドレスを期待していますが、今の OpenOCD の実装は 3バイトアドレスを渡しています。これでは動作しません。最初に動作確認したログを覚えているでしょうか?改めて見直すと、
$ 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
OpenOCD が 4バイトアドレスには対応していないよ、という警告を出していました。なるほど、やっと警告の意味がわかりました。
Unleashed の消去コマンドが動作しない理由はわかりましたが、まだいくつか不思議な点が残っています。
比較しやすいように SPI Flash のパラメータリストのうち HiFive1 と HiFive Unleashed に関わるところだけ抜粋しておきます。
// 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),
...
//★HiFive1 が搭載している Flash
FLASH_ID("issi is25lp032", 0x03, 0x00, 0x02, 0xd8, 0xc7, 0x0016609d, 0x100, 0x10000, 0x400000),
...
//★HiFive Unleashed が搭載している Flash
FLASH_ID("issi is25wp256d", 0x13, 0xec, 0x12, 0xdc, 0xc7, 0x0019709d, 0x100, 0x10000, 0x2000000),
...
結論としては OpenOCD の実装を変えないと HiFive Unleashed の SPI Flash の書き込み、消去は正常に行えないことがわかりました。手順では回避不可能です。
一覧が欲しくなってきたので作りました。
管理者: Katsuhiro Suzuki(katsuhiro( a t )katsuster.net)
This is Simple Diary 1.0
Copyright(C) Katsuhiro Suzuki 2006-2021.
Powered by PHP 5.2.17.
using GD bundled (2.0.34 compatible)(png support.)