FortranCL简单乘法

【FortranCL】FortranCL简单乘法

OpenCL简要介绍

OpenCL是面向由CPU、GPU和其他处理器组合构成的计算机进行编程的行业标准框架。OpenCL通过公布硬件来提供高度的可移植性（支持N卡、A卡、CPU等），而不是将硬件隐藏在精巧的抽象下，这表明OpenCL程序员必须显示的定义平台、上下文，以及在不同设备上的调度工作。
在一个典型的宿主机程序中，程序员不仅定义上下文和命令队列，定义内存和程序对象，还会构建宿主机上所需要的数据结构来支持应用。然后把重点转向命令队列，内存对象从宿主机移到设备上，内核参数关联到内存对象，然后提交到命令队列执行。内核完成工作时，计算中生成的内存对象可能会再复制到宿主机。
OpenCL框架基本工作流小结：

首先是一个定义上下文的宿主机程序，上图上下文包含两个Opencl设备，一个CPU一个GPU。接下来定义了命令队列，两个队列，一个面向GPU的有序命令队列，一个面向CPU的乱序命令队列。然后宿主机程序定义一个程序对象，这个程序对象编译后将为两个Opencl设备生成内核。接下来宿主机程序定义程序所需的内存对象，并将它们映射到内核的参数。最后，宿主机程序将命令放入命令队列来执行这些内核。

FortranCL简要介绍

• FortranCL是Fortran 90的OpenCL接口。它允许程序员直接从Fortran调用OpenCL。内核仍然用C语言编写。
  代码包github地址
• FortranCL的google官网只给出了一维向量和示例代码，虽然FortranCL接口尽可能接近OpenCL原来的API，也还是在参数等地方有些不同。例如函数clEnqueueNDRangeKernel在OpenCL中需要一个参数work_dim（说明工作项维度），而FortranCL是从global_work_size和local_work_size数组的维度获得的（它们必须具有相同的维度）。一些具体参数信息在FortranCL Wiki中有介绍。

构建FortranCL程序

• VS-IVF使用FortranCL静态库构建程序：
  • 创建一个Fortran控制台应用程序
  • 设置以下项目属性：
    • Fortran->语言->启用Fortran2003语义
    • Fortran->General->Additional Include Directories->解压缩FortranCL-x64.zip的目录
    • Linker->General-> Additional Include Directories->用于OpenCL安装的英特尔SDK的Lib\x64目录
    • Linker->Input->Additional Dependencies->OpenCL.lib
    • 讲解压缩的zip中的FortranCL.lib作为源文件添加到项目中
  • 注意：需要使用生成->配置管理 来创建新的x64平台配置
  • 项目目录结构（cl文件需要与.f90在同一文件夹下）：
  • 

代码示例

mult.f90

program main
use cl
implicit none
integer,parameter ::nx=1000,ny=1000
double precision :: ax(nx,ny),ay(nx,ny),res(nx,ny),res_cpu(nx,ny)
type(cl_platform_id)  ::platform
type(cl_device_id)    ::device
type(cl_context)      ::context
type(cl_command_queue) ::command_queue
type(cl_program)     ::prog
type(cl_kernel)      ::kernel
type(cl_mem)         ::cl_ax,cl_ay,cl_res
integer    :: num, ierr, irec ,i ,j,CPU_START,CPU_END,GPU_START,GPU_END,INI_START,INI_END,BUILDK_START,BUILDK_END,CBUFFER_START,CBUFFER_END,WBUFFER_START,WBUFFER_END,RUNK_START,RUNK_END
integer(8)   ::bytes,globalsize(2),localsize(2)
character(len = 100)  :: info
integer, parameter :: iunit = 10
integer, parameter :: source_length = 5000
character(len = source_length) :: source

  call RANDOM_NUMBER(ax)
  call RANDOM_NUMBER(ay)

  !=====================
  ! initialization
  !=====================

  call system_clock(GPU_START)
  ! get the platform id
  call system_clock(INI_START)
  call clgetplatformids(platform, num, ierr)
  if(ierr /= cl_success) stop "cannot get cl platform."

  ! get the device id
  call clgetdeviceids(platform, cl_device_type_all, device, num, ierr)
  if(ierr /= cl_success) stop "cannot get cl device."

  ! get the device name and print it
  call clgetdeviceinfo(device, cl_device_name, info, ierr)
  print*, "cl device: ", info

  !get max groupsize
  call clgetdeviceinfo(device, cl_device_max_work_group_size, info, ierr)
  ! create the context and the command queue
  context = clcreatecontext(platform, device, ierr)
  command_queue = clcreatecommandqueue(context, device, cl_queue_profiling_enable, ierr)
  call system_clock(INI_END)
  !=====================
  ! build the kernel
  !=====================
  call system_clock(BUILDK_START)
  ! read the source file  
  open(unit = iunit, file = 'D:\DongLY\vsproject\OpenCLmult\OpenCLmult\mult.cl', access='direct', status = 'old', action = 'read', iostat = ierr, recl = 1)
  write(*,*) ierr
  if (ierr /= 0) stop 'cannot open file mult.cl'

  source = ''
  irec = 1
  do
    read(unit = iunit, rec = irec, iostat = ierr) source(irec:irec)
    if (ierr /= 0) exit
    if(irec == source_length) stop 'Error: CL source file is too big'
    irec = irec + 1
  end do
  close(unit = iunit)
  
  ! create the program
  prog = clCreateProgramWithSource(context, source, ierr)
  if(ierr /= CL_SUCCESS) stop 'Error: cannot create program from source.'
  
  ! build
  call clBuildProgram(prog, '-cl-mad-enable', ierr)
  
  ! get the compilation log
  call clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG, source, irec) !将source作为返回错误信息的日志承载体
  if(len(trim(source)) > 0) write(*,*)trim(source) !若有错误信息打印出来
  !write(*,*) "ierr2",ierr
  if(ierr /= CL_SUCCESS) stop 'Error: program build failed.'
  ! finally get the kernel and release the program
  kernel = clCreateKernel(prog, 'mult', ierr)
  call clReleaseProgram(prog, ierr)
  call system_clock(BUILDK_END)
  
  !=====================
  ! RUN THE KERNEL
  !=====================
  bytes = int(nx, 8)*int(ny, 8)*8_8
  cl_ax = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, ierr)
  cl_ay = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, ierr)
  cl_res = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, ierr)
  call system_clock(WBUFFER_START)
  call clEnqueueWriteBuffer(command_queue, cl_ax, cl_bool(.true.), 0_8, bytes, ax(1,1), ierr)
  call clEnqueueWriteBuffer(command_queue, cl_ay, cl_bool(.true.), 0_8, bytes, ay(1,1), ierr)
  call system_clock(WBUFFER_END)
  call system_clock(CBUFFER_START)

  call clSetKernelArg(kernel, 0, cl_ax, ierr)
  call clSetKernelArg(kernel, 1, cl_ay, ierr)
  call clSetKernelArg(kernel, 2, cl_res, ierr)
  call clSetKernelArg(kernel, 3, nx, ierr)
  call clSetKernelArg(kernel, 4, ny, ierr)
  call system_clock(CBUFFER_END)
  
  localsize(1) = 32
  localsize(2) = 32
  globalsize(1) = int(ny, 8)
  globalsize(2) = int(nx, 8)
  if(mod(globalsize(1), localsize(1)) /= 0) globalsize(1) = globalsize(1) + localsize(1) - mod(globalsize(1), localsize(1)) 
  if(mod(globalsize(2), localsize(2)) /= 0) globalsize(2) = globalsize(2) + localsize(2) - mod(globalsize(2), localsize(2)) 
  write(*,*) "localsize",localsize
  write(*,*) "globalsize",globalsize

  ! execute the kernel
  call system_clock(RUNK_START)
  call clEnqueueNDRangeKernel(command_queue, kernel,globalsize, localsize, ierr)
  call clEnqueueReadBuffer(command_queue, cl_res, cl_bool(.true.), 0_8, bytes, res(1,1), ierr)
  write(*,*) ierr   ！打印出ierr可观察错误码，0为执行成功
  call system_clock(RUNK_END)
  call clFinish(command_queue, ierr)
  call clReleaseKernel(kernel, ierr)
  call clReleaseCommandQueue(command_queue, ierr)
  call clReleaseContext(context, ierr)
  call system_clock(GPU_END)
  
  !--------------------cpu----------------------------
  !---------------------------------------------------

  call system_clock(CPU_START)
   do j=2,ny
        do i=2,nx
            res_cpu(i,j) = ax(i,j)*ay(i,j)
           ! write(*,*) "res_cpu(",i,",",j,"),",(j-1)*nx+i-1,":",res_cpu(i,j)
        end do
    end do
  call system_clock(CPU_END)
 
  write(*,*) 'CPU time is ' , (CPU_END-CPU_START)/10000.
  write(*,*) 'GPU time is ' , (GPU_END-GPU_START)/10000.
  write(*,*) 'INIT TIME is ' , (INI_END-INI_START)/10000.
  write(*,*) 'BUILDK TIME is ' , (BUILDK_END-BUILDK_START)/10000.
  write(*,*) 'CREATE BUFFER TIME is ' , (CBUFFER_END-CBUFFER_START)/10000.
  write(*,*) 'WRITE BUFFER TIME is ' , (WBUFFER_END-WBUFFER_START)/10000.
  write(*,*) 'RUNK TIME is ' , (RUNK_END-RUNK_START)/10000.

pause

end program

mult.cl

__kernel void mult(__global double *ax, __global double *ay,__global double *res, const int nx, const int ny){
	int i = get_global_id(0);
	int j = get_global_id(1);
	int k = 1;
    if(i>=1&&i<=ny &&j>=1&&j<=nx)  res[i*nx + j] = ax[i*nx + j ] * ay[i*nx + j];
}

目前还存在一些问题，并且GPU时间明显慢于CPU（初始化创建上下文编译以及传输数据部分较为耗时）