saifthe1

/**
 * This kernel is invoked by VerletIntegrator to take one time step.
 */
class IntegrateVerletStepKernel : public KernelImpl {
public:
    static std::string Name() {
        return "IntegrateVerletStep";
    }
    IntegrateVerletStepKernel(std::string name, const Platform& platform) : KernelImpl(name, platform) {
    }

saifthe1

class Context;
class ContextImpl;

/**
 * An Integrator defines a method for simulating a System by integrating the equations of motion.
 * This is an abstract class.  Subclasses define particular integration methods.
 * 
 * Each Integrator object is bound to a particular Context which it integrates.  This connection
 * is specified by passing the Integrator as an argument to the constructor of the Context.
 */

saifthe1

namespace OpenMM {

/**
 * This is an Integrator which simulates a System using the velocity Verlet algorithm.
 */

class OPENMM_EXPORT VelocityVerletIntegrator : public Integrator {
public:
    /**
     * Create a VelocityVerletIntegrator.

saifthe1

#ifndef OPENMM_CONTROLTOOLS_H
#define OPENMM_CONTROLTOOLS_H

#include <stdio.h>
#include <iostream>
#include "openmm/Kernel.h"
#include "internal/windowsExport.h"

using std::string;

saifthe1

#ifndef OPENMM_MEASUREMENTTOOLS_H
#define OPENMM_MEASUREMENTTOOLS_H

#include<iostream>
#include<cmath>
#include<math.h>
#include "openmm/Tensor.h"
#include "openmm/Kernel.h"
#include "internal/windowsExport.h"

saifthe1

#include <iostream>
#include "openmm/Context.h"
#include "openmm/OpenMMException.h"
#include "openmm/internal/ContextImpl.h"
#include "openmm/kernels.h"
#include "openmm/ControlTools.h"


using namespace OpenMM;

saifthe1

void OpenCLMeasureBinVirialKernel::initialize(ContextImpl& impl){
        int pna = cl_.getPaddedNumAtoms();
        int numatoms = cl_.getNumAtoms();
    int numForceBuffers = cl_.getNumForceBuffers();
    
        numOfMolecules_ = cl_.getNumOfMolecules();
    
    cl::Kernel tempkernel = cl_.getNonbondedUtilities().getForceKernel();
    atomMasses_ = new OpenCLArray<mm_float4>(cl_,numOfMolecules_,"atomMasses",true);
    System& system = impl.getSystem();

saifthe1

/**
 * @todo: replace the preprocessor with dynamic values
 * using compiler prefix inside opencl
 */
#pragma OPENCL EXTENSION cl_khr_fp64 : enable


__kernel void binMomentum(
                          __global const float4* restrict velm,
                          __global const float4* restrict posq,

saifthe1

#ifdef SUPPORTS_DOUBLE_PRECISION
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#endif

/**
 * Perform the first step of velocity verlet integration.
 */

__kernel void velocityVerletPart1(int numAtoms, 
                                   __global const float* deltaT,

saifthe1

#ifdef SUPPORTS_64_BIT_ATOMICS
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
#endif

#define TILE_SIZE 32

// Cannot use float3 as OpenCL defines it to be 4 DWORD aligned. This would
// cause every element of array to have DWORD of padding to make it 4 DWORD
// aligned which wastes space and causes LDS bank conflicts as stride is no