いんちょき通信

PlatformとDeivce選択をOpenCLラップクラスに追加

「複数のOpenCLデバイスを使ってみる」で試した複数のデバイスを使うために、「外部ファイル読み取り関数をOpenCLラップクラスに追加」のOpenCLラップクラスにPlatform選択とDevice選択機能を追加した。

引数にPlatformを指定するインデックスとDeviceを指定するインデックスをとるコンストラクタを追加した。

        OpenCL::OpenCL(int platformIndex, int deviceIndex)
        {
            this->setPlatform();
            this->setContext(CL_DEVICE_TYPE_ALL, platformIndex);
            this->setDevice();
            this->queue_ = cl::CommandQueue(this->context_, this->devices_.at(deviceIndex), 0, &this->error_);
        }

この変更に伴って、setContext関数の引数を変更している。setContext関数は次の2種類が存在する。

• cl_int setContext()
• cl_int setContext(cl_device_type deviceType)

両方の引数にplatformIndexを追加し、次のように変更した。

• cl_int setContext(int platformIndex)
• cl_int setContext(cl_device_type deviceType, int platformIndex)

それぞれ、内部でcl_context_properties生成時のplataform指定を"0"固定から、platformIndexを使用するように変更している。それに伴い既存のコンストラクタ内でも明示的にplatformIndexを指定している。

変更後のソースとドキュメントはこちら。

• DECS_OpenCL_20120415をダウンロード
• Doclpp_20120415をダウンロード

このブログの開発用PCはこちら

 

外部ファイル読み取り関数をOpenCLラップクラスに追加

「OpenCLのソースを外部ファイルから読み込む」で試した外部のOpenCLソースを読み込む機能を「OpenCLをクラスでラップしてみた」、「OpenCLのラッパークラスのドキュメントを書いてみる」で書いたクラスに追加した。

• OpenCL.hpp

#ifndef __OPENCL_HPP__
#define __OPDNCL_HPP__
 
#if defined(__APPLE__) || defined(__MACOSX)
    #include <OpenCL/cl.hpp> // For Mac OSX
#else
    #include <CL/cl.hpp> // For Windows/Linux
#endif
#include <iostream>

#if defined(__ENABLE_BOOST__)
    #include <boost/filesystem/fstream.hpp>
#else
    #include <fstream>
#endif

namespace DECS{
    namespace OpenCL{

        #if defined(__ENABLE_BOOST__)
            typedef boost::filesystem::path PATH;
            typedef boost::filesystem::ifstream IFSTREAM;
        #else
            typedef std::string PATH;
            typedef std::ifstream IFSTREAM;
        #endif

        const cl_device_type DEFAULT_DEVICE_TYPE = CL_DEVICE_TYPE_GPU;

        class OpenCL{
        private:
            cl_int error_;
            std::vector platforms_;
            std::vector devices_;
            cl::Context context_;
            cl::CommandQueue queue_;
            cl::Program::Sources source_;
            cl::Program program_;
            cl::Kernel kernel_;
            std::vector argument_;
        protected:
            cl_int setPlatform();
            cl_int setContext();
            cl_int setContext(cl_device_type deviceType);
            cl_int setDevice();
        public:
            OpenCL();
            OpenCL(cl_device_type deviceType);
            ~OpenCL();
            void setSource(std::string& source);
            int setSourceFromFile(PATH srcFilePath);
            cl_int buildProgram();
            cl_int setKernel(std::string functionName);
            cl_int setKernelInputArgument(cl_int index, float* argument, const int n);
            cl_int setKernelOutputArgument(cl_int index, const int n);
            cl_int enqueueNDRange(const int n);
            cl_int readBuffer(cl_int index, float* output, const int n);
        };
    }
}

#endif // __OPENCL_HPP__

太字部分が今回追加した部分となる。プリプロセッサマクロでBoost C++ Libraryのfilesystemを使用するかどうかを選択できるようにした。"__ENABLE_BOOST__"を定義すると"boost::filesystem"を使用する。定義しない場合は標準ライブラリのstd::stringとstd::ifstreamを使用してソースファイルを読み込む。setSourceFronFileの中身は次のようになっている。

        int OpenCL::setSourceFromFile(PATH srcFilePath)
        {
            IFSTREAM srcFile(srcFilePath);
            if(!srcFile){
                return 1;
            }
            std::string src;
            std::string buf;
            while(std::getline(srcFile, buf)){
                src.append(buf);
            }
            this->setSource(src);
            return 0;
        }

前回「OpenCLのソースを外部ファイルから読み込む」で試したものとほとんど変更はない。関数内部でsetSourceを呼び出し、プライベート変数にソースを保存している。

この変更を含むすべてのソースはこちら→ DECS_OpenCL_20120414をダウンロード

前回の「OpenCLのラッパークラスのドキュメントを書いてみる」で書いたドキュメントには既に含まれている。

このブログの開発用PCはこちら

 

OpenCLのラッパークラスのドキュメントを書いてみる

「OpenCLをクラスでラップしてみた」で作成したOpenCLのラッパークラスのドキュメントをDoxygenを使って作成した。Doxygenについては「doxygenを使ってドキュメントを作成してみる」や「doxygenで複数言語のドキュメントを生成してみる」あたりでも話題にしている。

今回は複数言語の切り替えは特に行っておらず、基本的に英語で記載している。

Doxygenコメント入りのヘッダファイルが次のようになる。

• OpenCL.hpp

/******************************************************************************
* Copyright (c) 2012 INCHOKI Studio
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
*  1. Redistributions of source code must retain the above copyright notice,
*     this list of conditions and the following disclaimer.
*  2. Redistributions in binary form must reproduce the above copyright notice,
*     this list of conditions and the following disclaimer in the documentation
*     and/or other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
* LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
* CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
* SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
* INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
* CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE
*
******************************************************************************/

/*! \file
 *  \brief DECS CL for OpenCL C++ Bindings 1.1
 *  \author INCHOKI Studio
 *
 *  \version 0.1
 *  \date April 2012
 */

/*! \mainpage
 *  \section intro Introduction
 *  OpenCL enable to give effectivness to various kinds of software. It is good
 *  acpect of OpenCL. However, It has complex prosedure to use it for hobby
 *  programings and to go through its advantages. Many lines of code must be to
 *  written before the caluculations which you want are begun. The C++ classes
 *  including in this file help to use OpenCL easily.
 *
 *  The interface for wrapping OpenCL is contained with a single header file
 *  DECS/OpenCL/OpenCL.hpp and all definitions are within the namespace 
 *  DECS::CL. However, some definitions included DECS/OpenCL/OpenCL.hpp are
 *  implemented in a binary file, for example decsopencl.lib on Windows. The
 *  binary library should be lined when the software is built.
 *
 *  For detail or other information see:
 *
 *      WASABI Tokyo (Author's Blog in English)
 *      http://www.wasabi-tokyo.net/
 *
 *      INCHOKI Journal (Author's Blog in Japanese)
 *      http://blog.inchoki.com/
 *
 *  \section example Example
 *  The following example shows a simple use case for the DECS CL
 *  for OpenCL C++ Bindings.
 *
 *  \code
 * 
 *  #include "OpenCL.hpp"
 *  #include <iostream>
 *  #include <cstdlib>
 *
 *  const int nElements = 9000000;
 *  float input1[nElements];
 *  float input2[nElements];
 *  float output[nElements];
 *
 *  int main(int argc, char* argv[])
 *  {
 *      for(int i = 0; i < nElements; i++){
 *          input1[i] = (float)i * 10.0f;
 *          input2[i] = (float)i / 20.0f;
 *          output[i] = 0.0f;
 *      }
 *
 *      try{
 *          DECS::OpenCL::OpenCL ocl;
 *          ocl.setSourceFromFile("addVector.cl");
 *          ocl.buildProgram();
 *          ocl.setKernel(std::string("addVector"));
 *          ocl.setKernelInputArgument(0, input1, nElements);
 *          ocl.setKernelInputArgument(1, input2, nElements);
 *          ocl.setKernelOutputArgument(2, nElements);
 *          ocl.enqueueNDRange(nElements);
 *          ocl.readBuffer(2, output, nElements);
 *
 *          for(int i = 0; i < 20; i++){
 *              std::cout << "input1[" << i << "], input2[" << i << "], output[" << i << "] : ";
 *              std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl;
 *          }
 *
 *      }catch(cl::Error err){
 *           std::cerr << "ERROR: " << err.what() << "(" << err.err() << ")" << std::endl;
 *  
 *      }
 *
 *      return EXIT_SUCCESS;
 *  }
 *
 *  \endcode
 */
#ifndef __OPENCL_HPP__
#define __OPDNCL_HPP__
 
#if defined(__APPLE__) || defined(__MACOSX)
    #include <OpenCL/cl.hpp> // For Mac OSX
#else
    #include <CL/cl.hpp> // For Windows/Linux
#endif
#include <iostream>

#if defined(__ENABLE_BOOST__)
    #include <boost/filesystem/fstream.hpp>
#else
    #include <fstream>
#endif

/*! \namespace DECS
 *  \brief The top level namespace for Discrete Element Calculation Suite(DECS)
 *
 *  DECS(\em D escrete \em E lement \em C alculation \em S uite) is the set of
 *  utilities and libraries for Granular and Powder pysical simulation. Their
 *  definitions are contained within this namespace or sub-namespace under this
 *  namespace.
 */
namespace DECS{

/*! \namespace OpenCL
 *  \brief The DECS CL for OpenCL C++ Bindings is defined within this namespace
 *
 *  All Classes, structures, constants and macros are defined within this
 *  namespace. This namespace is one of sub-namespace of DECS.
 */
    namespace OpenCL{

        #if defined(__ENABLE_BOOST__)
            typedef boost::filesystem::path PATH;
            typedef boost::filesystem::ifstream IFSTREAM;
        #else
            typedef std::string PATH;
            typedef std::ifstream IFSTREAM;
        #endif

/*! \var cl_device_type DEFAULT_DEVICE_TYPE
 *  \brief This variable is the constant default device type.
 *
 *  DEFAULT_DEVICE_TYPE is defined for the default constructor of OpenCL Class.
 *  CL_DEVICE_TYPE_GPU is set in this variable. When the default constructor of
 *  OpenCL Class is used, only GPU devices are selected.
 */
        const cl_device_type DEFAULT_DEVICE_TYPE = CL_DEVICE_TYPE_GPU;

/*! \class OpenCL
 *  \brief This class is the main class of DECS CL for OpenCL C++ Bindings.
 *
 *  The OpenCL class supply to the series of procedures to use OpenCL in the
 *  softwares from initialization to finalization. The initialization process,
 *  for example getting platforms, contexts, devices and creating command
 *  queues is run in constructor.
 *
 *  \todo Preparing the function to load binary program (Priority: High)
 *  \todo Preparing the function to make kernel with binary program
 *  (Priority: High)
 */
        class OpenCL{
        private:
            cl_int error_;
            std::vector platforms_;
            std::vector devices_;
            cl::Context context_;
            cl::CommandQueue queue_;
            cl::Program::Sources source_;
            cl::Program program_;
            cl::Kernel kernel_;
            std::vector argument_;
        protected:
            cl_int setPlatform();
            cl_int setContext();
            cl_int setContext(cl_device_type deviceType);
            cl_int setDevice();
        public:

/*! \fn OpenCL()
 *  \brief The default constructor
 *
 *  This is the default constructor of DECS::OpenCL::OpenCL class. This
 *  constructor is initialize the class as following steps.
 *   -# All platforms set into the private variabe 
 *      std::vector platforms_ using the setPlatform
 *      protected function.
 *   -# The Context with DEFAULT_DEVICE_TYPE set into the private
 *      variable cl::Context context_ using the setContext protected
 *      function.
 *   -# All devices related with context_ set into the private varible 
 *      std::vector device_ using the setDevice protected
 *      function.
 *   -# The Command Queue related the first device in device_ set into
 *      the private varible cl::CommandQueue queue_.
 *
 *  \todo Change due to select the device creating the command queue (Priority: High)
 *  \todo Change due to select the platform creating the context (Priority: Low)
 */

            OpenCL();

/*! \fn OpenCL(cl_device_type deviceType)
 *  \brief The constructor with the device type
 *
 *  This is the constructor of DECS::OpenCL::OpenCL class with the device type.
 *  This constructor is initialize the class as following steps.
 *   -# All platforms set into the private variabe 
 *      std::vector platforms_ using the setPlatform
 *      protected function.
 *   -# The Context with the argument deviceType set into the private
 *      variable cl::Context context_ using the setContext protected
 *      function.
 *   -# All devices related with context_ set into the private varible 
 *      std::vector device_ using the setDevice protected
 *      function.
 *   -# The Command Queue related the first device in device_ set into
 *      the private varible cl::CommandQueue queue_.
 *
 *  @param[in] deviceType The device type of context.
 *  
 *  All parameters are defined in OpenCL.
 *
 *  For detail see: The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  \todo Change due to select the device creating the command queue (Priority: High)
 *  \todo Change due to select the platform creating the context (Priority: Low)
 */
            OpenCL(cl_device_type deviceType);

/*! \fn ~OpenCL()
 *  \brief The default destructor
 *
 *  The default destructor does not have the specific process at this time. 
 */
            ~OpenCL();

/*! \fn setSource(std::string& source)
 *  \brief The function to set source strings for the program
 *
 *  This function set the source strings into the private variable
 *  cl::Program::Sources source_. cl::Program::Sources is
 *  std::vector for source string and its size. When this function
 *  call second time and more, the source strings push back the std::vector.
 *  The size of std::vector is managed in this function, so the user don't need
 *  to conserned with the size.
 *
 *  @param[in] source The source string refference.
 */

            void setSource(std::string& source);

/*! \fn setSourceFromFile(std::string& srcFile)
 *  \brief The function to load and set a source file
 *
 *  This function load a souce from source file and set its contents into the
 *  private variable source_. This function call the function setSource
 *  in its internal process. If you define the macro "__ENABLE_BOOST__", the
 *  type "PATH" is "boost::filesystem::path". And "std::string" is used in the
 *  default settings.
 *
 *  @param[in] srcFilePath The path to source file.
 *  \return 0: Success, 1: "srcFilePath cannot be opened
 */
            int setSourceFromFile(PATH srcFilePath);

/*! \fn buildProgram()
 *  \brief The function to build the program
 *
 *  This fucntion initialize the private variable program_ and build the
 *  programs from the sources stored in the private variable source_. If an
 *  error occur in initializing program, this function return the error code of
 *  cl::Program and throw the exception. If an error occur in building program,
 *  this function return the error code of cl::Program::Build and throw the
 *  exception. These exception is defined in OpenCL C++ Bindings.
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  \return The error code of cl::Program or cl::Program::Build.
 */

            cl_int buildProgram();

/*! \fn setKernel(std::string functionName)
 *  \brief The function to set kernel
 *
 *  This function set the kernel with the OpenCL function included in the
 *  private variable program_instance into the private variable cl::Kernel
 *  kernel_. If an error occur in making the kernel instance, this function
 *  return the error code of cl::Kernel and throw the OpenCL C++ Bindings
 *  exception. 
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] functionName The function name to set into kernel.
 *  \return The error code of cl::Kernel.
 *
 *  \todo Change due to use binary program (Priority: High)
 */

            cl_int setKernel(std::string functionName);

/*! \fn setKernelInputArgument(cl_int index, float* argument, const int n)
 *  \brief The function to set the input argument into the private variable kernel_
 *
 *  This Function set the input argument for kernel into the private variable kernel_.
 *  The memory object is create based on the pointer of argument and the number
 *  of element of argument. In actual, this memory object is set into the kernel.
 *  If a error occur in making the memory object, this function return the error
 *  code of cl::Buffer and throw the exception. And if a error occur in setting
 *  the argument into the kernel, this function return the error code of
 *  cl::kernel::setArg and throw the exception.
 *
 *  \warning At this time this function can deal with only float type. And this 
 *  function can only deal with the memory object with only CL_MEM_READ_ONLY.

 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] index The index of argument order of the kernel.
 *  @param[in] argument The pointer of a argument for the kernel.
 *  @param[in] n The number of element of the argument.
 *  \return The error code of cl::Buffer or cl::Kernel::setArg.
 *
 *  \todo Change due to deal with other type argument (Priority: High)
 *  \todo Change due to switch multi kernels (Priority: Mid)
 */

            cl_int setKernelInputArgument(cl_int index, float* argument, const int n);

/*! \fn setKernelOutputArgument(cl_int index, const int n)
 *  \brief The function to set the output argument into the private variable kernel_
 *
 *  This Function set the output argument for kernel into the private variable kernel_.
 *  The memory object is create based on the pointer of argument and the number
 *  of element of argument. In actual, this memory object is set into the kernel.
 *  If a error occur in making the memory object, this function return the error
 *  code of cl::Buffer and throw the exception. And if a error occur in setting
 *  the argument into the kernel, this function return the error code of
 *  cl::kernel::setArg and throw the exception.
 *
 *  \warning At this time this function can deal with only float type. And this 
 *  function can only deal with the memory object with only CL_MEM_WRITE_ONLY.
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] index The index of argument order of the kernel.
 *  @param[in] n The number of element of the argument.
 *  \return The error code of cl::Buffer or cl::Kernel::setArg.
 *
 *  \todo Change due to deal with other type argument (Priority: High)
 *  \todo Change due to switch multi kernels (Priority: Mid)
 */
            cl_int setKernelOutputArgument(cl_int index, const int n);

/*! \fn enqueueNDRange(const int n)
 *
 *  \brief The function to enqueue the kernel
 *
 *  This function enqueue the private variable kernel_ into the command queue
 *  which created as the private variable queue_. If a error occur, this
 *  function return the error code of cl::CommandQueue::enqueueNDRangeKernel,
 *  and throw the exception of its.
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] n The number of index of the arguments.
 *
 *  \return The error code of cl::CommandQueue::enqueueNDRangeKernel
 */
            cl_int enqueueNDRange(const int n);

/*! \fn readBuffer(cl_int index, float* output, const int n)
 *  \brief The fucntion to read the buffer
 *
 *  This function read the output from the buffer into the argument
 *  "output". If a error occur duaring the reading buffer, this function
 *  return the error code of cl::CommandQueue::enqueueReadBuffer and throw
 *  the exception of cl::CommandQueue::enqueueReadBuffer and throw.

 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  \warning At this time this function can deal with only float type. 
 *
 *  @param[in] index The index of argument order of the kernel.
 *  @param[out] output The variable to store the output.
 *  @param[in] The number of element of the argument.
 *
 *  \return The error code of cl::CommandQueue::enqueueReadBuffer
 */
            cl_int readBuffer(cl_int index, float* output, const int n);
        };
    }
}

#endif // __OPENCL_HPP__

これからDoxygenを使って生成したドキュメントはこちら(doclpp_201204014.pdf)。

このブログの開発用PCはこちら

複数のOpenCLデバイスを使ってみる

IntelのOpenCLドライバもあるらしいので、インストールした流れで、AMD RadeonとCore i7両方を使って計算をしてみることにした。

そのまえに、ちゃんとIntelドライバを認識しているか、AMD Accelerated Parallel Processing (APP) SDKに付属している"clinfo.exe"で確認した。

C:\Program Files (x86)\AMD APP\bin\x86_64>clinfo.exe
Number of platforms:                             2
  Platform Profile:                              FULL_PROFILE
  Platform Version:                              OpenCL 1.1 AMD-APP (831.4)
  Platform Name:                                 AMD Accelerated Parallel Proces
sing
  Platform Vendor:                               Advanced Micro Devices, Inc.
  Platform Extensions:                           cl_khr_icd cl_amd_event_callbac
k cl_amd_offline_devices cl_khr_d3d10_sharing
  Platform Profile:                              FULL_PROFILE
  Platform Version:                              OpenCL 1.1
  Platform Name:                                 Intel(R) OpenCL
  Platform Vendor:                               Intel(R) Corporation
  Platform Extensions:                           cl_khr_fp64 cl_khr_global_int32
_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomi
cs cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_intel_pr
intf cl_ext_device_fission cl_intel_immediate_execution cl_khr_gl_sharing cl_khr
_icd

こんな感じの出力が出て、確かにIntel提供のプラットフォームが増えている。

「OpenCLで実際に計算してみる」のコードを単純にAMD用とIntel用に二重化して動かしてみた。

#include <cstdio>
#include <cmath>
#include <boost/chrono.hpp>
#include <boost/chrono/duration.hpp>

using namespace boost;

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif //__APPLE__

const int nElements = 9000000;
const int maxDevices  = 10;

float input1[nElements];
float input2[nElements];
float output[nElements];

void addVector(float* input1, float* input2, float* output){
	for(int i = 0; i < nElements; i++){
		output[i] = sin(input1[i]) * sin(input2[i]);
		output[i] = cos(output[i]);
		output[i] = pow(output[i], output[i]);
	}
}

int main(int argc, char* argv[])
{
    cl_int status;

    cl_platform_id platforms[10];
    cl_uint num_platforms;
    status = clGetPlatformIDs(sizeof(platforms) / sizeof(platforms[0]),
		platforms,
		&num_platforms);

    std::cout << "Number Of Platforms : " << num_platforms << std::endl;

    cl_context_properties amdProp[]
        = {CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[0], 0};

    cl_context_properties intelProp[]
        = {CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[1], 0};

    cl_context amdContext = clCreateContextFromType(amdProp,
        CL_DEVICE_TYPE_GPU,
        NULL,
        NULL,
        &status);

    cl_context intelContext = clCreateContextFromType(intelProp,
        CL_DEVICE_TYPE_CPU,
        NULL,
        NULL,
        &status);

    cl_device_id amdDev[maxDevices];
    size_t size_return;
    status = clGetContextInfo(amdContext,
		CL_CONTEXT_DEVICES,
		sizeof(amdDev),
		amdDev,
		&size_return);

    cl_device_id intelDev[maxDevices];
    status = clGetContextInfo(intelContext,
		CL_CONTEXT_DEVICES,
		sizeof(intelDev),
		intelDev,
		&size_return);

    cl_command_queue amdQueue = clCreateCommandQueue(amdContext,
        amdDev[0],
        0,
        &status);

    cl_command_queue intelQueue = clCreateCommandQueue(intelContext,
        intelDev[0],
        0,
        &status);

    const char *sources[] = {
        "__kernel void\n\
        addVector(__global const float *input1,\n\
        __global const float *input2,\n\
        __global float *output)\n\
        {\n\
        int index = get_global_id(0);\n\
		output[index] = sin(input1[index]) * sin(input2[index]);\n\
		output[index] = cos(output[index]);\n\
		output[index] = pow(output[index], output[index]);\n\
    }\n"};

    
    cl_program amdProgram = clCreateProgramWithSource(amdContext,
        1,
        (const char**)&sources,
        NULL,
        &status);

    cl_program intelProgram = clCreateProgramWithSource(intelContext,
        1,
        (const char**)&sources,
        NULL,
        &status);

    status = clBuildProgram(amdProgram, 1, amdDev, NULL, NULL, NULL);
    status = clBuildProgram(intelProgram, 1, intelDev, NULL, NULL, NULL);
    clUnloadCompiler();

    cl_kernel amdKernel = clCreateKernel(amdProgram, "addVector", &status);
    cl_kernel intelKernel = clCreateKernel(intelProgram, "addVector", &status);

    for (int i = 0; i < nElements; i++) {
        input1[i] = (float)i * 10.0f;
        input2[i] = (float)i / 20.0f;
        output[i] = 0.0f;
    }

    cl_mem amdMemInput1 = clCreateBuffer(amdContext,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		sizeof(cl_float) * nElements,
		input1,
		&status);

    cl_mem amdMemInput2 = clCreateBuffer(amdContext,
        CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        sizeof(cl_float) * nElements,
        input2,
        &status);

    cl_mem amdMemOutput = clCreateBuffer(amdContext,
        CL_MEM_WRITE_ONLY,
        sizeof(cl_float) * nElements,
        NULL,
        &status);

    cl_mem intelMemInput1 = clCreateBuffer(intelContext,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		sizeof(cl_float) * nElements,
		input1,
		&status);

    cl_mem intelMemInput2 = clCreateBuffer(intelContext,
        CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        sizeof(cl_float) * nElements,
        input2,
        &status);

    cl_mem intelMemOutput = clCreateBuffer(intelContext,
        CL_MEM_WRITE_ONLY,
        sizeof(cl_float) * nElements,
        NULL,
        &status);
    status = clSetKernelArg(amdKernel, 0, sizeof(cl_mem), (void *)&amdMemInput1);
    status = clSetKernelArg(amdKernel, 1, sizeof(cl_mem), (void *)&amdMemInput2);
    status = clSetKernelArg(amdKernel, 2, sizeof(cl_mem), (void *)&amdMemOutput);

    status = clSetKernelArg(intelKernel, 0, sizeof(cl_mem), (void *)&intelMemInput1);
    status = clSetKernelArg(intelKernel, 1, sizeof(cl_mem), (void *)&intelMemInput2);
    status = clSetKernelArg(intelKernel, 2, sizeof(cl_mem), (void *)&intelMemOutput);

	 boost::chrono::steady_clock::time_point start =
		 boost::chrono::steady_clock::now();

    size_t globalSize[] = {nElements};
    status = clEnqueueNDRangeKernel(amdQueue,
        amdKernel,
        1,
        NULL,
        globalSize,
        0, 0, NULL, NULL);

    status = clEnqueueReadBuffer(amdQueue,
        amdMemOutput,
        CL_TRUE,
        0,
        sizeof(cl_float) * nElements,
        output, 0, NULL, NULL);

	boost::chrono::duration sec =
		boost::chrono::steady_clock::now() - start;

    std::cout << "input1, input2, output" << std::endl;
    for (int i = 0; i < 100; i++) {
        std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl;
    }
	
    std::cout << "OpenCL addVector(AMD) " << sec.count() << " seconds\n";

	 boost::chrono::steady_clock::time_point start_i =
		 boost::chrono::steady_clock::now();

    status = clEnqueueNDRangeKernel(intelQueue,
        intelKernel,
        1,
        NULL,
        globalSize,
        0, 0, NULL, NULL);

    status = clEnqueueReadBuffer(intelQueue,
        intelMemOutput,
        CL_TRUE,
        0,
        sizeof(cl_float) * nElements,
        output, 0, NULL, NULL);

	boost::chrono::duration sec_i =
		boost::chrono::steady_clock::now() - start_i;

    std::cout << "input1, input2, output" << std::endl;
    for (int i = 0; i < 100; i++) {
        std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl;
    }
	
    std::cout << "OpenCL addVector(intel) " << sec_i.count() << " seconds\n";

    clReleaseMemObject(amdMemOutput);
    clReleaseMemObject(amdMemInput2);
    clReleaseMemObject(amdMemInput1);
    clReleaseKernel(amdKernel);
    clReleaseProgram(amdProgram);
    clReleaseCommandQueue(amdQueue);
    clReleaseContext(amdContext);
    clReleaseMemObject(intelMemOutput);
    clReleaseMemObject(intelMemInput2);
    clReleaseMemObject(intelMemInput1);
    clReleaseKernel(intelKernel);
    clReleaseProgram(intelProgram);
    clReleaseCommandQueue(intelQueue);
    clReleaseContext(intelContext);

	boost::chrono::steady_clock::time_point start2 =
		boost::chrono::steady_clock::now();

	addVector(input1, input2, output);

	boost::chrono::duration sec2 =
		boost::chrono::steady_clock::now() - start2;

	std::cout << "Normal addVector " << sec2.count() << " seconds\n";
    return 0;
}

amd*はRadeon用、intel*はCore i7用として、それぞれ実行時間を計測した。

実行時間部分の結果は以下のようになる。

OpenCL addVector(AMD) 0.129915 seconds
OpenCL addVector(intel) 0.114104 seconds
Normal addVector 1.12876 seconds

何回か実行したが、若干Intelプラットフォームの方が実行時間は短かった。

マルチスレッド化して、スレッドごとにデバイスを割り当てるなんてこともできそうな気がする。このあたりの結果を「OpenCLをクラスでラップしてみた」のクラスにでも反映していこうかなと思う。

このブログの開発用PCはこちら

微小時間刻み運動の要素パラメータ用データストア

前回の「微小時間刻み運動の要素をクラス化してみる」で微小時間刻み運動の要素をクラス化してみたが、どうにも使い勝手が悪そうなので、何かできないかと色々試した見ることにした。

ひとまず、データの柔軟性と将来的にOpenCLで並列処理をする際のデータの受け渡しなどを考えて、要素クラスからデータの実体を分離してみることにした。メインロジックとしては要素クラスを通してデータを操作するが、並列処理をする場合は全要素についてフラットなデータ空間の方が扱いやすいのではないかという予想してみた。なんか前にチラッと見たデザインパターンのフライウェイトとかにも近いような気がする。当然データを分離しているので、タイムステップの更新の際に実データの移動は発生せず、フライウェイトパターンと同じような利点もあると予想している。

実際のコードは次のようになっている。

• constant.hpp

#ifndef __CONSTANT_HPP__
#define __CONSTANT_HPP__

namespace DECS{
    // Constants for General
    typedef const unsigned int id_t;
    typedef const unsigned int dim_t;
    id_t DEFAULT_DIMENSION = 3;

    enum AXIS{
        X,
        Y,
        Z
    };

    enum RAXIS{
        RX,
        RY,
        RZ
    };

    // Constants for DataStore
    typedef unsigned int n_parameter_t;
    const n_parameter_t DEFAULT_PARAMETER_NUMBERS = 7;
    typedef unsigned int data_size_t;
    const data_size_t DEFAULT_DATASTORE_SIZE = 10;
    typedef unsigned int index_t;
}
#endif // __CONSTANT_HPP__

• DataStore.hpp

#ifndef __DATASTORE_HPP__
#define __DATASTORE_HPP__

#include <decs/constant.hpp>

namespace DECS{

    class DataStore{
    private:
        data_size_t     size_;
        data_size_t     wholeSize_;
        dim_t           dim_;
        n_parameter_t   nParam_;
        double*         data_;
    protected:
        void    fillZero();
    public:
        DataStore();
        DataStore(data_size_t size, dim_t dim, n_parameter_t nParam);
        DataStore(const DataStore& obj);
        ~DataStore();

        double* getStorePtr();
        double* getStorePtr(id_t id, index_t index);

        DataStore operator=(const DataStore& obj) const;
    };
}

#endif // __DATASTORE_HPP__

• DataStore.cpp

#include <decs/DataStore.hpp>

namespace DECS{
    DataStore::DataStore() : size_(DEFAULT_DATASTORE_SIZE),
                             dim_(DEFAULT_DIMENSION),
                             nParam_(DEFAULT_PARAMETER_NUMBERS)
    {
        this->wholeSize_ = this->size_ * this->dim_ * this->nParam_;
        this->data_ = new double[this->wholeSize_];
        this->fillZero();
    }

    DataStore::DataStore(data_size_t size, dim_t dim, n_parameter_t nParam) :
        size_(size), dim_(dim), nParam_(nParam)
    {
        this->wholeSize_ = this->size_ * this->dim_ * this->nParam_;
        this->data_ = new double[this->wholeSize_];
        this->fillZero();
    }

    DataStore::DataStore(const DataStore& obj) :
        size_(obj.size_), dim_(obj.dim_), nParam_(obj.nParam_), wholeSize_(obj.wholeSize_)
    {
        this->wholeSize_ = this->size_ * this->dim_ * this->nParam_;
        this->data_ = new double[this->wholeSize_];
        for(data_size_t i = 0; i < this->wholeSize_; i++){
            this->data_[i] = obj.data_[i];
        }
    }

    DataStore::~DataStore()
    {
        delete [] this->data_;
    }

    void    DataStore::fillZero()
    {
        for(data_size_t i = 0; i < this->wholeSize_; i++){
            this->data_[i] = 0.0;
        }
    }

    double* DataStore::getStorePtr()
    {
        return this->data_;
    }

    double* DataStore::getStorePtr(id_t id, index_t index)
    {
        if(id <= this->size_ && index <= this->nParam_){
            return this->data_ + id * this->nParam_ * this->dim_
                + index * this->dim_;
        }
        return 0;
    }

    DataStore DataStore::operator=(const DataStore& obj) const
    {
        DataStore rtn(obj);
        return rtn;
    }
}

constant.hppは単純に定数や列挙体をまとめて定義しているだけのヘッダーファイルとなっている。データ自体はdata_として1次元の配列として保持していて、getStorePtr(id_t　id, index_t index)でid番目の要素のindex番目の3次元ベクトル(デフォルトでは)パラメータのポインターを渡すようにした。

また、OpenCLにフラットなデータを渡せるようにgetStorePtr()でdata_の先頭のポインターを渡すようにしている。ただ、ポインターを渡してしまうならクラスじゃなくて構造体でいいのではないかとう気もする。うーん。

コピーコンストラクタとオペレータ"="は結果出力時にデータを対比することを想定している。

なんかおとなしくstd::vectorを使えばいいじゃないかという気もするが、std::vectorってそのままじゃOpenCLカーネルには渡せない気がするし、いちいちデータを作っていたらそれはそれでオーバーヘッドになるなと思うしで、色々迷うところだ。

このブログの開発用PCはこちら