Fortran の配列を Swift 経由で GPU に送り計算する

Swift の文法もよくわからないまま、GPU へ突撃。GPU 用の機械語を吐き出すコンパイラの言語は C++14 になっています。

プログラム

以下では GPU 用のプログラム（拡張子 .metal）、 Swift 用のプログラム（拡張子 .swift）、Fortran 用のプログラム（拡張子 .90）をその順にコンパイル＆リンクして実行します。

なおエラーチェックをしないことで手抜きで書きました。また M1 Mac を前提として比較的新しい API や共有メモリーを無チェックで使っています。

xcode はエラーチェックの不備などがあると修正候補を示してくれるので、言われるままにクリックしてみました。びっくりマークやはてなマークが意味もよくわからぬままプイプイ挿入されてエラーが消えましたw　ちんねりむっつりマニュアル類を眺めて沈潜して黙考せねばならない古典プログラミング環境に対して、なんとなく浮草の如く漂ってゆけるモダンなプログラミング環境の片鱗を味わう事ができました。

手順メモ

metal 絡みのプログラムを Xcode を使わずにコマンドラインでコンパイルする方法がネット上で調べてもよく分かりませんでしたが、適当にやったら出来ました。

xcrun -sdk macosx metal mymetal.metal 

swiftc mymetal.swift -emit-library

gfortran mymetal.f90 libmymetal.dylib

GPU 用

GPU での計算を行うプログラム。１次元の配列構造を用います。

#include <metal_stdlib>

kernel void mykernel(device float* arr [[buffer(0)]],
                     constant int32_t &nsize [[buffer(1)]],
                     uint pos [[thread_position_in_grid]])
{
    if (pos < (uint)nsize) {
        arr[pos] = metal::sqrt(arr[pos]);
    }
}

Swift 用

Fortran から配列を受け取り、計算に用いるバイト列としてパラメータとともに GPU に送りこみます。また GPU 用のコンパイル済み機械語ルーチンも送りこんで計算させます。

import Metal

@_cdecl("mymetal")
public func mymetaln(n: Int32, px: UnsafeMutablePointer<Float32>) {
    let device = MTLCreateSystemDefaultDevice()
    let cmdQueue = device?.makeCommandQueue()
    let cmdBuff = cmdQueue?.makeCommandBuffer()
    let encoder = cmdBuff?.makeComputeCommandEncoder()
    
    // make pipeline for GPU binary
    let lib = device?.makeDefaultLibrary()
    let kernelSub: MTLFunction! = lib?.makeFunction(name: "mykernel")   
    let pipeline = try! device?.makeComputePipelineState(function: kernelSub)
    encoder?.setComputePipelineState(pipeline!)
   
    // make memory buffers
    var nsize = Int(n)
    let nlen = MemoryLayout<Float32>.size * nsize
    let buff0 = device?.makeBuffer(bytes: px, length: nlen, options: [.storageModeShared, .cpuCacheModeWriteCombined])
    encoder?.setBuffer(buff0, offset: 0, index: 0)
    let buff1 = device?.makeBuffer(bytes: &nsize, length: MemoryLayout<Int>.size, options: .storageModeShared)
    encoder?.setBuffer(buff1, offset: 0, index: 1)
       
    // group GPU threads
    //    non-uniform thread group assumed
    let w = pipeline?.threadExecutionWidth
    let perGroup = MTLSize(width: w!   , height: 1, depth: 1)
    let perGrid  = MTLSize(width: nsize, height: 1, depth: 1)
    encoder?.dispatchThreads(perGrid, threadsPerThreadgroup: perGroup)
    encoder?.endEncoding()
    
    // run GPU kernel
    cmdBuff?.commit()
    cmdBuff?.waitUntilCompleted()
    
    // to array
    let pbuff = buff0?.contents().bindMemory(to: Float32.self, capacity: nsize)
    for i in 0 ..< n {
        px[Int(i)] = pbuff![Int(i)]
    }

    print("GPU calculation finished")
}

Fortran 用

module test_m
    implicit none
    interface
        subroutine metal_calc(n, x) bind(c, name = 'mymetal')
            integer, value :: n
            real, intent(in out) :: x(n)
        end subroutine metal_calc
    end interface
end module test_m


program test
    use :: test_m
    implicit none
    integer, parameter :: n = 101
    integer :: i
    real :: x0(n), x1(n), eps = epsilon(0.0) 
    
    x0 = [(real(i), i = 0, n - 1)] 
    x1 = sqrt(x0)
    ! calclation by C++(Metal GPU) via Swift
    call metal_calc(n, x0)
    print *, 'By GPU (Metal C++ via Swift)'
    print '(5es15.7)', x0
    print *
    print *, 'By CPU (Fortran)'
    print '(5es15.7)', x1
    print *
    print '(5es15.7)', x1 - x0

    print *
    print *, 'Machine epsilon =', eps
    print *, '   2 epsilin,       4 epsilon,       8 epsilon'
    print *, 2 * eps, 4 * eps, 8 * eps
end program test

実行結果

0~100 までの平方根を Fortran 側と GPU 側で計算し結果を比較します。

所々、計算結果にわずかに違いがあって、その差の絶対値が machine epsilon の 2, 4, 8 倍になっています。丸めモードの違いにしては大きいような気もします。GPU 側は雑に計算しているのではないかと妄想します。この点に関しては心の片隅に留めておきたいと思います。

[a] M1:~/fortran/swift% xcrun -sdk macosx metal mymetal.metal 
[a] M1:~/fortran/swift% swiftc mymetal.swift -emit-library
[a] M1:~/fortran/swift% gfortran mymetal.f90 libmymetal.dylib 
[a] M1:~/fortran/swift% ./a.out                               
GPU calculation finished
 By GPU (Metal C++ via Swift)
  0.0000000E+00  1.0000000E+00  1.4142135E+00  1.7320508E+00  2.0000000E+00
  2.2360680E+00  2.4494898E+00  2.6457512E+00  2.8284271E+00  3.0000000E+00
  3.1622777E+00  3.3166249E+00  3.4641016E+00  3.6055512E+00  3.7416573E+00
  3.8729832E+00  4.0000000E+00  4.1231060E+00  4.2426405E+00  4.3588991E+00
  4.4721360E+00  4.5825758E+00  4.6904159E+00  4.7958317E+00  4.8989797E+00
  5.0000000E+00  5.0990195E+00  5.1961527E+00  5.2915025E+00  5.3851647E+00
  5.4772253E+00  5.5677643E+00  5.6568542E+00  5.7445626E+00  5.8309517E+00
  5.9160800E+00  6.0000000E+00  6.0827627E+00  6.1644139E+00  6.2449985E+00
  6.3245554E+00  6.4031243E+00  6.4807410E+00  6.5574389E+00  6.6332498E+00
  6.7082043E+00  6.7823300E+00  6.8556547E+00  6.9282031E+00  7.0000005E+00
  7.0710683E+00  7.1414289E+00  7.2111025E+00  7.2801104E+00  7.3484697E+00
  7.4161992E+00  7.4833145E+00  7.5498343E+00  7.6157737E+00  7.6811457E+00
  7.7459664E+00  7.8102503E+00  7.8740077E+00  7.9372540E+00  8.0000000E+00
  8.0622578E+00  8.1240387E+00  8.1853523E+00  8.2462120E+00  8.3066244E+00
  8.3666000E+00  8.4261503E+00  8.4852810E+00  8.5440035E+00  8.6023254E+00
  8.6602535E+00  8.7177982E+00  8.7749643E+00  8.8317604E+00  8.8881941E+00
  8.9442720E+00  9.0000000E+00  9.0553856E+00  9.1104336E+00  9.1651516E+00
  9.2195444E+00  9.2736187E+00  9.3273792E+00  9.3808317E+00  9.4339809E+00
  9.4868336E+00  9.5393925E+00  9.5916634E+00  9.6436510E+00  9.6953602E+00
  9.7467947E+00  9.7979593E+00  9.8488579E+00  9.8994951E+00  9.9498739E+00
  1.0000000E+01

 By CPU (Fortran)
  0.0000000E+00  1.0000000E+00  1.4142135E+00  1.7320508E+00  2.0000000E+00
  2.2360680E+00  2.4494898E+00  2.6457512E+00  2.8284271E+00  3.0000000E+00
  3.1622777E+00  3.3166249E+00  3.4641016E+00  3.6055512E+00  3.7416575E+00
  3.8729835E+00  4.0000000E+00  4.1231055E+00  4.2426405E+00  4.3588991E+00
  4.4721360E+00  4.5825758E+00  4.6904159E+00  4.7958317E+00  4.8989797E+00
  5.0000000E+00  5.0990195E+00  5.1961522E+00  5.2915025E+00  5.3851647E+00
  5.4772258E+00  5.5677643E+00  5.6568542E+00  5.7445626E+00  5.8309517E+00
  5.9160800E+00  6.0000000E+00  6.0827627E+00  6.1644139E+00  6.2449980E+00
  6.3245554E+00  6.4031243E+00  6.4807405E+00  6.5574384E+00  6.6332498E+00
  6.7082038E+00  6.7823300E+00  6.8556547E+00  6.9282031E+00  7.0000000E+00
  7.0710678E+00  7.1414285E+00  7.2111025E+00  7.2801099E+00  7.3484693E+00
  7.4161983E+00  7.4833150E+00  7.5498343E+00  7.6157732E+00  7.6811457E+00
  7.7459669E+00  7.8102498E+00  7.8740077E+00  7.9372540E+00  8.0000000E+00
  8.0622578E+00  8.1240387E+00  8.1853523E+00  8.2462111E+00  8.3066235E+00
  8.3666000E+00  8.4261494E+00  8.4852810E+00  8.5440035E+00  8.6023254E+00
  8.6602545E+00  8.7177982E+00  8.7749643E+00  8.8317604E+00  8.8881941E+00
  8.9442720E+00  9.0000000E+00  9.0553856E+00  9.1104336E+00  9.1651516E+00
  9.2195444E+00  9.2736187E+00  9.3273792E+00  9.3808317E+00  9.4339809E+00
  9.4868326E+00  9.5393925E+00  9.5916634E+00  9.6436510E+00  9.6953602E+00
  9.7467947E+00  9.7979593E+00  9.8488579E+00  9.8994951E+00  9.9498739E+00
  1.0000000E+01

  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  2.3841858E-07
  2.3841858E-07  0.0000000E+00 -4.7683716E-07  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00 -4.7683716E-07  0.0000000E+00  0.0000000E+00
  4.7683716E-07  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00 -4.7683716E-07
  0.0000000E+00  0.0000000E+00 -4.7683716E-07 -4.7683716E-07  0.0000000E+00
 -4.7683716E-07  0.0000000E+00  0.0000000E+00  0.0000000E+00 -4.7683716E-07
 -4.7683716E-07 -4.7683716E-07  0.0000000E+00 -4.7683716E-07 -4.7683716E-07
 -9.5367432E-07  4.7683716E-07  0.0000000E+00 -4.7683716E-07  0.0000000E+00
  4.7683716E-07 -4.7683716E-07  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00 -9.5367432E-07 -9.5367432E-07
  0.0000000E+00 -9.5367432E-07  0.0000000E+00  0.0000000E+00  0.0000000E+00
  9.5367432E-07  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
 -9.5367432E-07  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00  0.0000000E+00
  0.0000000E+00

 Machine epsilon =   1.19209290E-07
    2 epsilin,       4 epsilon,       8 epsilon
   2.38418579E-07   4.76837158E-07   9.53674316E-07
[a] M1:~/fortran/swift%

感想

一応、アイデアの原理的な検証を達しましたが、やっつけは好ましくないので少しづつマニュアル類を読んで正則な書き方が出来る様になれたらよかろうかと思います。

Swift で GPU を利用する部分は、昔のバッチジョブのジョブカードを書く類推で、プログラムの代わりにカーネル・サブルーチンを、入出力装置を装置番号に結び付ける代わりにサブルーチン引数を引数 index 番号に、それぞれパイプラインとメモリーバッファ上で結び付けて、ジョブカードに書く代わりに cmdbuffer 上にエンコーダーで積んで並べて、commit でジョブのキューにで投入と考えるとまぁなんとなくイメージがつかめなくもないです。つまりここでの swift は JCL (Job Control Language) なのでエラーで即死はむしろ吉かとｗ

それにつけても Fortran は気持ちいいですね。