FortranCL简单乘法

本文主要是介绍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) :: sourcecall RANDOM_NUMBER(ax)call RANDOM_NUMBER(ay)!=====================! initialization!=====================call system_clock(GPU_START)! get the platform idcall system_clock(INI_START)call clgetplatformids(platform, num, ierr)if(ierr /= cl_success) stop "cannot get cl platform."! get the device idcall clgetdeviceids(platform, cl_device_type_all, device, num, ierr)if(ierr /= cl_success) stop "cannot get cl device."! get the device name and print itcall clgetdeviceinfo(device, cl_device_name, info, ierr)print*, "cl device: ", info!get max groupsizecall clgetdeviceinfo(device, cl_device_max_work_group_size, info, ierr)! create the context and the command queuecontext = 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(*,*) ierrif (ierr /= 0) stop 'cannot open file mult.cl'source = ''irec = 1doread(unit = iunit, rec = irec, iostat = ierr) source(irec:irec)if (ierr /= 0) exitif(irec == source_length) stop 'Error: CL source file is too big'irec = irec + 1end doclose(unit = iunit)! create the programprog = clCreateProgramWithSource(context, source, ierr)if(ierr /= CL_SUCCESS) stop 'Error: cannot create program from source.'! buildcall clBuildProgram(prog, '-cl-mad-enable', ierr)! get the compilation logcall clGetProgramBuildInfo(prog, device, CL_PROGRAM_BUILD_LOG, source, irec) !将source作为返回错误信息的日志承载体if(len(trim(source)) > 0) write(*,*)trim(source) !若有错误信息打印出来!write(*,*) "ierr2",ierrif(ierr /= CL_SUCCESS) stop 'Error: program build failed.'! finally get the kernel and release the programkernel = clCreateKernel(prog, 'mult', ierr)call clReleaseProgram(prog, ierr)call system_clock(BUILDK_END)!=====================! RUN THE KERNEL!=====================bytes = int(nx, 8)*int(ny, 8)*8_8cl_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) = 32localsize(2) = 32globalsize(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",localsizewrite(*,*) "globalsize",globalsize! execute the kernelcall 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,nydo i=2,nxres_cpu(i,j) = ax(i,j)*ay(i,j)! write(*,*) "res_cpu(",i,",",j,"),",(j-1)*nx+i-1,":",res_cpu(i,j)end doend docall 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.pauseend 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];
}