Continuing implementation of OpenCL platform

platforms/opencl/CMakeLists.txt

@@ -82,4 +82,8 @@ ENDFOREACH(subdir)
 
 INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/src)
 
+# Install kernel files that will be loaded at runtime.
+FILE(GLOB OPENCL_KERNELS   ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/*.cl src/kernels/*.h)
+INSTALL_FILES(/lib/plugins/opencl FILES ${OPENCL_KERNELS})
+
 SUBDIRS (sharedTarget)

platforms/opencl/include/OpenCLPlatform.h

@@ -49,7 +49,7 @@ public:
         return name;
     }
     double getSpeed() const {
-        return 100;
+        return 0; // TODO Increase this.  Currently set to 0 so it will never be selected automatically.
     }
     bool supportsDoublePrecision() const;
     const std::string& getPropertyValue(const Context& context, const std::string& property) const;

platforms/opencl/src/OpenCLArray.h

@@ -51,11 +51,25 @@ public:
      *                          the OpenCL Buffer
      */
     OpenCLArray(OpenCLContext& context, int size, const std::string& name, bool createHostBuffer = false) :
-            context(context), size(size), name(name), local(createHostBuffer ? size : 0) {
+            context(context), size(size), name(name), local(createHostBuffer ? size : 0), ownsBuffer(true) {
         buffer = new cl::Buffer(context.getContext(), CL_MEM_READ_WRITE, size*sizeof(T));
     }
+    /**
+     * Create an OpenCLArray object the uses a preexisting Buffer.
+     *
+     * @param context           the context for which to create the array
+     * @param buffer            the OpenCL Buffer this object encapsulates
+     * @param size              the number of elements in the array
+     * @param name              the name of the array
+     * @param createHostBuffer  specifies whether to create a buffer in host memory for copying data to and from
+     *                          the OpenCL Buffer
+     */
+    OpenCLArray(OpenCLContext& context, cl::Buffer* buffer, int size, const std::string& name, bool createHostBuffer = false) :
+            context(context), buffer(buffer), size(size), name(name), local(createHostBuffer ? size : 0), ownsBuffer(false) {
+    }
     ~OpenCLArray() {
-        delete buffer;
+        if (ownsBuffer)
+            delete buffer;
     }
     const T& operator[](int index) const {
         return local[index];
@@ -63,6 +77,18 @@ public:
     T& operator[](int index) {
         return local[index];
     }
+    /**
+     * Get the size of the array.
+     */
+    int getSize() {
+        return size;
+    }
+    /**
+     * Get the OpenCL Buffer object.
+     */
+    cl::Buffer& getDeviceBuffer() {
+        return *buffer;
+    }
     /**
      * Get a pointer to the host buffer.
      */
@@ -114,6 +140,7 @@ private:
     cl::Buffer* buffer;
     std::vector<T> local;
     int size;
+    bool ownsBuffer;
     std::string name;
 };
 

platforms/opencl/src/OpenCLContext.cpp

@@ -26,25 +26,87 @@
 
 #include "OpenCLContext.h"
 #include "OpenCLArray.h"
+#include "openmm/Platform.h"
+#include <fstream>
+#include <iostream>
 
 using namespace OpenMM;
+using namespace std;
 
 OpenCLContext::OpenCLContext(int numParticles, int platformIndex, int deviceIndex) {
     // TODO Select the platform and device correctly
-    context = new cl::Context(CL_DEVICE_TYPE_CPU);
-    queue = new cl::CommandQueue(getContext(), getContext().getInfo<CL_CONTEXT_DEVICES>()[0]);
-    posq = new OpenCLArray<cl_float4>(*this, numParticles, "posq", true);
-    velm = new OpenCLArray<cl_float4>(*this, numParticles, "velm", true);
-    force = new OpenCLArray<cl_float4>(*this, numParticles, "force", true);
-    atomIndex = new OpenCLArray<cl_int>(*this, numParticles, "atomIndex", true);
+    context = cl::Context(CL_DEVICE_TYPE_CPU);
+    device = context.getInfo<CL_CONTEXT_DEVICES>()[0];
+    queue = cl::CommandQueue(context, device);
+    numAtoms = numParticles;
+    paddedNumAtoms = TileSize*((numParticles+TileSize-1)/TileSize);
+    numAtomBlocks = (paddedNumAtoms+(TileSize-1))/TileSize;
+    numTiles = numAtomBlocks*(numAtomBlocks+1)/2;
+    numThreadBlocks = 8*device.getInfo<CL_DEVICE_MAX_COMPUTE_UNITS>();
+    forceBufferPerWarp = true;
+    numForceBuffers = numThreadBlocks*ThreadBlockSize/TileSize;
+    if (numForceBuffers >= numAtomBlocks) {
+        // For small systems, it is more efficient to have one force buffer per block of 32 atoms instead of one per warp.
+
+        forceBufferPerWarp = false;
+        numForceBuffers = numAtomBlocks;
+    }
+    posq = new OpenCLArray<cl_float4>(*this, paddedNumAtoms, "posq", true);
+    velm = new OpenCLArray<cl_float4>(*this, paddedNumAtoms, "velm", true);
+    forceBuffers = new OpenCLArray<cl_float4>(*this, paddedNumAtoms*numForceBuffers, "forceBuffers", false);
+    force = new OpenCLArray<cl_float4>(*this, &forceBuffers->getDeviceBuffer(), paddedNumAtoms, "force", true);
+    atomIndex = new OpenCLArray<cl_int>(*this, paddedNumAtoms, "atomIndex", true);
+    for (int i = 0; i < paddedNumAtoms; ++i)
+        atomIndex->set(i, i);
+    atomIndex->upload();
+
+    // Create utility kernels that are used in multiple places.
+
+    utilities = createProgram(loadSourceFromFile("utilities.cl"));
+    clearBufferKernel = cl::Kernel(utilities, "clearBuffer");
 }
 
 OpenCLContext::~OpenCLContext() {
-    delete context;
-    delete queue;
     delete posq;
     delete velm;
     delete force;
     delete atomIndex;
 }
 
+string OpenCLContext::loadSourceFromFile(const string& filename) const {
+    ifstream file((Platform::getDefaultPluginsDirectory()+"/opencl/"+filename).c_str());
+    if (!file.is_open())
+        throw OpenMMException("Unable to load kernel: "+filename);
+    string kernel;
+    string line;
+    while (!file.eof()) {
+        getline(file, line);
+        kernel += line;
+        kernel += '\n';
+    }
+    file.close();
+    return kernel;
+}
+
+cl::Program OpenCLContext::createProgram(const std::string source) {
+    cl::Program::Sources sources(1, make_pair(source.c_str(), source.size()));
+    cl::Program program(context, sources);
+    try {
+        program.build(vector<cl::Device>(1, device));
+    } catch (cl::Error err) {
+        throw OpenMMException("Error compiling kernel: "+program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device));
+    }
+    return program;
+}
+
+void OpenCLContext::clearBuffer(OpenCLArray<float>& array) {
+    clearBufferKernel.setArg<cl::Buffer>(0, array.getDeviceBuffer());
+    clearBufferKernel.setArg<cl_int>(1, array.getSize());
+    queue.enqueueNDRangeKernel(clearBufferKernel, cl::NullRange, cl::NDRange(numThreadBlocks*ThreadBlockSize), cl::NDRange(ThreadBlockSize));
+}
+
+void OpenCLContext::clearBuffer(OpenCLArray<cl_float4>& array) {
+    clearBufferKernel.setArg<cl::Buffer>(0, array.getDeviceBuffer());
+    clearBufferKernel.setArg<cl_int>(1, array.getSize()*4);
+    queue.enqueueNDRangeKernel(clearBufferKernel, cl::NullRange, cl::NDRange(numThreadBlocks*ThreadBlockSize), cl::NDRange(ThreadBlockSize));
+}
\ No newline at end of file

platforms/opencl/src/OpenCLContext.h

@@ -41,19 +41,21 @@ class OpenCLArray;
 
 class OpenCLContext {
 public:
+    static const int ThreadBlockSize = 64;
+    static const int TileSize = 32;
     OpenCLContext(int numParticles, int platformIndex, int deviceIndex);
     ~OpenCLContext();
     /**
      * Get the cl::Context associated with this object.
      */
     cl::Context& getContext() {
-        return *context;
+        return context;
     }
     /**
      * Get the cl::CommandQueue associated with this object.
      */
     cl::CommandQueue& getQueue() {
-        return *queue;
+        return queue;
     }
     /**
      * Get the array which contains the position and charge of each atom.
@@ -73,18 +75,51 @@ public:
     OpenCLArray<cl_float4>& getForce() {
         return *force;
     }
+    /**
+     * Get the array which contains the buffers in which forces are computed.
+     */
+    OpenCLArray<cl_float4>& getForceBuffers() {
+        return *forceBuffers;
+    }
     /**
      * Get the array which contains the index of each atom.
      */
     OpenCLArray<cl_int>& getAtomIndex() {
         return *atomIndex;
     }
+    /**
+     * Load OpenCL source code from a file in the kernels directory.
+     */
+    std::string loadSourceFromFile(const std::string& filename) const;
+    /**
+     * Create an OpenCL Program from source code.
+     */
+    cl::Program createProgram(const std::string source);
+    /**
+     * Set all elements of an array to 0.
+     */
+    void clearBuffer(OpenCLArray<float>& array);
+    /**
+     * Set all elements of an array to 0.
+     */
+    void clearBuffer(OpenCLArray<cl_float4>& array);
+    int numAtoms;
+    int paddedNumAtoms;
+    int numAtomBlocks;
+    int numTiles;
+    int numThreadBlocks;
+    int numForceBuffers;
+    bool forceBufferPerWarp;
 private:
-    cl::Context* context;
-    cl::CommandQueue* queue;
+    cl::Context context;
+    cl::Device device;
+    cl::CommandQueue queue;
+    cl::Program utilities;
+    cl::Kernel clearBufferKernel;
     OpenCLArray<cl_float4>* posq;
     OpenCLArray<cl_float4>* velm;
     OpenCLArray<cl_float4>* force;
+    OpenCLArray<cl_float4>* forceBuffers;
     OpenCLArray<cl_int>* atomIndex;
 };
 

platforms/opencl/src/OpenCLKernelFactory.cpp

@@ -33,10 +33,10 @@ using namespace OpenMM;
 
 KernelImpl* OpenCLKernelFactory::createKernelImpl(std::string name, const Platform& platform, ContextImpl& context) const {
     OpenCLPlatform::PlatformData& data = *static_cast<OpenCLPlatform::PlatformData*>(context.getPlatformData());
-//    if (name == InitializeForcesKernel::Name())
-//        return new OpenCLInitializeForcesKernel(name, platform);
-//    if (name == UpdateTimeKernel::Name())
-//        return new OpenCLUpdateTimeKernel(name, platform, data);
+    if (name == InitializeForcesKernel::Name())
+        return new OpenCLInitializeForcesKernel(name, platform, data);
+    if (name == UpdateTimeKernel::Name())
+        return new OpenCLUpdateTimeKernel(name, platform, data);
 //    if (name == CalcHarmonicBondForceKernel::Name())
 //        return new OpenCLCalcHarmonicBondForceKernel(name, platform, data, context.getSystem());
 //    if (name == CalcHarmonicAngleForceKernel::Name())
@@ -51,8 +51,8 @@ KernelImpl* OpenCLKernelFactory::createKernelImpl(std::string name, const Platfo
 //        return new OpenCLCalcCustomNonbondedForceKernel(name, platform, data, context.getSystem());
 //    if (name == CalcGBSAOBCForceKernel::Name())
 //        return new OpenCLCalcGBSAOBCForceKernel(name, platform, data);
-//    if (name == IntegrateVerletStepKernel::Name())
-//        return new OpenCLIntegrateVerletStepKernel(name, platform, data);
+    if (name == IntegrateVerletStepKernel::Name())
+        return new OpenCLIntegrateVerletStepKernel(name, platform, data);
 //    if (name == IntegrateLangevinStepKernel::Name())
 //        return new OpenCLIntegrateLangevinStepKernel(name, platform, data);
 //    if (name == IntegrateBrownianStepKernel::Name())
@@ -63,8 +63,8 @@ KernelImpl* OpenCLKernelFactory::createKernelImpl(std::string name, const Platfo
 //        return new OpenCLIntegrateVariableLangevinStepKernel(name, platform, data);
 //    if (name == ApplyAndersenThermostatKernel::Name())
 //        return new OpenCLApplyAndersenThermostatKernel(name, platform, data);
-//    if (name == CalcKineticEnergyKernel::Name())
-//        return new OpenCLCalcKineticEnergyKernel(name, platform);
+    if (name == CalcKineticEnergyKernel::Name())
+        return new OpenCLCalcKineticEnergyKernel(name, platform);
 //    if (name == RemoveCMMotionKernel::Name())
 //        return new OpenCLRemoveCMMotionKernel(name, platform, data);
     throw OpenMMException((std::string("Tried to create kernel with illegal kernel name '")+name+"'").c_str());

platforms/opencl/src/OpenCLKernels.h

@@ -33,55 +33,57 @@
 
 namespace OpenMM {
 
-///**
-// * This kernel is invoked at the start of each force evaluation to clear the forces.
-// */
-//class OpenCLInitializeForcesKernel : public InitializeForcesKernel {
-//public:
-//    OpenCLInitializeForcesKernel(std::string name, const Platform& platform) : InitializeForcesKernel(name, platform) {
-//    }
-//    /**
-//     * Initialize the kernel.
-//     *
-//     * @param system     the System this kernel will be applied to
-//     */
-//    void initialize(const System& system);
-//    /**
-//     * Execute the kernel.
-//     *
-//     * @param context    the context in which to execute this kernel
-//     */
-//    void execute(ContextImpl& context);
-//};
-//
-///**
-// * This kernel is invoked to get or set the current time.
-// */
-//class OpenCLUpdateTimeKernel : public UpdateTimeKernel {
-//public:
-//    OpenCLUpdateTimeKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data) : UpdateTimeKernel(name, platform), data(data) {
-//    }
-//    /**
-//     * Initialize the kernel.
-//     *
-//     * @param system     the System this kernel will be applied to
-//     */
-//    void initialize(const System& system);
-//    /**
-//     * Get the current time (in picoseconds).
-//     *
-//     * @param context    the context in which to execute this kernel
-//     */
-//    double getTime(const ContextImpl& context) const;
-//    /**
-//     * Set the current time (in picoseconds).
-//     *
-//     * @param context    the context in which to execute this kernel
-//     */
-//    void setTime(ContextImpl& context, double time);
-//private:
-//    OpenCLPlatform::PlatformData& data;
-//};
+/**
+ * This kernel is invoked at the start of each force evaluation to clear the forces.
+ */
+class OpenCLInitializeForcesKernel : public InitializeForcesKernel {
+public:
+    OpenCLInitializeForcesKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data) : InitializeForcesKernel(name, platform), data(data) {
+    }
+    /**
+     * Initialize the kernel.
+     *
+     * @param system     the System this kernel will be applied to
+     */
+    void initialize(const System& system);
+    /**
+     * Execute the kernel.
+     *
+     * @param context    the context in which to execute this kernel
+     */
+    void execute(ContextImpl& context);
+private:
+   OpenCLPlatform::PlatformData& data;
+};
+
+/**
+ * This kernel is invoked to get or set the current time.
+ */
+class OpenCLUpdateTimeKernel : public UpdateTimeKernel {
+public:
+    OpenCLUpdateTimeKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data) : UpdateTimeKernel(name, platform), data(data) {
+    }
+    /**
+     * Initialize the kernel.
+     *
+     * @param system     the System this kernel will be applied to
+     */
+    void initialize(const System& system);
+    /**
+     * Get the current time (in picoseconds).
+     *
+     * @param context    the context in which to execute this kernel
+     */
+    double getTime(const ContextImpl& context) const;
+    /**
+     * Set the current time (in picoseconds).
+     *
+     * @param context    the context in which to execute this kernel
+     */
+    void setTime(ContextImpl& context, double time);
+private:
+    OpenCLPlatform::PlatformData& data;
+};
 //
 ///**
 // * This kernel is invoked by HarmonicBondForce to calculate the forces acting on the system and the energy of the system.
@@ -321,34 +323,34 @@ namespace OpenMM {
 //private:
 //    OpenCLPlatform::PlatformData& data;
 //};
-//
-///**
-// * This kernel is invoked by VerletIntegrator to take one time step.
-// */
-//class OpenCLIntegrateVerletStepKernel : public IntegrateVerletStepKernel {
-//public:
-//    OpenCLIntegrateVerletStepKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data) : IntegrateVerletStepKernel(name, platform), data(data) {
-//    }
-//    ~OpenCLIntegrateVerletStepKernel();
-//    /**
-//     * Initialize the kernel.
-//     *
-//     * @param system     the System this kernel will be applied to
-//     * @param integrator the VerletIntegrator this kernel will be used for
-//     */
-//    void initialize(const System& system, const VerletIntegrator& integrator);
-//    /**
-//     * Execute the kernel.
-//     *
-//     * @param context    the context in which to execute this kernel
-//     * @param integrator the VerletIntegrator this kernel is being used for
-//     */
-//    void execute(ContextImpl& context, const VerletIntegrator& integrator);
-//private:
-//    OpenCLPlatform::PlatformData& data;
-//    double prevStepSize;
-//};
-//
+
+/**
+ * This kernel is invoked by VerletIntegrator to take one time step.
+ */
+class OpenCLIntegrateVerletStepKernel : public IntegrateVerletStepKernel {
+public:
+    OpenCLIntegrateVerletStepKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data) : IntegrateVerletStepKernel(name, platform), data(data) {
+    }
+    ~OpenCLIntegrateVerletStepKernel();
+    /**
+     * Initialize the kernel.
+     *
+     * @param system     the System this kernel will be applied to
+     * @param integrator the VerletIntegrator this kernel will be used for
+     */
+    void initialize(const System& system, const VerletIntegrator& integrator);
+    /**
+     * Execute the kernel.
+     *
+     * @param context    the context in which to execute this kernel
+     * @param integrator the VerletIntegrator this kernel is being used for
+     */
+    void execute(ContextImpl& context, const VerletIntegrator& integrator);
+private:
+    OpenCLPlatform::PlatformData& data;
+    double prevStepSize;
+};
+
 ///**
 // * This kernel is invoked by LangevinIntegrator to take one time step.
 // */
@@ -484,30 +486,30 @@ namespace OpenMM {
 //    OpenCLPlatform::PlatformData& data;
 //    double prevTemp, prevFrequency, prevStepSize;
 //};
-//
-///**
-// * This kernel is invoked to calculate the kinetic energy of the system.
-// */
-//class OpenCLCalcKineticEnergyKernel : public CalcKineticEnergyKernel {
-//public:
-//    OpenCLCalcKineticEnergyKernel(std::string name, const Platform& platform) : CalcKineticEnergyKernel(name, platform) {
-//    }
-//    /**
-//     * Initialize the kernel.
-//     *
-//     * @param system     the System this kernel will be applied to
-//     */
-//    void initialize(const System& system);
-//    /**
-//     * Execute the kernel.
-//     *
-//     * @param context    the context in which to execute this kernel
-//     */
-//    double execute(ContextImpl& context);
-//private:
-//    std::vector<double> masses;
-//};
-//
+
+/**
+ * This kernel is invoked to calculate the kinetic energy of the system.
+ */
+class OpenCLCalcKineticEnergyKernel : public CalcKineticEnergyKernel {
+public:
+    OpenCLCalcKineticEnergyKernel(std::string name, const Platform& platform) : CalcKineticEnergyKernel(name, platform) {
+    }
+    /**
+     * Initialize the kernel.
+     *
+     * @param system     the System this kernel will be applied to
+     */
+    void initialize(const System& system);
+    /**
+     * Execute the kernel.
+     *
+     * @param context    the context in which to execute this kernel
+     */
+    double execute(ContextImpl& context);
+private:
+    std::vector<double> masses;
+};
+
 ///**
 // * This kernel is invoked to remove center of mass motion from the system.
 // */

platforms/opencl/src/OpenCLPlatform.cpp

@@ -45,8 +45,8 @@ extern "C" void initOpenMMPlugin() {
 
 OpenCLPlatform::OpenCLPlatform() {
     OpenCLKernelFactory* factory = new OpenCLKernelFactory();
-//    registerKernelFactory(InitializeForcesKernel::Name(), factory);
-//    registerKernelFactory(UpdateTimeKernel::Name(), factory);
+    registerKernelFactory(InitializeForcesKernel::Name(), factory);
+    registerKernelFactory(UpdateTimeKernel::Name(), factory);
 //    registerKernelFactory(CalcHarmonicBondForceKernel::Name(), factory);
 //    registerKernelFactory(CalcHarmonicAngleForceKernel::Name(), factory);
 //    registerKernelFactory(CalcPeriodicTorsionForceKernel::Name(), factory);
@@ -54,13 +54,13 @@ OpenCLPlatform::OpenCLPlatform() {
 //    registerKernelFactory(CalcNonbondedForceKernel::Name(), factory);
 //    registerKernelFactory(CalcCustomNonbondedForceKernel::Name(), factory);
 //    registerKernelFactory(CalcGBSAOBCForceKernel::Name(), factory);
-//    registerKernelFactory(IntegrateVerletStepKernel::Name(), factory);
+    registerKernelFactory(IntegrateVerletStepKernel::Name(), factory);
 //    registerKernelFactory(IntegrateLangevinStepKernel::Name(), factory);
 //    registerKernelFactory(IntegrateBrownianStepKernel::Name(), factory);
 //    registerKernelFactory(IntegrateVariableVerletStepKernel::Name(), factory);
 //    registerKernelFactory(IntegrateVariableLangevinStepKernel::Name(), factory);
 //    registerKernelFactory(ApplyAndersenThermostatKernel::Name(), factory);
-//    registerKernelFactory(CalcKineticEnergyKernel::Name(), factory);
+    registerKernelFactory(CalcKineticEnergyKernel::Name(), factory);
 //    registerKernelFactory(RemoveCMMotionKernel::Name(), factory);
     platformProperties.push_back(OpenCLPlatformIndex());
     platformProperties.push_back(OpenCLDeviceIndex());

platforms/opencl/src/kernels/utilities.cl

+__kernel void clearBuffer(__global float* buffer, int size) {
+    int index = get_global_id(0);
+    int step = get_global_size(0);
+    __global float4* buffer4 = (__global float4*) buffer;
+    int sizeDiv4 = size/4;
+    while (index < sizeDiv4) {
+        buffer4[index] = (float4) (0.0f);
+        index += step;
+    }
+    if (get_global_id(0) == 0)
+        for (int i = sizeDiv4*4; i < size; i++)
+            buffer[i] = 0.0f;
+}