fortran66のブログ

fortran について書きます。

Fortran 用 CUDA インターフェース

CUDA

イスラエルの企業 GASS 社がFortran 用 CUDA インターフェースを作っているようです。開発中のようですがメールするともらえます。マニュアル無しです。バイナリエディタでライブラリを覗けば、サブルーチン名は分かりますが・・
http://www.gass-ltd.co.il/en/Default.aspx
FORTRAN CUDA
http://www.gass-ltd.co.il/en/products/Fortran.aspx

FORTRAN77 っぽいバインディングです。著作権的にサンプル例を出していいのかよく分かりませんが、能書きを読む限りだしてもよさそうです。

コンパイラ・オプションとしては、小文字化＆アンダースコアの付加が必要のようです。ライブラリも Fortran 用と C 用のバッティングを避けるには少し小手先の工夫が必要です。この辺が C 言語との絡みの嫌なところ。Fortran2003 の C-binding が早くすっきりした形で普及してもらいたいです。

/names:lowercase /assume:underscore /libs:dll /threads

  • sample1: Copyright 2008 Company for Advanced Supercomputing Solutions Ltd (GASS).
C* 
* Copyright 2008 Company for Advanced Supercomputing Solutions Ltd
* (GASS).
* All rights reserved.
* http://www.gass-ltd.co.il
*
* NOTICE TO USER:
*
* This source code is subject to GASS ownership rights under U.S. and
* international Copyright laws.  Users and possessors of this source
* code are hereby granted a nonexclusive, royalty-free license to use this
* code in individual and commercial software.
*
* GASS MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
* CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
* IMPLIED WARRANTY OF ANY KIND.  GASS DISCLAIMS ALL WARRANTIES WITH
* REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
* MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR
* PURPOSE.
* IN NO EVENT SHALL GASS BE LIABLE FOR ANY SPECIAL, INDIRECT,
* INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM
* LOSS OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT,
* NEGLIGENCE OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE
* USE OR PERFORMANCE OF THIS SOURCE CODE.
*
* U.S. Government End Users.   This source code is a "commercial item"
* as that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of
* "commercial computer  software"  and "commercial computer software
* documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT
* 1995) and is provided to the U.S. Government only as a commercial end item.
* Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
* 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
* source code with only those rights set forth herein.
*
* Any use of this source code in individual and commercial software
* must include, in the user documentation and internal comments to the code,
* the above Disclaimer and U.S. Government End Users Notice.
*
      program cuda_fortran_test
      integer dev
      character*256 n
      print *,'Existing CUDA devices:'

      call cuInit(0)
      idevices = 0
      call cuDeviceGetCount(idevices)
      do i=1,idevices
            call cuDeviceGet(dev, i-1)
            call cuDeviceGetName(n, 256, dev)
            print *,'Device ',i-1,': ',n
      enddo
    • 実行結果

  • sample2: Copyright 2008 Company for Advanced Supercomputing Solutions Ltd (GASS).
C*
* Copyright 2008 Company for Advanced Supercomputing Solutions Ltd
* (GASS).
* All rights reserved.
* http://www.gass-ltd.co.il
*
* NOTICE TO USER:
*
* This source code is subject to GASS ownership rights under U.S. and
* international Copyright laws.  Users and possessors of this source
* code are hereby granted a nonexclusive, royalty-free license to use this
* code in individual and commercial software.
*
* GASS MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
* CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
* IMPLIED WARRANTY OF ANY KIND.  GASS DISCLAIMS ALL WARRANTIES WITH
* REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
* MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR
* PURPOSE.
* IN NO EVENT SHALL GASS BE LIABLE FOR ANY SPECIAL, INDIRECT,
* INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM
* LOSS OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT,
* NEGLIGENCE OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE
* USE OR PERFORMANCE OF THIS SOURCE CODE.
*
* U.S. Government End Users.   This source code is a "commercial item"
* as that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of
* "commercial computer  software"  and "commercial computer software
* documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT
* 1995) and is provided to the U.S. Government only as a commercial end item.
* Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
* 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
* source code with only those rights set forth herein.
*
* Any use of this source code in individual and commercial software
* must include, in the user documentation and internal comments to the code,
* the above Disclaimer and U.S. Government End Users Notice.
*
      program cuda_fortran_run_module
      integer data1(128,16), data2(16), res(128,16)

* Initialize the GPU, use the 1st (0)
      call cuInit(0)
      call cuDeviceGet(idev, 0)
      call cuCtxCreate(ictx, 0, idev)

* Load module and function
      call cuModuleLoad(imod,
     . 'C:\Users\Haru\Documents\Visual Studio 2005\Projects\CudaIl\ForCu
     .da2\test2.cubin')
      call cuModuleGetFunction(ifunc, imod, 'scale')

* Fill arrays
      do i=1,16
         do j=1,128
             data1(j,i) = 128*i + j
         enddo
      enddo
      do i=1,16
             data2(i) = 16-i 
      enddo

* Allocate device pointers
      call cuMemAlloc(idata1_ptr, 16*128*4)
      call cuMemAlloc(idata2_ptr, 16*4)

* Copy data to device
      call cuMemcpyHtoD(idata1_ptr, data1, 16*128*4)
      call cuMemcpyHtoD(idata2_ptr, data2, 16*4)

* Set function parameters
* Function parameters size is 16, because we use 2 pointers in 64 bit
* each being 8 bytes. For 32 platform, this will be 8.
      call cuParamSeti(ifunc, 0, idata1_ptr)
      call cuParamSeti(ifunc, 8, idata2_ptr)
      call cuParamSetSize(ifunc, 16)

* Launch the calculation on the GPU, use the 'y' axis of the block
      call cuFuncSetBlockShape(ifunc, 128, 1, 1)
      call cuLaunchGrid(ifunc, 1, 16, 1)
*      Utility function to check what result returned the last call to
*      a CUDA driver function
*      call GetLastCUDAResult(ires)
*      print *,ires

* Copy the results back
      call cuMemcpyDtoH(res, idata1_ptr, 16*128*4)

* Release all resources
      call cuMemFree(idata1_ptr)
      call cuMemFree(idata2_ptr)

* Verify results
      call verify_data(data1, data2, res)

c      return
      end program

* Verify the results from the GPU
      subroutine verify_data(data1, data2, res)
      integer data1(128,16), data2(16), res(128,16)
      integer test
      
      test = 1
      do i=1,16
         do j=1,128
             if(res(j,i).ne.(data1(j,i)*data2(i)))then
                  print *,res(j,i),(data1(j,i)*data2(i))
                  test=0
             endif
         enddo
      enddo

      if(test.eq.0)then
          print *,'Results do not match'
      else
          print *,'Execution was OK'
      endif

      return
      end

/*
 * Copyright 2008 Company for Advanced Supercomputing Solutions Ltd (GASS).
 * All rights reserved.
 * http://www.gass-ltd.co.il
 *
 * NOTICE TO USER:
 *
 * This source code is subject to GASS ownership rights under U.S. and
 * international Copyright laws.  Users and possessors of this source code
 * are hereby granted a nonexclusive, royalty-free license to use this code
 * in individual and commercial software.
 *
 * GASS MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE
 * CODE FOR ANY PURPOSE.  IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR
 * IMPLIED WARRANTY OF ANY KIND.  GASS DISCLAIMS ALL WARRANTIES WITH
 * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF
 * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
 * IN NO EVENT SHALL GASS BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL,
 * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS
 * OF USE, DATA OR PROFITS,  WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE
 * OR OTHER TORTIOUS ACTION,  ARISING OUT OF OR IN CONNECTION WITH THE USE
 * OR PERFORMANCE OF THIS SOURCE CODE.
 *
 * U.S. Government End Users.   This source code is a "commercial item" as
 * that term is defined at  48 C.F.R. 2.101 (OCT 1995), consisting  of
 * "commercial computer  software"  and "commercial computer software
 * documentation" as such terms are  used in 48 C.F.R. 12.212 (SEPT 1995)
 * and is provided to the U.S. Government only as a commercial end item.
 * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through
 * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the
 * source code with only those rights set forth herein.
 *
 * Any use of this source code in individual and commercial software must
 * include, in the user documentation and internal comments to the code,
 * the above Disclaimer and U.S. Government End Users Notice.
 */


/*
 * A simple example that scales values in data1, with the corresponding value
 * from data2.
 */
extern "C" __global__ void scale(int *data1, int *data2)
{
	int idx = blockDim.x * blockIdx.y + threadIdx.x;
	data1[idx] *= data2[blockIdx.y];
}
    • 実行結果

Open64

ところで、CUDA コンパイラの吐くアセンブリ・コードのようなものは、オープンソースの Open64 プロジェクトの成果物を土台にしているようです。このプロジェクトでは Fortran9x コンパイラも作っていたはずなので、NVIDIA も本気を出してくれれば CUDA 用の Fortran コンパイラがすぐにでも出るのではないかと期待してしまいます。
http://www.open64.net/