Merge branch 'master' of https://bitbucket.org/erwincoumans/bullet3_experiments

Conflicts:
	btgui/OpenGLWindow/GLInstancingRenderer.cpp

btgui/MidiTest/LICENSE.txt

@@ -0,0 +1,29 @@
+    RtMidi WWW site: http://music.mcgill.ca/~gary/rtmidi/
+
+    RtMidi: realtime MIDI i/o C++ classes
+    Copyright (c) 2003-2012 Gary P. Scavone
+
+    Permission is hereby granted, free of charge, to any person
+    obtaining a copy of this software and associated documentation files
+    (the "Software"), to deal in the Software without restriction,
+    including without limitation the rights to use, copy, modify, merge,
+    publish, distribute, sublicense, and/or sell copies of the Software,
+    and to permit persons to whom the Software is furnished to do so,
+    subject to the following conditions:
+
+    The above copyright notice and this permission notice shall be
+    included in all copies or substantial portions of the Software.
+
+    Any person wishing to distribute modifications to the Software is
+    asked to send the modifications to the original developer so that
+    they can be incorporated into the canonical version.  This is,
+    however, not a binding provision of this license.
+
+    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+    EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+    MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+    IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR
+    ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF
+    CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION
+    WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+*/

btgui/MidiTest/RtError.h

@@ -0,0 +1,60 @@
+/************************************************************************/
+/*! \class RtError
+    \brief Exception handling class for RtAudio & RtMidi.
+
+    The RtError class is quite simple but it does allow errors to be
+    "caught" by RtError::Type. See the RtAudio and RtMidi
+    documentation to know which methods can throw an RtError.
+
+*/
+/************************************************************************/
+
+#ifndef RTERROR_H
+#define RTERROR_H
+
+#include <exception>
+#include <iostream>
+#include <string>
+
+class RtError : public std::exception
+{
+ public:
+  //! Defined RtError types.
+  enum Type {
+    WARNING,           /*!< A non-critical error. */
+    DEBUG_WARNING,     /*!< A non-critical error which might be useful for debugging. */
+    UNSPECIFIED,       /*!< The default, unspecified error type. */
+    NO_DEVICES_FOUND,  /*!< No devices found on system. */
+    INVALID_DEVICE,    /*!< An invalid device ID was specified. */
+    MEMORY_ERROR,      /*!< An error occured during memory allocation. */
+    INVALID_PARAMETER, /*!< An invalid parameter was specified to a function. */
+    INVALID_USE,       /*!< The function was called incorrectly. */
+    DRIVER_ERROR,      /*!< A system driver error occured. */
+    SYSTEM_ERROR,      /*!< A system error occured. */
+    THREAD_ERROR       /*!< A thread error occured. */
+  };
+
+  //! The constructor.
+  RtError( const std::string& message, Type type = RtError::UNSPECIFIED ) throw() : message_(message), type_(type) {}
+ 
+  //! The destructor.
+  virtual ~RtError( void ) throw() {}
+
+  //! Prints thrown error message to stderr.
+  virtual void printMessage( void ) const throw() { std::cerr << '\n' << message_ << "\n\n"; }
+
+  //! Returns the thrown error message type.
+  virtual const Type& getType(void) const throw() { return type_; }
+
+  //! Returns the thrown error message string.
+  virtual const std::string& getMessage(void) const throw() { return message_; }
+
+  //! Returns the thrown error message as a c-style string.
+  virtual const char* what( void ) const throw() { return message_.c_str(); }
+
+ protected:
+  std::string message_;
+  Type type_;
+};
+
+#endif

btgui/MidiTest/RtMidi.h

@@ -0,0 +1,675 @@
+/**********************************************************************/
+/*! \class RtMidi
+    \brief An abstract base class for realtime MIDI input/output.
+
+    This class implements some common functionality for the realtime
+    MIDI input/output subclasses RtMidiIn and RtMidiOut.
+
+    RtMidi WWW site: http://music.mcgill.ca/~gary/rtmidi/
+
+    RtMidi: realtime MIDI i/o C++ classes
+    Copyright (c) 2003-2012 Gary P. Scavone
+
+    Permission is hereby granted, free of charge, to any person
+    obtaining a copy of this software and associated documentation files
+    (the "Software"), to deal in the Software without restriction,
+    including without limitation the rights to use, copy, modify, merge,
+    publish, distribute, sublicense, and/or sell copies of the Software,
+    and to permit persons to whom the Software is furnished to do so,
+    subject to the following conditions:
+
+    The above copyright notice and this permission notice shall be
+    included in all copies or substantial portions of the Software.
+
+    Any person wishing to distribute modifications to the Software is
+    asked to send the modifications to the original developer so that
+    they can be incorporated into the canonical version.  This is,
+    however, not a binding provision of this license.
+
+    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+    EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+    MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+    IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR
+    ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF
+    CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION
+    WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+*/
+/**********************************************************************/
+
+/*!
+  \file RtMidi.h
+ */
+
+// RtMidi: Version 2.0.1
+
+#ifndef RTMIDI_H
+#define RTMIDI_H
+
+#include "RtError.h"
+#include <string>
+#include <vector>
+
+class RtMidi
+{
+ public:
+
+  //! MIDI API specifier arguments.
+  enum Api {
+    UNSPECIFIED,    /*!< Search for a working compiled API. */
+    MACOSX_CORE,    /*!< Macintosh OS-X Core Midi API. */
+    LINUX_ALSA,     /*!< The Advanced Linux Sound Architecture API. */
+    UNIX_JACK,      /*!< The Jack Low-Latency MIDI Server API. */
+    WINDOWS_MM,     /*!< The Microsoft Multimedia MIDI API. */
+    WINDOWS_KS,     /*!< The Microsoft Kernel Streaming MIDI API. */
+    RTMIDI_DUMMY    /*!< A compilable but non-functional API. */
+  };
+
+  //! A static function to determine the available compiled MIDI APIs.
+  /*!
+    The values returned in the std::vector can be compared against
+    the enumerated list values.  Note that there can be more than one
+    API compiled for certain operating systems.
+  */
+  static void getCompiledApi( std::vector<RtMidi::Api> &apis );
+
+  //! Pure virtual openPort() function.
+  virtual void openPort( unsigned int portNumber = 0, const std::string portName = std::string( "RtMidi" ) ) = 0;
+
+  //! Pure virtual openVirtualPort() function.
+  virtual void openVirtualPort( const std::string portName = std::string( "RtMidi" ) ) = 0;
+
+  //! Pure virtual getPortCount() function.
+  virtual unsigned int getPortCount() = 0;
+
+  //! Pure virtual getPortName() function.
+  virtual std::string getPortName( unsigned int portNumber = 0 ) = 0;
+
+  //! Pure virtual closePort() function.
+  virtual void closePort( void ) = 0;
+
+  //! A basic error reporting function for RtMidi classes.
+  static void error( RtError::Type type, std::string errorString );
+
+ protected:
+
+  RtMidi() {};
+  virtual ~RtMidi() {};
+};
+
+/**********************************************************************/
+/*! \class RtMidiIn
+    \brief A realtime MIDI input class.
+
+    This class provides a common, platform-independent API for
+    realtime MIDI input.  It allows access to a single MIDI input
+    port.  Incoming MIDI messages are either saved to a queue for
+    retrieval using the getMessage() function or immediately passed to
+    a user-specified callback function.  Create multiple instances of
+    this class to connect to more than one MIDI device at the same
+    time.  With the OS-X and Linux ALSA MIDI APIs, it is also possible
+    to open a virtual input port to which other MIDI software clients
+    can connect.
+
+    by Gary P. Scavone, 2003-2012.
+*/
+/**********************************************************************/
+
+// **************************************************************** //
+//
+// RtMidiIn and RtMidiOut class declarations.
+//
+// RtMidiIn / RtMidiOut are "controllers" used to select an available
+// MIDI input or output interface.  They present common APIs for the
+// user to call but all functionality is implemented by the classes
+// MidiInApi, MidiOutApi and their subclasses.  RtMidiIn and RtMidiOut
+// each create an instance of a MidiInApi or MidiOutApi subclass based
+// on the user's API choice.  If no choice is made, they attempt to
+// make a "logical" API selection.
+//
+// **************************************************************** //
+
+class MidiInApi;
+class MidiOutApi;
+
+class RtMidiIn : public RtMidi
+{
+ public:
+
+  //! User callback function type definition.
+  typedef void (*RtMidiCallback)( double timeStamp, std::vector<unsigned char> *message, void *userData);
+
+  //! Default constructor that allows an optional api, client name and queue size.
+  /*!
+    An assert will be fired if a MIDI system initialization
+    error occurs.  The queue size defines the maximum number of
+    messages that can be held in the MIDI queue (when not using a
+    callback function).  If the queue size limit is reached,
+    incoming messages will be ignored.
+
+    If no API argument is specified and multiple API support has been
+    compiled, the default order of use is JACK, ALSA (Linux) and CORE,
+    Jack (OS-X).
+  */
+  RtMidiIn( RtMidi::Api api=UNSPECIFIED,
+            const std::string clientName = std::string( "RtMidi Input Client"),
+            unsigned int queueSizeLimit = 100 );
+
+  //! If a MIDI connection is still open, it will be closed by the destructor.
+  ~RtMidiIn ( void );
+
+  //! Returns the MIDI API specifier for the current instance of RtMidiIn.
+  RtMidi::Api getCurrentApi( void );
+
+  //! Open a MIDI input connection.
+  /*!
+    An optional port number greater than 0 can be specified.
+    Otherwise, the default or first port found is opened.
+  */
+  void openPort( unsigned int portNumber = 0, const std::string portName = std::string( "RtMidi Input" ) );
+
+  //! Create a virtual input port, with optional name, to allow software connections (OS X and ALSA only).
+  /*!
+    This function creates a virtual MIDI input port to which other
+    software applications can connect.  This type of functionality
+    is currently only supported by the Macintosh OS-X and Linux ALSA
+    APIs (the function does nothing for the other APIs).
+  */
+  void openVirtualPort( const std::string portName = std::string( "RtMidi Input" ) );
+
+  //! Set a callback function to be invoked for incoming MIDI messages.
+  /*!
+    The callback function will be called whenever an incoming MIDI
+    message is received.  While not absolutely necessary, it is best
+    to set the callback function before opening a MIDI port to avoid
+    leaving some messages in the queue.
+  */
+  void setCallback( RtMidiCallback callback, void *userData = 0 );
+
+  //! Cancel use of the current callback function (if one exists).
+  /*!
+    Subsequent incoming MIDI messages will be written to the queue
+    and can be retrieved with the \e getMessage function.
+  */
+  void cancelCallback();
+
+  //! Close an open MIDI connection (if one exists).
+  void closePort( void );
+
+  //! Return the number of available MIDI input ports.
+  unsigned int getPortCount();
+
+  //! Return a string identifier for the specified MIDI input port number.
+  /*!
+    An empty string is returned if an invalid port specifier is provided.
+  */
+  std::string getPortName( unsigned int portNumber = 0 );
+
+  //! Specify whether certain MIDI message types should be queued or ignored during input.
+  /*!
+    o      By default, MIDI timing and active sensing messages are ignored
+    during message input because of their relative high data rates.
+    MIDI sysex messages are ignored by default as well.  Variable
+    values of "true" imply that the respective message type will be
+    ignored.
+  */
+  void ignoreTypes( bool midiSysex = true, bool midiTime = true, bool midiSense = true );
+
+  //! Fill the user-provided vector with the data bytes for the next available MIDI message in the input queue and return the event delta-time in seconds.
+  /*!
+    This function returns immediately whether a new message is
+    available or not.  A valid message is indicated by a non-zero
+    vector size.  An assert is fired if an error occurs during
+    message retrieval or an input connection was not previously
+    established.
+  */
+  double getMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void openMidiApi( RtMidi::Api api, const std::string clientName, unsigned int queueSizeLimit );
+  MidiInApi *rtapi_;
+
+};
+
+/**********************************************************************/
+/*! \class RtMidiOut
+    \brief A realtime MIDI output class.
+
+    This class provides a common, platform-independent API for MIDI
+    output.  It allows one to probe available MIDI output ports, to
+    connect to one such port, and to send MIDI bytes immediately over
+    the connection.  Create multiple instances of this class to
+    connect to more than one MIDI device at the same time.  With the
+    OS-X and Linux ALSA MIDI APIs, it is also possible to open a
+    virtual port to which other MIDI software clients can connect.
+
+    by Gary P. Scavone, 2003-2012.
+*/
+/**********************************************************************/
+
+class RtMidiOut : public RtMidi
+{
+ public:
+
+  //! Default constructor that allows an optional client name.
+  /*!
+    An exception will be thrown if a MIDI system initialization error occurs.
+
+    If no API argument is specified and multiple API support has been
+    compiled, the default order of use is JACK, ALSA (Linux) and CORE,
+    Jack (OS-X).
+  */
+  RtMidiOut( RtMidi::Api api=UNSPECIFIED,
+             const std::string clientName = std::string( "RtMidi Output Client") );
+
+  //! The destructor closes any open MIDI connections.
+  ~RtMidiOut( void );
+
+  //! Returns the MIDI API specifier for the current instance of RtMidiOut.
+  RtMidi::Api getCurrentApi( void );
+
+  //! Open a MIDI output connection.
+  /*!
+      An optional port number greater than 0 can be specified.
+      Otherwise, the default or first port found is opened.  An
+      exception is thrown if an error occurs while attempting to make
+      the port connection.
+  */
+  void openPort( unsigned int portNumber = 0, const std::string portName = std::string( "RtMidi Output" ) );
+
+  //! Close an open MIDI connection (if one exists).
+  void closePort( void );
+
+  //! Create a virtual output port, with optional name, to allow software connections (OS X and ALSA only).
+  /*!
+      This function creates a virtual MIDI output port to which other
+      software applications can connect.  This type of functionality
+      is currently only supported by the Macintosh OS-X and Linux ALSA
+      APIs (the function does nothing with the other APIs).  An
+      exception is thrown if an error occurs while attempting to create
+      the virtual port.
+  */
+  void openVirtualPort( const std::string portName = std::string( "RtMidi Output" ) );
+
+  //! Return the number of available MIDI output ports.
+  unsigned int getPortCount( void );
+
+  //! Return a string identifier for the specified MIDI port type and number.
+  /*!
+      An empty string is returned if an invalid port specifier is provided.
+  */
+  std::string getPortName( unsigned int portNumber = 0 );
+
+  //! Immediately send a single message out an open MIDI output port.
+  /*!
+      An exception is thrown if an error occurs during output or an
+      output connection was not previously established.
+  */
+  void sendMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void openMidiApi( RtMidi::Api api, const std::string clientName );
+  MidiOutApi *rtapi_;
+};
+
+
+// **************************************************************** //
+//
+// MidiInApi / MidiOutApi class declarations.
+//
+// Subclasses of MidiInApi and MidiOutApi contain all API- and
+// OS-specific code necessary to fully implement the RtMidi API.
+//
+// Note that MidiInApi and MidiOutApi are abstract base classes and
+// cannot be explicitly instantiated.  RtMidiIn and RtMidiOut will
+// create instances of a MidiInApi or MidiOutApi subclass.
+//
+// **************************************************************** //
+
+class MidiInApi
+{
+ public:
+
+  MidiInApi( unsigned int queueSizeLimit );
+  virtual ~MidiInApi( void );
+  virtual RtMidi::Api getCurrentApi( void ) = 0;
+  virtual void openPort( unsigned int portNumber, const std::string portName ) = 0;
+  virtual void openVirtualPort( const std::string portName ) = 0;
+  virtual void closePort( void ) = 0;
+  void setCallback( RtMidiIn::RtMidiCallback callback, void *userData );
+  void cancelCallback( void );
+  virtual unsigned int getPortCount( void ) = 0;
+  virtual std::string getPortName( unsigned int portNumber ) = 0;
+  virtual void ignoreTypes( bool midiSysex, bool midiTime, bool midiSense );
+  double getMessage( std::vector<unsigned char> *message );
+
+  // A MIDI structure used internally by the class to store incoming
+  // messages.  Each message represents one and only one MIDI message.
+  struct MidiMessage { 
+    std::vector<unsigned char> bytes; 
+    double timeStamp;
+
+    // Default constructor.
+  MidiMessage()
+  :bytes(0), timeStamp(0.0) {}
+  };
+
+  struct MidiQueue {
+    unsigned int front;
+    unsigned int back;
+    unsigned int size;
+    unsigned int ringSize;
+    MidiMessage *ring;
+
+    // Default constructor.
+  MidiQueue()
+  :front(0), back(0), size(0), ringSize(0) {}
+  };
+
+  // The RtMidiInData structure is used to pass private class data to
+  // the MIDI input handling function or thread.
+  struct RtMidiInData {
+    MidiQueue queue;
+    MidiMessage message;
+    unsigned char ignoreFlags;
+    bool doInput;
+    bool firstMessage;
+    void *apiData;
+    bool usingCallback;
+    void *userCallback;
+    void *userData;
+    bool continueSysex;
+
+    // Default constructor.
+  RtMidiInData()
+  : ignoreFlags(7), doInput(false), firstMessage(true),
+      apiData(0), usingCallback(false), userCallback(0), userData(0),
+      continueSysex(false) {}
+  };
+
+ protected:
+  virtual void initialize( const std::string& clientName ) = 0;
+  RtMidiInData inputData_;
+
+  void *apiData_;
+  bool connected_;
+  std::string errorString_;
+};
+
+class MidiOutApi
+{
+ public:
+
+  MidiOutApi( void );
+  virtual ~MidiOutApi( void );
+  virtual RtMidi::Api getCurrentApi( void ) = 0;
+  virtual void openPort( unsigned int portNumber, const std::string portName ) = 0;
+  virtual void openVirtualPort( const std::string portName ) = 0;
+  virtual void closePort( void ) = 0;
+  virtual unsigned int getPortCount( void ) = 0;
+  virtual std::string getPortName( unsigned int portNumber ) = 0;
+  virtual void sendMessage( std::vector<unsigned char> *message ) = 0;
+
+ protected:
+  virtual void initialize( const std::string& clientName ) = 0;
+
+  void *apiData_;
+  bool connected_;
+  std::string errorString_;
+};
+
+// **************************************************************** //
+//
+// Inline RtMidiIn and RtMidiOut definitions.
+//
+// **************************************************************** //
+
+inline RtMidi::Api RtMidiIn :: getCurrentApi( void ) { return rtapi_->getCurrentApi(); }
+inline void RtMidiIn :: openPort( unsigned int portNumber, const std::string portName ) { return rtapi_->openPort( portNumber, portName ); }
+inline void RtMidiIn :: openVirtualPort( const std::string portName ) { return rtapi_->openVirtualPort( portName ); }
+inline void RtMidiIn :: closePort( void ) { return rtapi_->closePort(); }
+inline void RtMidiIn :: setCallback( RtMidiCallback callback, void *userData ) { return rtapi_->setCallback( callback, userData ); }
+inline void RtMidiIn :: cancelCallback( void ) { return rtapi_->cancelCallback(); }
+inline unsigned int RtMidiIn :: getPortCount( void ) { return rtapi_->getPortCount(); }
+inline std::string RtMidiIn :: getPortName( unsigned int portNumber ) { return rtapi_->getPortName( portNumber ); }
+inline void RtMidiIn :: ignoreTypes( bool midiSysex, bool midiTime, bool midiSense ) { return rtapi_->ignoreTypes( midiSysex, midiTime, midiSense ); }
+inline double RtMidiIn :: getMessage( std::vector<unsigned char> *message ) { return rtapi_->getMessage( message ); }
+
+inline RtMidi::Api RtMidiOut :: getCurrentApi( void )  { return rtapi_->getCurrentApi(); }
+inline void RtMidiOut :: openPort( unsigned int portNumber, const std::string portName ) { return rtapi_->openPort( portNumber, portName ); }
+inline void RtMidiOut :: openVirtualPort( const std::string portName ) { return rtapi_->openVirtualPort( portName ); }
+inline void RtMidiOut :: closePort( void ) { return rtapi_->closePort(); }
+inline unsigned int RtMidiOut :: getPortCount( void ) { return rtapi_->getPortCount(); }
+inline std::string RtMidiOut :: getPortName( unsigned int portNumber ) { return rtapi_->getPortName( portNumber ); }
+inline void RtMidiOut :: sendMessage( std::vector<unsigned char> *message ) { return rtapi_->sendMessage( message ); }
+
+// **************************************************************** //
+//
+// MidiInApi and MidiOutApi subclass prototypes.
+//
+// **************************************************************** //
+
+#if !defined(__LINUX_ALSA__) && !defined(__UNIX_JACK__) && !defined(__MACOSX_CORE__) && !defined(__WINDOWS_MM__) && !defined(__WINDOWS_KS__)
+  #define __RTMIDI_DUMMY__
+#endif
+
+#if defined(__MACOSX_CORE__)
+
+class MidiInCore: public MidiInApi
+{
+ public:
+  MidiInCore( const std::string clientName, unsigned int queueSizeLimit );
+  ~MidiInCore( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::MACOSX_CORE; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+class MidiOutCore: public MidiOutApi
+{
+ public:
+  MidiOutCore( const std::string clientName );
+  ~MidiOutCore( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::MACOSX_CORE; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+  void sendMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+#endif
+
+#if defined(__UNIX_JACK__)
+
+class MidiInJack: public MidiInApi
+{
+ public:
+  MidiInJack( const std::string clientName, unsigned int queueSizeLimit );
+  ~MidiInJack( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::UNIX_JACK; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+class MidiOutJack: public MidiOutApi
+{
+ public:
+  MidiOutJack( const std::string clientName );
+  ~MidiOutJack( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::UNIX_JACK; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+  void sendMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+#endif
+
+#if defined(__LINUX_ALSA__)
+
+class MidiInAlsa: public MidiInApi
+{
+ public:
+  MidiInAlsa( const std::string clientName, unsigned int queueSizeLimit );
+  ~MidiInAlsa( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::LINUX_ALSA; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+class MidiOutAlsa: public MidiOutApi
+{
+ public:
+  MidiOutAlsa( const std::string clientName );
+  ~MidiOutAlsa( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::LINUX_ALSA; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+  void sendMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+#endif
+
+#if defined(__WINDOWS_MM__)
+
+class MidiInWinMM: public MidiInApi
+{
+ public:
+  MidiInWinMM( const std::string clientName, unsigned int queueSizeLimit );
+  ~MidiInWinMM( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::WINDOWS_MM; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+class MidiOutWinMM: public MidiOutApi
+{
+ public:
+  MidiOutWinMM( const std::string clientName );
+  ~MidiOutWinMM( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::WINDOWS_MM; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+  void sendMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+#endif
+
+#if defined(__WINDOWS_KS__)
+
+class MidiInWinKS: public MidiInApi
+{
+ public:
+  MidiInWinKS( const std::string clientName, unsigned int queueSizeLimit );
+  ~MidiInWinKS( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::WINDOWS_KS; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+class MidiOutWinKS: public MidiOutApi
+{
+ public:
+  MidiOutWinKS( const std::string clientName );
+  ~MidiOutWinKS( void );
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::WINDOWS_KS; };
+  void openPort( unsigned int portNumber, const std::string portName );
+  void openVirtualPort( const std::string portName );
+  void closePort( void );
+  unsigned int getPortCount( void );
+  std::string getPortName( unsigned int portNumber );
+  void sendMessage( std::vector<unsigned char> *message );
+
+ protected:
+  void initialize( const std::string& clientName );
+};
+
+#endif
+
+#if defined(__RTMIDI_DUMMY__)
+aa
+class MidiInDummy: public MidiInApi
+{
+ public:
+ MidiInDummy( const std::string clientName, unsigned int queueSizeLimit ) : MidiInApi( queueSizeLimit ) { errorString_ = "MidiInDummy: This class provides no functionality."; RtMidi::error( RtError::WARNING, errorString_ ); };
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::RTMIDI_DUMMY; };
+  void openPort( unsigned int portNumber, const std::string portName ) {};
+  void openVirtualPort( const std::string portName ) {};
+  void closePort( void ) {};
+  unsigned int getPortCount( void ) { return 0; };
+  std::string getPortName( unsigned int portNumber ) { return ""; };
+
+ protected:
+  void initialize( const std::string& clientName ) {};
+};
+
+class MidiOutDummy: public MidiOutApi
+{
+ public:
+  MidiOutDummy( const std::string clientName ) { errorString_ = "MidiOutDummy: This class provides no functionality."; RtMidi::error( RtError::WARNING, errorString_ ); };
+  RtMidi::Api getCurrentApi( void ) { return RtMidi::RTMIDI_DUMMY; };
+  void openPort( unsigned int portNumber, const std::string portName ) {};
+  void openVirtualPort( const std::string portName ) {};
+  void closePort( void ) {};
+  unsigned int getPortCount( void ) { return 0; };
+  std::string getPortName( unsigned int portNumber ) { return ""; };
+  void sendMessage( std::vector<unsigned char> *message ) {};
+
+ protected:
+  void initialize( const std::string& clientName ) {};
+};
+
+#endif
+
+#endif

btgui/MidiTest/cmidiin.cpp

@@ -0,0 +1,111 @@
+//*****************************************//
+//  cmidiin.cpp
+//  by Gary Scavone, 2003-2004.
+//
+//  Simple program to test MIDI input and
+//  use of a user callback function.
+//
+//*****************************************//
+
+#include <iostream>
+#include <cstdlib>
+#include "RtMidi.h"
+
+void usage( void ) {
+  // Error function in case of incorrect command-line
+  // argument specifications.
+  std::cout << "\nuseage: cmidiin <port>\n";
+  std::cout << "    where port = the device to use (default = 0).\n\n";
+  exit( 0 );
+}
+
+void mycallback( double deltatime, std::vector< unsigned char > *message, void *userData )
+{
+  unsigned int nBytes = message->size();
+  for ( unsigned int i=0; i<nBytes; i++ )
+    std::cout << "Byte " << i << " = " << (int)message->at(i) << ", ";
+  if ( nBytes > 0 )
+    std::cout << "stamp = " << deltatime << std::endl;
+}
+
+// This function should be embedded in a try/catch block in case of
+// an exception.  It offers the user a choice of MIDI ports to open.
+// It returns false if there are no ports available.
+bool chooseMidiPort( RtMidiIn *rtmidi );
+
+int main( int argc, char *argv[] )
+{
+  RtMidiIn *midiin = 0;
+
+  // Minimal command-line check.
+  if ( argc > 2 ) usage();
+
+
+    // RtMidiIn constructor
+    midiin = new RtMidiIn();
+
+    // Call function to select port.
+    if ( chooseMidiPort( midiin ) == false ) goto cleanup;
+
+    // Set our callback function.  This should be done immediately after
+    // opening the port to avoid having incoming messages written to the
+    // queue instead of sent to the callback function.
+    midiin->setCallback( &mycallback );
+
+    // Don't ignore sysex, timing, or active sensing messages.
+    midiin->ignoreTypes( false, false, false );
+
+    std::cout << "\nReading MIDI input ... press <enter> to quit.\n";
+    char input;
+    std::cin.get(input);
+	getchar();
+
+
+ cleanup:
+
+  delete midiin;
+
+  return 0;
+}
+
+bool chooseMidiPort( RtMidiIn *rtmidi )
+{
+	  /*
+
+	std::cout << "\nWould you like to open a virtual input port? [y/N] ";
+
+  std::string keyHit;
+  std::getline( std::cin, keyHit );
+  if ( keyHit == "y" ) {
+    rtmidi->openVirtualPort();
+    return true;
+  }
+  */
+
+  std::string portName;
+  unsigned int i = 0, nPorts = rtmidi->getPortCount();
+  if ( nPorts == 0 ) {
+    std::cout << "No input ports available!" << std::endl;
+    return false;
+  }
+
+  if ( nPorts == 1 ) {
+    std::cout << "\nOpening " << rtmidi->getPortName() << std::endl;
+  }
+  else {
+    for ( i=0; i<nPorts; i++ ) {
+      portName = rtmidi->getPortName(i);
+      std::cout << "  Input port #" << i << ": " << portName << '\n';
+    }
+
+    do {
+      std::cout << "\nChoose a port number: ";
+      std::cin >> i;
+    } while ( i >= nPorts );
+  }
+
+//  std::getline( std::cin, keyHit );  // used to clear out stdin
+  rtmidi->openPort( i );
+
+  return true;
+}

btgui/MidiTest/premake4.lua

@@ -0,0 +1,35 @@
+
+	project "rtMidiTest"
+		
+	kind "ConsoleApp"
+	
+--	defines {  }
+	
+	targetdir "../../bin"
+	
+	includedirs 
+	{
+		".",
+	}
+
+			
+--	links { }
+	
+	
+	files {
+		"**.cpp",
+		"**.h"
+	}
+	if os.is("Windows") then
+		links {"winmm"}
+		defines {"__WINDOWS_MM__", "WIN32"}
+	end
+
+	if os.is("Linux") then 
+	end
+
+	if os.is("MacOSX") then
+		links{"CoreAudio.framework", "coreMIDI.framework", "Cocoa.framework"}
+		defines {"__MACOSX_CORE__"}
+		print ("hi!")
+	end

btgui/OpenGLWindow/GLInstancingRenderer.cpp

@@ -209,7 +209,6 @@ void btDefaultMouseMoveCallback( float x, float y)
 
 void btDefaultKeyboardCallback(int key, int state)
 {
-	//printf("world2\n");
 }
 
 

build/premake4.lua

@@ -91,22 +91,23 @@
 
 
 	if not _OPTIONS["ios"] then
-		include "../opencl/vector_add_simplified"
+		include "../demo/gpudemo"
-		include "../opencl/vector_add"
+	include "../btgui/MidiTest"
-		include "../opencl/basic_initialize"
+--		include "../opencl/vector_add_simplified"
-		include "../opencl/parallel_primitives/host"
+--		include "../opencl/vector_add"
-		include "../opencl/parallel_primitives/test"
+--		include "../opencl/basic_initialize"
-		include "../opencl/parallel_primitives/benchmark"
+--		include "../opencl/parallel_primitives/host"
-		include "../opencl/lds_bank_conflict"
+--		include "../opencl/parallel_primitives/test"
-		include "../opencl/reduce"
+--		include "../opencl/parallel_primitives/benchmark"
-		include "../opencl/gpu_broadphase/test"
+--		include "../opencl/lds_bank_conflict"
-		include "../opencl/gpu_sat/test"
+--		include "../opencl/reduce"
+--		include "../opencl/gpu_broadphase/test"
+--		include "../opencl/gpu_sat/test"
 		include "../btgui/Gwen"
 		include "../btgui/GwenOpenGLTest"
-		include "../btgui/OpenGLTrueTypeFont"
+--		include "../btgui/OpenGLTrueTypeFont"
-		include "../btgui/OpenGLWindow"
+--		include "../btgui/OpenGLWindow"
-		include "../demo/gpudemo"
+--		include "../demo/ObjLoader"
-		include "../demo/ObjLoader"
 		
 		
 	end

build/stringify.bat

@@ -13,6 +13,8 @@ premake4 --file=stringifyKernel.lua --kernelfile="../opencl/gpu_broadphase/kerne
 
 premake4 --file=stringifyKernel.lua --kernelfile="../opencl/gpu_sat/kernels/sat.cl" --headerfile="../opencl/gpu_sat/kernels/satKernels.h" --stringname="satKernelsCL" stringify
 premake4 --file=stringifyKernel.lua --kernelfile="../opencl/gpu_sat/kernels/satClipHullContacts.cl" --headerfile="../opencl/gpu_sat/kernels/satClipHullContacts.h" --stringname="satClipKernelsCL" stringify
+premake4 --file=stringifyKernel.lua --kernelfile="../opencl/gpu_sat/kernels/primitiveContacts.cl" --headerfile="../opencl/gpu_sat/kernels/primitiveContacts.h" --stringname="primitiveContactsKernelsCL" stringify
+
 premake4 --file=stringifyKernel.lua --kernelfile="../opencl/gpu_sat/kernels/bvhTraversal.cl" --headerfile="../opencl/gpu_sat/kernels/bvhTraversal.h" --stringname="bvhTraversalKernelCL" stringify
 
 

demo/gpudemo/GpuDemo.h

@@ -32,20 +32,22 @@ public:
             float gapZ;
             GLInstancingRenderer*   m_instancingRenderer;
 			class btgWindowInterface*	m_window;
+			class GwenUserInterface*	m_gui;
 
             ConstructionInfo()
                     :useOpenCL(true),
                     preferredOpenCLPlatformIndex(-1),
                     preferredOpenCLDeviceIndex(-1),
-					arraySizeX(10),
+					arraySizeX(25),
-		arraySizeY(20),
+		arraySizeY(23),
-		arraySizeZ(10),
+		arraySizeZ(23),
 		m_useConcaveMesh(false),
 		gapX(14.3),
 		gapY(14.0),
 		gapZ(14.3),
                     m_instancingRenderer(0),
-					m_window(0)
+					m_window(0),
+					m_gui(0)
             {
             }
     };

demo/gpudemo/main_opengl3core.cpp

@@ -29,6 +29,7 @@
 #include "rigidbody/ConcaveScene.h"
 #include "rigidbody/GpuConvexScene.h"
 #include "rigidbody/GpuCompoundScene.h"
+#include "rigidbody/GpuSphereScene.h"
 
 //#include "BroadphaseBenchmark.h"
 
@@ -64,25 +65,25 @@ btAlignedObjectArray<const char*> demoNames;
 int selectedDemo = 0;
 GpuDemo::CreateFunc* allDemos[]=
 {
+	GpuSphereScene::MyCreateFunc,
 	GpuConvexScene::MyCreateFunc,
 	ConcaveScene::MyCreateFunc,
 
-	GpuConvexScene::MyCreateFunc,
 	GpuCompoundScene::MyCreateFunc,
+	PairBench::MyCreateFunc,	
 
 
-	GpuRigidBodyDemo::MyCreateFunc,
+	//GpuRigidBodyDemo::MyCreateFunc,
 
 	//BroadphaseBenchmark::CreateFunc,
 	//GpuBoxDemo::CreateFunc,
 
-	PairBench::MyCreateFunc,	
 
 
-	ParticleDemo::MyCreateFunc,
+	//ParticleDemo::MyCreateFunc,
+	
 	
 	
-	//SpheresDemo::CreateFunc,
 	//GpuCompoundDemo::CreateFunc,
 	//EmptyDemo::CreateFunc,
 };
@@ -381,7 +382,11 @@ int main(int argc, char* argv[])
 	
 	args.GetCmdLineArgument("selected_demo",selectedDemo);
 
-	useNewBatchingKernel = args.CheckCmdLineFlag("new_batching");
+		
+	if (args.CheckCmdLineFlag("new_batching"))
+	{
+		useNewBatchingKernel = true;
+	}
 	bool benchmark=args.CheckCmdLineFlag("benchmark");
 	dump_timings=args.CheckCmdLineFlag("dump_timings");
 	ci.useOpenCL = !args.CheckCmdLineFlag("disable_opencl");
@@ -563,6 +568,7 @@ int main(int argc, char* argv[])
 
 		ci.m_instancingRenderer = new GLInstancingRenderer(maxObjectCapacity);//render.getInstancingRenderer();
 		ci.m_window = window;
+		ci.m_gui = gui;
 		ci.m_instancingRenderer->init();
 		ci.m_instancingRenderer->InitShaders();
 		
@@ -615,11 +621,6 @@ int main(int argc, char* argv[])
 
 			window->startRendering();
 
-			char msg[1024];
-			int numInstances = 0;//ci.m_instancingRenderer->getNumInstances();
-			sprintf(msg,"Num objects = %d",numInstances);
-			gui->setStatusBarMessage(msg,true);
-
 			glClearColor(0.6,0.6,0.6,1);
 			glClear(GL_COLOR_BUFFER_BIT| GL_DEPTH_BUFFER_BIT|GL_STENCIL_BUFFER_BIT);
 			glEnable(GL_DEPTH_TEST);

demo/gpudemo/rigidbody/GpuConvexScene.cpp

@@ -16,6 +16,8 @@
 #include "gpu_rigidbody/host/btGpuNarrowPhase.h"
 #include "gpu_rigidbody/host/btConfig.h"
 #include "GpuRigidBodyDemoInternalData.h"
+#include "../gwenUserInterface.h"
+
 
 void GpuConvexScene::setupScene(const ConstructionInfo& ci)
 {
@@ -86,4 +88,9 @@ void GpuConvexScene::setupScene(const ConstructionInfo& ci)
 	m_instancingRenderer->setCameraTargetPosition(camPos);
 	m_instancingRenderer->setCameraDistance(120);
 
+
+	char msg[1024];
+	int numInstances = index;
+	sprintf(msg,"Num objects = %d",numInstances);
+	ci.m_gui->setStatusBarMessage(msg,true);
 }

demo/gpudemo/rigidbody/GpuSphereScene.cpp

@@ -0,0 +1,151 @@
+#include "GpuSphereScene.h"
+#include "GpuRigidBodyDemo.h"
+#include "BulletCommon/btQuickprof.h"
+#include "OpenGLWindow/ShapeData.h"
+
+#include "OpenGLWindow/GLInstancingRenderer.h"
+#include "BulletCommon/btQuaternion.h"
+#include "OpenGLWindow/btgWindowInterface.h"
+#include "gpu_broadphase/host/btGpuSapBroadphase.h"
+#include "../GpuDemoInternalData.h"
+#include "basic_initialize/btOpenCLUtils.h"
+#include "OpenGLWindow/OpenGLInclude.h"
+#include "OpenGLWindow/GLInstanceRendererInternalData.h"
+#include "parallel_primitives/host/btLauncherCL.h"
+#include "gpu_rigidbody/host/btGpuRigidBodyPipeline.h"
+#include "gpu_rigidbody/host/btGpuNarrowPhase.h"
+#include "gpu_rigidbody/host/btConfig.h"
+#include "GpuRigidBodyDemoInternalData.h"
+#include "../gwenUserInterface.h"
+
+
+
+
+void GpuSphereScene::setupScene(const ConstructionInfo& ci)
+{
+	int strideInBytes = 9*sizeof(float);
+	int numVertices = sizeof(cube_vertices)/strideInBytes;
+	int numIndices = sizeof(cube_indices)/sizeof(int);
+	//int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+	
+	int group=1;
+	int mask=1;
+	int index=0;
+	
+	if (1)
+	{
+		int shapeId = ci.m_instancingRenderer->registerShape(&cube_vertices[0],numVertices,cube_indices,numIndices);
+		btVector4 scaling(120,2,120,1);
+		int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
+		btVector3 normal(0,-1,0);
+		float constant=2;
+		
+		//int colIndex = m_data->m_np->registerPlaneShape(normal,constant);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
+		btVector4 position(0,50,0,0);
+		//btQuaternion orn(0,0,0,1);
+		btQuaternion orn(btVector3(1,0,0),0.3);
+
+		btVector4 color(0,0,1,1);
+		
+		int id = ci.m_instancingRenderer->registerGraphicsInstance(shapeId,position,orn,color,scaling);
+		int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(0.f,position,orn,colIndex,index);
+				
+		index++;
+	}
+
+	
+
+	{
+
+
+
+		int prevGraphicsShapeIndex = -1;
+		float radius = 1;
+		if (radius>=100)
+		{
+			int numVertices = sizeof(detailed_sphere_vertices)/strideInBytes;
+			int numIndices = sizeof(detailed_sphere_indices)/sizeof(int);
+			prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&detailed_sphere_vertices[0],numVertices,detailed_sphere_indices,numIndices);
+		} else
+		{
+			bool usePointSprites = true;
+			if (usePointSprites)
+			{
+				int numVertices = sizeof(point_sphere_vertices)/strideInBytes;
+				int numIndices = sizeof(point_sphere_indices)/sizeof(int);
+				prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&point_sphere_vertices[0],numVertices,point_sphere_indices,numIndices,BT_GL_POINTS);
+			} else
+			{
+				if (radius>=10)
+				{
+					int numVertices = sizeof(medium_sphere_vertices)/strideInBytes;
+					int numIndices = sizeof(medium_sphere_indices)/sizeof(int);
+					prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&medium_sphere_vertices[0],numVertices,medium_sphere_indices,numIndices);
+				} else
+				{
+					int numVertices = sizeof(low_sphere_vertices)/strideInBytes;
+					int numIndices = sizeof(low_sphere_indices)/sizeof(int);
+					prevGraphicsShapeIndex = ci.m_instancingRenderer->registerShape(&low_sphere_vertices[0],numVertices,low_sphere_indices,numIndices);
+				}
+			}
+		}
+
+
+
+
+		btVector4 colors[4] = 
+	{
+		btVector4(1,0,0,1),
+		btVector4(0,1,0,1),
+		btVector4(0,1,1,1),
+		btVector4(1,1,0,1),
+	};
+		
+
+
+
+
+
+
+
+
+		int curColor = 0;
+		
+		//int colIndex = m_data->m_np->registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
+		int colIndex = m_data->m_np->registerSphereShape(radius);//>registerConvexHullShape(&cube_vertices[0],strideInBytes,numVertices, scaling);
+		for (int i=0;i<ci.arraySizeX;i++)
+		{
+			for (int j=0;j<ci.arraySizeY;j++)
+			{
+				for (int k=0;k<ci.arraySizeZ;k++)
+				{
+					float mass = 1.f;
+
+					//btVector3 position((j&1)+i*2.2,1+j*2.,(j&1)+k*2.2);
+					btVector3 position(i*radius*3,j*radius*3,k*radius*3);
+					
+					btQuaternion orn(0,0,0,1);
+				
+					btVector4 color = colors[curColor];
+					curColor++;
+					curColor&=3;
+					btVector4 scaling(radius,radius,radius,1);
+					int id = ci.m_instancingRenderer->registerGraphicsInstance(prevGraphicsShapeIndex,position,orn,color,scaling);
+					int pid = m_data->m_rigidBodyPipeline->registerPhysicsInstance(mass,position,orn,colIndex,index);
+				
+					index++;
+				}
+			}
+		}
+	}
+	float camPos[4]={ci.arraySizeX,ci.arraySizeY/2,ci.arraySizeZ,0};
+	//float camPos[4]={1,12.5,1.5,0};
+	m_instancingRenderer->setCameraTargetPosition(camPos);
+	m_instancingRenderer->setCameraDistance(150);
+
+
+	char msg[1024];
+	int numInstances = index;
+	sprintf(msg,"Num objects = %d",numInstances);
+	ci.m_gui->setStatusBarMessage(msg,true);
+}

demo/gpudemo/rigidbody/GpuSphereScene.h

@@ -0,0 +1,27 @@
+#ifndef GPU_SPHERE_SCENE_H
+#define GPU_SPHERE_SCENE_H
+
+#include "GpuRigidBodyDemo.h"
+
+class GpuSphereScene : public GpuRigidBodyDemo
+{
+public:
+
+	GpuSphereScene(){}
+	virtual ~GpuSphereScene(){}
+	virtual const char* getName()
+	{
+		return "GRBSphere";
+	}
+
+	static GpuDemo* MyCreateFunc()
+	{
+		GpuDemo* demo = new GpuSphereScene;
+		return demo;
+	}
+
+	virtual void setupScene(const ConstructionInfo& ci);
+
+};
+
+#endif //GPU_SPHERE_SCENE_H

opencl/basic_initialize/btOpenCLUtils.cpp

@@ -549,6 +549,7 @@ cl_program btOpenCLUtils_compileCLProgramFromString(cl_context clContext, cl_dev
 		strippedName = strip2(clFileNameForCaching,"\\");
 		strippedName = strip2(strippedName,"/");
 
+
 #ifdef _WIN32
 		sprintf_s(binaryFileName,BT_MAX_STRING_LENGTH,"cache/%s.%s.%s.bin",strippedName, deviceName,driverVersion );
 #else

opencl/gpu_rigidbody/host/Solver.cpp

@@ -17,7 +17,7 @@ subject to the following restrictions:
 #include "Solver.h"
 
 ///useNewBatchingKernel  is a rewritten kernel using just a single thread of the warp, for experiments
-bool useNewBatchingKernel = false;
+bool useNewBatchingKernel = true;
 
 #define SOLVER_SETUP_KERNEL_PATH "opencl/gpu_rigidbody/kernels/solverSetup.cl"
 #define SOLVER_SETUP2_KERNEL_PATH "opencl/gpu_rigidbody/kernels/solverSetup2.cl"

opencl/gpu_rigidbody/host/btConfig.h

@@ -18,7 +18,7 @@ struct	btConfig
 	int m_maxTriConvexPairCapacity;
 
 	btConfig()
-		:m_maxConvexBodies(128*1024),
+		:m_maxConvexBodies(32*1024),
 		m_maxConvexShapes(8192),
 		m_maxVerticesPerFace(64),
 		m_maxFacesPerShape(64),

opencl/gpu_rigidbody/host/btGpuBatchingPgsSolver.cpp

@@ -37,8 +37,8 @@ enum
 };
 
 
-bool gpuBatchContacts = true;
+bool gpuBatchContacts = true;//true;
-bool gpuSolveConstraint = true;
+bool gpuSolveConstraint = true;//true;
 
 
 struct	btGpuBatchingPgsSolverInternalData

opencl/gpu_rigidbody/host/btGpuJacobiSolver.cpp

@@ -66,7 +66,7 @@ btGpuJacobiSolver::btGpuJacobiSolver(cl_context ctx, cl_device_id device, cl_com
 	const char* additionalMacros="";
 	const char* solverUtilsSource = solverUtilsCL;
 	{
-		cl_program solverUtilsProg= btOpenCLUtils::compileCLProgramFromString( ctx, device, 0, &pErrNum,additionalMacros, SOLVER_UTILS_KERNEL_PATH,true);
+		cl_program solverUtilsProg= btOpenCLUtils::compileCLProgramFromString( ctx, device, solverUtilsSource, &pErrNum,additionalMacros, SOLVER_UTILS_KERNEL_PATH);
 		btAssert(solverUtilsProg);
 		m_data->m_countBodiesKernel =  btOpenCLUtils::compileCLKernelFromString( ctx, device, solverUtilsSource, "CountBodiesKernel", &pErrNum, solverUtilsProg,additionalMacros );
 		btAssert(m_data->m_countBodiesKernel);
@@ -921,4 +921,438 @@ void  btGpuJacobiSolver::solveGroup(btOpenCLArray<btRigidBodyCL>* bodies,btOpenC
 
 
 
+}
+
+
+void  btGpuJacobiSolver::solveGroupMixed(btOpenCLArray<btRigidBodyCL>* bodiesGPU,btOpenCLArray<btInertiaCL>* inertiasGPU,btOpenCLArray<btContact4>* manifoldPtrGPU,const btJacobiSolverInfo& solverInfo)
+{
+
+	btAlignedObjectArray<btRigidBodyCL> bodiesCPU;
+	bodiesGPU->copyToHost(bodiesCPU);
+	btAlignedObjectArray<btInertiaCL> inertiasCPU;
+	inertiasGPU->copyToHost(inertiasCPU);
+	btAlignedObjectArray<btContact4> manifoldPtrCPU;
+	manifoldPtrGPU->copyToHost(manifoldPtrCPU);
+	
+	int numBodiesCPU = bodiesGPU->size();
+	int numManifoldsCPU = manifoldPtrGPU->size();
+	BT_PROFILE("btGpuJacobiSolver::solveGroupMixed");
+
+	btAlignedObjectArray<unsigned int> bodyCount;
+	bodyCount.resize(numBodiesCPU);
+	for (int i=0;i<numBodiesCPU;i++)
+		bodyCount[i] = 0;
+
+	btAlignedObjectArray<btInt2> contactConstraintOffsets;
+	contactConstraintOffsets.resize(numManifoldsCPU);
+
+
+	for (int i=0;i<numManifoldsCPU;i++)
+	{
+		int pa = manifoldPtrCPU[i].m_bodyAPtrAndSignBit;
+		int pb = manifoldPtrCPU[i].m_bodyBPtrAndSignBit;
+
+		bool isFixedA = (pa <0) || (pa == solverInfo.m_fixedBodyIndex);
+		bool isFixedB = (pb <0) || (pb == solverInfo.m_fixedBodyIndex);
+
+		int bodyIndexA = manifoldPtrCPU[i].getBodyA();
+		int bodyIndexB = manifoldPtrCPU[i].getBodyB();
+
+		if (!isFixedA)
+		{
+			contactConstraintOffsets[i].x = bodyCount[bodyIndexA];
+			bodyCount[bodyIndexA]++;
+		}
+		if (!isFixedB)
+		{
+			contactConstraintOffsets[i].y = bodyCount[bodyIndexB];
+			bodyCount[bodyIndexB]++;
+		} 
+	}
+
+	btAlignedObjectArray<unsigned int> offsetSplitBodies;
+	offsetSplitBodies.resize(numBodiesCPU);
+	unsigned int totalNumSplitBodiesCPU;
+	m_data->m_scan->executeHost(bodyCount,offsetSplitBodies,numBodiesCPU,&totalNumSplitBodiesCPU);
+	int numlastBody = bodyCount[numBodiesCPU-1];
+	totalNumSplitBodiesCPU += numlastBody;
+
+		int numBodies = bodiesGPU->size();
+	int numManifolds = manifoldPtrGPU->size();
+
+	m_data->m_bodyCount->resize(numBodies);
+	
+	unsigned int val=0;
+	btInt2 val2;
+	val2.x=0;
+	val2.y=0;
+
+	 {
+		BT_PROFILE("m_filler");
+		m_data->m_contactConstraintOffsets->resize(numManifolds);
+		m_data->m_filler->execute(*m_data->m_bodyCount,val,numBodies);
+		
+	
+		m_data->m_filler->execute(*m_data->m_contactConstraintOffsets,val2,numManifolds);
+	}
+
+	{
+		BT_PROFILE("m_countBodiesKernel");
+		btLauncherCL launcher(this->m_queue,m_data->m_countBodiesKernel);
+		launcher.setBuffer(manifoldPtrGPU->getBufferCL());
+		launcher.setBuffer(m_data->m_bodyCount->getBufferCL());
+		launcher.setBuffer(m_data->m_contactConstraintOffsets->getBufferCL());
+		launcher.setConst(numManifolds);
+		launcher.setConst(solverInfo.m_fixedBodyIndex);
+		launcher.launch1D(numManifolds);
+	}
+
+	unsigned int totalNumSplitBodies=0;
+	m_data->m_offsetSplitBodies->resize(numBodies);
+	m_data->m_scan->execute(*m_data->m_bodyCount,*m_data->m_offsetSplitBodies,numBodies,&totalNumSplitBodies);
+	totalNumSplitBodies+=m_data->m_bodyCount->at(numBodies-1);
+
+	if (totalNumSplitBodies != totalNumSplitBodiesCPU)
+	{
+		printf("error in totalNumSplitBodies!\n");
+	}
+
+	int numContacts = manifoldPtrGPU->size();
+	m_data->m_contactConstraints->resize(numContacts);
+
+	
+	{
+		BT_PROFILE("contactToConstraintSplitKernel");
+		btLauncherCL launcher( m_queue, m_data->m_contactToConstraintSplitKernel);
+		launcher.setBuffer(manifoldPtrGPU->getBufferCL());
+		launcher.setBuffer(bodiesGPU->getBufferCL());
+		launcher.setBuffer(inertiasGPU->getBufferCL());
+		launcher.setBuffer(m_data->m_contactConstraints->getBufferCL());
+		launcher.setBuffer(m_data->m_bodyCount->getBufferCL());
+        launcher.setConst(numContacts);
+		launcher.setConst(solverInfo.m_deltaTime);
+		launcher.setConst(solverInfo.m_positionDrift);
+		launcher.setConst(solverInfo.m_positionConstraintCoeff);
+		launcher.launch1D( numContacts, 64 );
+		clFinish(m_queue);
+	}
+
+
+
+	btAlignedObjectArray<btGpuConstraint4> contactConstraints;
+	contactConstraints.resize(numManifoldsCPU);
+
+	for (int i=0;i<numManifoldsCPU;i++)
+	{
+		ContactToConstraintKernel(&manifoldPtrCPU[0],&bodiesCPU[0],&inertiasCPU[0],&contactConstraints[0],numManifoldsCPU,
+			solverInfo.m_deltaTime,
+			solverInfo.m_positionDrift,
+			solverInfo.m_positionConstraintCoeff,
+			i, bodyCount);
+	}
+	int maxIter = solverInfo.m_numIterations;
+
+
+	btAlignedObjectArray<btVector3> deltaLinearVelocities;
+	btAlignedObjectArray<btVector3> deltaAngularVelocities;
+	deltaLinearVelocities.resize(totalNumSplitBodiesCPU);
+	deltaAngularVelocities.resize(totalNumSplitBodiesCPU);
+	for (int i=0;i<totalNumSplitBodiesCPU;i++)
+	{
+		deltaLinearVelocities[i].setZero();
+		deltaAngularVelocities[i].setZero();
+	}
+
+	m_data->m_deltaLinearVelocities->resize(totalNumSplitBodies);
+	m_data->m_deltaAngularVelocities->resize(totalNumSplitBodies);
+
+
+	
+	{
+		BT_PROFILE("m_clearVelocitiesKernel");
+		btLauncherCL launch(m_queue,m_data->m_clearVelocitiesKernel);
+		launch.setBuffer(m_data->m_deltaAngularVelocities->getBufferCL());
+		launch.setBuffer(m_data->m_deltaLinearVelocities->getBufferCL());
+		launch.setConst(totalNumSplitBodies);
+		launch.launch1D(totalNumSplitBodies);
+	}
+	
+
+		///!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
+
+
+	m_data->m_contactConstraints->copyToHost(contactConstraints);
+	m_data->m_offsetSplitBodies->copyToHost(offsetSplitBodies);
+	m_data->m_contactConstraintOffsets->copyToHost(contactConstraintOffsets);
+	m_data->m_deltaLinearVelocities->copyToHost(deltaLinearVelocities);
+	m_data->m_deltaAngularVelocities->copyToHost(deltaAngularVelocities);
+
+	for (int iter = 0;iter<maxIter;iter++)
+	{
+
+				{
+			BT_PROFILE("m_solveContactKernel");
+			btLauncherCL launcher( m_queue, m_data->m_solveContactKernel );
+			launcher.setBuffer(m_data->m_contactConstraints->getBufferCL());
+			launcher.setBuffer(bodiesGPU->getBufferCL());
+			launcher.setBuffer(inertiasGPU->getBufferCL());
+			launcher.setBuffer(m_data->m_contactConstraintOffsets->getBufferCL());
+			launcher.setBuffer(m_data->m_offsetSplitBodies->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaLinearVelocities->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaAngularVelocities->getBufferCL());
+			launcher.setConst(solverInfo.m_deltaTime);
+			launcher.setConst(solverInfo.m_positionDrift);
+			launcher.setConst(solverInfo.m_positionConstraintCoeff);
+			launcher.setConst(solverInfo.m_fixedBodyIndex);
+			launcher.setConst(numManifolds);
+
+			launcher.launch1D(numManifolds);
+			clFinish(m_queue);
+		}
+
+
+		int i=0;
+		for( i=0; i<numManifoldsCPU; i++)
+		{
+
+			float frictionCoeff = contactConstraints[i].getFrictionCoeff();
+			int aIdx = (int)contactConstraints[i].m_bodyA;
+			int bIdx = (int)contactConstraints[i].m_bodyB;
+			btRigidBodyCL& bodyA = bodiesCPU[aIdx];
+			btRigidBodyCL& bodyB = bodiesCPU[bIdx];
+
+			btVector3 zero(0,0,0);
+			
+			btVector3* dlvAPtr=&zero;
+			btVector3* davAPtr=&zero;
+			btVector3* dlvBPtr=&zero;
+			btVector3* davBPtr=&zero;
+			
+			if (bodyA.getInvMass())
+			{
+				int bodyOffsetA = offsetSplitBodies[aIdx];
+				int constraintOffsetA = contactConstraintOffsets[i].x;
+				int splitIndexA = bodyOffsetA+constraintOffsetA;
+				dlvAPtr = &deltaLinearVelocities[splitIndexA];
+				davAPtr = &deltaAngularVelocities[splitIndexA];
+			}
+
+			if (bodyB.getInvMass())
+			{
+				int bodyOffsetB = offsetSplitBodies[bIdx];
+				int constraintOffsetB = contactConstraintOffsets[i].y;
+				int splitIndexB= bodyOffsetB+constraintOffsetB;
+				dlvBPtr =&deltaLinearVelocities[splitIndexB];
+				davBPtr = &deltaAngularVelocities[splitIndexB];
+			}
+
+
+
+			{
+				float maxRambdaDt[4] = {FLT_MAX,FLT_MAX,FLT_MAX,FLT_MAX};
+				float minRambdaDt[4] = {0.f,0.f,0.f,0.f};
+
+				solveContact( contactConstraints[i], (btVector3&)bodyA.m_pos, (btVector3&)bodyA.m_linVel, (btVector3&)bodyA.m_angVel, bodyA.m_invMass, inertiasCPU[aIdx].m_invInertiaWorld, 
+					(btVector3&)bodyB.m_pos, (btVector3&)bodyB.m_linVel, (btVector3&)bodyB.m_angVel, bodyB.m_invMass, inertiasCPU[bIdx].m_invInertiaWorld,
+					maxRambdaDt, minRambdaDt , *dlvAPtr,*davAPtr,*dlvBPtr,*davBPtr		);
+
+
+			}
+		}
+
+		
+		{
+			BT_PROFILE("average velocities");
+			btLauncherCL launcher( m_queue, m_data->m_averageVelocitiesKernel);
+			launcher.setBuffer(bodiesGPU->getBufferCL());
+			launcher.setBuffer(m_data->m_offsetSplitBodies->getBufferCL());
+			launcher.setBuffer(m_data->m_bodyCount->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaLinearVelocities->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaAngularVelocities->getBufferCL());
+			launcher.setConst(numBodies);
+			launcher.launch1D(numBodies);
+			clFinish(m_queue);
+		}
+
+		//easy
+		for (int i=0;i<numBodiesCPU;i++)
+		{
+			if (bodiesCPU[i].getInvMass())
+			{
+				int bodyOffset = offsetSplitBodies[i];
+				int count = bodyCount[i];
+				float factor = 1.f/float(count);
+				btVector3 averageLinVel;
+				averageLinVel.setZero();
+				btVector3 averageAngVel;
+				averageAngVel.setZero();
+				for (int j=0;j<count;j++)
+				{
+					averageLinVel += deltaLinearVelocities[bodyOffset+j]*factor;
+					averageAngVel += deltaAngularVelocities[bodyOffset+j]*factor;
+				}
+				for (int j=0;j<count;j++)
+				{
+					deltaLinearVelocities[bodyOffset+j] = averageLinVel;
+					deltaAngularVelocities[bodyOffset+j] = averageAngVel;
+				}
+			}
+		}
+//	m_data->m_deltaAngularVelocities->copyFromHost(deltaAngularVelocities);
+	//m_data->m_deltaLinearVelocities->copyFromHost(deltaLinearVelocities);
+	m_data->m_deltaAngularVelocities->copyToHost(deltaAngularVelocities);
+	m_data->m_deltaLinearVelocities->copyToHost(deltaLinearVelocities);
+
+#if 0
+
+		{
+			BT_PROFILE("m_solveFrictionKernel");
+			btLauncherCL launcher( m_queue, m_data->m_solveFrictionKernel);
+			launcher.setBuffer(m_data->m_contactConstraints->getBufferCL());
+			launcher.setBuffer(bodiesGPU->getBufferCL());
+			launcher.setBuffer(inertiasGPU->getBufferCL());
+			launcher.setBuffer(m_data->m_contactConstraintOffsets->getBufferCL());
+			launcher.setBuffer(m_data->m_offsetSplitBodies->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaLinearVelocities->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaAngularVelocities->getBufferCL());
+			launcher.setConst(solverInfo.m_deltaTime);
+			launcher.setConst(solverInfo.m_positionDrift);
+			launcher.setConst(solverInfo.m_positionConstraintCoeff);
+			launcher.setConst(solverInfo.m_fixedBodyIndex);
+			launcher.setConst(numManifolds);
+
+			launcher.launch1D(numManifolds);
+			clFinish(m_queue);
+		}
+
+		//solve friction
+
+		for(int i=0; i<numManifoldsCPU; i++)
+		{
+			float maxRambdaDt[4] = {FLT_MAX,FLT_MAX,FLT_MAX,FLT_MAX};
+			float minRambdaDt[4] = {0.f,0.f,0.f,0.f};
+
+			float sum = 0;
+			for(int j=0; j<4; j++)
+			{
+				sum +=contactConstraints[i].m_appliedRambdaDt[j];
+			}
+			float frictionCoeff = contactConstraints[i].getFrictionCoeff();
+			int aIdx = (int)contactConstraints[i].m_bodyA;
+			int bIdx = (int)contactConstraints[i].m_bodyB;
+			btRigidBodyCL& bodyA = bodiesCPU[aIdx];
+			btRigidBodyCL& bodyB = bodiesCPU[bIdx];
+
+			btVector3 zero(0,0,0);
+			
+			btVector3* dlvAPtr=&zero;
+			btVector3* davAPtr=&zero;
+			btVector3* dlvBPtr=&zero;
+			btVector3* davBPtr=&zero;
+			
+			if (bodyA.getInvMass())
+			{
+				int bodyOffsetA = offsetSplitBodies[aIdx];
+				int constraintOffsetA = contactConstraintOffsets[i].x;
+				int splitIndexA = bodyOffsetA+constraintOffsetA;
+				dlvAPtr = &deltaLinearVelocities[splitIndexA];
+				davAPtr = &deltaAngularVelocities[splitIndexA];
+			}
+
+			if (bodyB.getInvMass())
+			{
+				int bodyOffsetB = offsetSplitBodies[bIdx];
+				int constraintOffsetB = contactConstraintOffsets[i].y;
+				int splitIndexB= bodyOffsetB+constraintOffsetB;
+				dlvBPtr =&deltaLinearVelocities[splitIndexB];
+				davBPtr = &deltaAngularVelocities[splitIndexB];
+			}
+
+			for(int j=0; j<4; j++)
+			{
+				maxRambdaDt[j] = frictionCoeff*sum;
+				minRambdaDt[j] = -maxRambdaDt[j];
+			}
+
+			solveFriction( contactConstraints[i], (btVector3&)bodyA.m_pos, (btVector3&)bodyA.m_linVel, (btVector3&)bodyA.m_angVel, bodyA.m_invMass,inertiasCPU[aIdx].m_invInertiaWorld, 
+				(btVector3&)bodyB.m_pos, (btVector3&)bodyB.m_linVel, (btVector3&)bodyB.m_angVel, bodyB.m_invMass, inertiasCPU[bIdx].m_invInertiaWorld,
+				maxRambdaDt, minRambdaDt , *dlvAPtr,*davAPtr,*dlvBPtr,*davBPtr);
+
+		}
+
+		{
+			BT_PROFILE("average velocities");
+			btLauncherCL launcher( m_queue, m_data->m_averageVelocitiesKernel);
+			launcher.setBuffer(bodiesGPU->getBufferCL());
+			launcher.setBuffer(m_data->m_offsetSplitBodies->getBufferCL());
+			launcher.setBuffer(m_data->m_bodyCount->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaLinearVelocities->getBufferCL());
+			launcher.setBuffer(m_data->m_deltaAngularVelocities->getBufferCL());
+			launcher.setConst(numBodies);
+			launcher.launch1D(numBodies);
+			clFinish(m_queue);
+		}
+
+		//easy
+		for (int i=0;i<numBodiesCPU;i++)
+		{
+			if (bodiesCPU[i].getInvMass())
+			{
+				int bodyOffset = offsetSplitBodies[i];
+				int count = bodyCount[i];
+				float factor = 1.f/float(count);
+				btVector3 averageLinVel;
+				averageLinVel.setZero();
+				btVector3 averageAngVel;
+				averageAngVel.setZero();
+				for (int j=0;j<count;j++)
+				{
+					averageLinVel += deltaLinearVelocities[bodyOffset+j]*factor;
+					averageAngVel += deltaAngularVelocities[bodyOffset+j]*factor;
+				}
+				for (int j=0;j<count;j++)
+				{
+					deltaLinearVelocities[bodyOffset+j] = averageLinVel;
+					deltaAngularVelocities[bodyOffset+j] = averageAngVel;
+				}
+			}
+		}
+
+#endif
+
+	}
+
+	{
+		BT_PROFILE("update body velocities");
+		btLauncherCL launcher( m_queue, m_data->m_updateBodyVelocitiesKernel);
+		launcher.setBuffer(bodiesGPU->getBufferCL());
+		launcher.setBuffer(m_data->m_offsetSplitBodies->getBufferCL());
+		launcher.setBuffer(m_data->m_bodyCount->getBufferCL());
+		launcher.setBuffer(m_data->m_deltaLinearVelocities->getBufferCL());
+		launcher.setBuffer(m_data->m_deltaAngularVelocities->getBufferCL());
+		launcher.setConst(numBodies);
+		launcher.launch1D(numBodies);
+		clFinish(m_queue);
+	}
+
+
+	//easy
+	for (int i=0;i<numBodiesCPU;i++)
+	{
+		if (bodiesCPU[i].getInvMass())
+		{
+			int bodyOffset = offsetSplitBodies[i];
+			int count = bodyCount[i];
+			if (count)
+			{
+				bodiesCPU[i].m_linVel += deltaLinearVelocities[bodyOffset];
+				bodiesCPU[i].m_angVel += deltaAngularVelocities[bodyOffset];
+			}
+		}
+	}
+
+
+//	bodiesGPU->copyFromHost(bodiesCPU);
+
+
 }

opencl/gpu_rigidbody/host/btGpuJacobiSolver.h

@@ -46,6 +46,7 @@ public:
 
 	void  solveGroupHost(btRigidBodyCL* bodies,btInertiaCL* inertias,int numBodies,btContact4* manifoldPtr, int numManifolds,btTypedConstraint** constraints,int numConstraints,const btJacobiSolverInfo& solverInfo);
 	void  solveGroup(btOpenCLArray<btRigidBodyCL>* bodies,btOpenCLArray<btInertiaCL>* inertias,btOpenCLArray<btContact4>* manifoldPtr,const btJacobiSolverInfo& solverInfo);
+	void  solveGroupMixed(btOpenCLArray<btRigidBodyCL>* bodies,btOpenCLArray<btInertiaCL>* inertias,btOpenCLArray<btContact4>* manifoldPtr,const btJacobiSolverInfo& solverInfo);
 
 };
 #endif //BT_GPU_JACOBI_SOLVER_H

opencl/gpu_rigidbody/host/btGpuNarrowPhase.cpp

@@ -196,6 +196,85 @@ int	btGpuNarrowPhase::allocateCollidable()
 
 
 
+
+
+int		btGpuNarrowPhase::registerSphereShape(float radius)
+{
+	int collidableIndex = allocateCollidable();
+
+	btCollidable& col = getCollidableCpu(collidableIndex);
+	col.m_shapeType = SHAPE_SPHERE;
+	col.m_shapeIndex = 0;
+	col.m_radius = radius;
+	
+	if (col.m_shapeIndex>=0)
+	{
+		btSapAabb aabb;
+		btVector3 myAabbMin(-radius,-radius,-radius);
+		btVector3 myAabbMax(radius,radius,radius);
+
+		aabb.m_min[0] = myAabbMin[0];//s_convexHeightField->m_aabb.m_min.x;
+		aabb.m_min[1] = myAabbMin[1];//s_convexHeightField->m_aabb.m_min.y;
+		aabb.m_min[2] = myAabbMin[2];//s_convexHeightField->m_aabb.m_min.z;
+		aabb.m_minIndices[3] = 0;
+
+		aabb.m_max[0] = myAabbMax[0];//s_convexHeightField->m_aabb.m_max.x;
+		aabb.m_max[1] = myAabbMax[1];//s_convexHeightField->m_aabb.m_max.y;
+		aabb.m_max[2] = myAabbMax[2];//s_convexHeightField->m_aabb.m_max.z;
+		aabb.m_signedMaxIndices[3] = 0;
+
+		m_data->m_localShapeAABBCPU->push_back(aabb);
+		m_data->m_localShapeAABBGPU->push_back(aabb);
+		clFinish(m_queue);
+	}
+	
+	return collidableIndex;
+}
+
+
+int btGpuNarrowPhase::registerFace(const btVector3& faceNormal, float faceConstant)
+{
+	int faceOffset = m_data->m_convexFaces.size();
+	btGpuFace& face = m_data->m_convexFaces.expand();
+	face.m_plane[0] = faceNormal.getX();
+	face.m_plane[1] = faceNormal.getY();
+	face.m_plane[2] = faceNormal.getZ();
+	face.m_plane[3] = faceConstant;
+	m_data->m_convexFacesGPU->copyFromHost(m_data->m_convexFaces);
+	return faceOffset;
+}
+
+int		btGpuNarrowPhase::registerPlaneShape(const btVector3& planeNormal, float planeConstant)
+{
+	int collidableIndex = allocateCollidable();
+
+	btCollidable& col = getCollidableCpu(collidableIndex);
+	col.m_shapeType = SHAPE_PLANE;
+	col.m_shapeIndex = registerFace(planeNormal,planeConstant);
+	col.m_radius = planeConstant;
+	
+	if (col.m_shapeIndex>=0)
+	{
+		btSapAabb aabb;
+		aabb.m_min[0] = -1e30f;
+		aabb.m_min[1] = -1e30f;
+		aabb.m_min[2] = -1e30f;
+		aabb.m_minIndices[3] = 0;
+		
+		aabb.m_max[0] = 1e30f;
+		aabb.m_max[1] = 1e30f;
+		aabb.m_max[2] = 1e30f;
+		aabb.m_signedMaxIndices[3] = 0;
+
+		m_data->m_localShapeAABBCPU->push_back(aabb);
+		m_data->m_localShapeAABBGPU->push_back(aabb);
+		clFinish(m_queue);
+	}
+	
+	return collidableIndex;
+}
+
+
 int btGpuNarrowPhase::registerConvexHullShape(btConvexUtility* convexPtr,btCollidable& col)
 {
 	m_data->m_convexData->resize(m_data->m_numAcceleratedShapes+1);

opencl/gpu_rigidbody/host/btGpuNarrowPhase.h

@@ -31,6 +31,8 @@ public:
 
 	virtual ~btGpuNarrowPhase(void);
 
+	int		registerSphereShape(float radius);
+	int		registerPlaneShape(const btVector3& planeNormal, float planeConstant);
 
 	int registerCompoundShape(btAlignedObjectArray<btGpuChildShape>* childShapes);
 	int registerFace(const btVector3& faceNormal, float faceConstant);

opencl/gpu_rigidbody/kernels/solveContact.cl

@@ -237,7 +237,7 @@ void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4
 
 void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)
 {
-	*linear = -n;
+	*linear = mymake_float4(-n.xyz,0.f);
 	*angular0 = -cross3(r0, n);
 	*angular1 = cross3(r1, n);
 }

opencl/gpu_rigidbody/kernels/solveContact.h

@@ -239,7 +239,7 @@ static const char* solveContactCL= \
 "\n"
 "void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)\n"
 "{\n"
-"	*linear = -n;\n"
+"	*linear = mymake_float4(-n.xyz,0.f);\n"
 "	*angular0 = -cross3(r0, n);\n"
 "	*angular1 = cross3(r1, n);\n"
 "}\n"

opencl/gpu_rigidbody/kernels/solveFriction.cl

@@ -237,7 +237,7 @@ void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4
 
 void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)
 {
-	*linear = -n;
+	*linear = mymake_float4(-n.xyz,0.f);
 	*angular0 = -cross3(r0, n);
 	*angular1 = cross3(r1, n);
 }

opencl/gpu_rigidbody/kernels/solveFriction.h

@@ -239,7 +239,7 @@ static const char* solveFrictionCL= \
 "\n"
 "void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)\n"
 "{\n"
-"	*linear = -n;\n"
+"	*linear = mymake_float4(-n.xyz,0.f);\n"
 "	*angular0 = -cross3(r0, n);\n"
 "	*angular1 = cross3(r1, n);\n"
 "}\n"

opencl/gpu_rigidbody/kernels/solverSetup.cl

@@ -435,7 +435,7 @@ typedef struct
 
 void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)
 {
-	*linear = -n;
+	*linear = make_float4(-n.xyz,0.f);
 	*angular0 = -cross3(r0, n);
 	*angular1 = cross3(r1, n);
 }

opencl/gpu_rigidbody/kernels/solverSetup.h

@@ -437,7 +437,7 @@ static const char* solverSetupCL= \
 "\n"
 "void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)\n"
 "{\n"
-"	*linear = -n;\n"
+"	*linear = make_float4(-n.xyz,0.f);\n"
 "	*angular0 = -cross3(r0, n);\n"
 "	*angular1 = cross3(r1, n);\n"
 "}\n"

opencl/gpu_rigidbody/kernels/solverUtils.cl

@@ -462,7 +462,7 @@ __global float4* deltaLinearVelocities, __global float4* deltaAngularVelocities,
 
 void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)
 {
-	*linear = -n;
+	*linear = make_float4(-n.xyz,0.f);
 	*angular0 = -cross3(r0, n);
 	*angular1 = cross3(r1, n);
 }
@@ -537,10 +537,12 @@ void solveContact(__global Constraint4* cs,
 		setLinearAndAngular( -cs->m_linear, r0, r1, &linear, &angular0, &angular1 );
 	
 
+
 		float rambdaDt = calcRelVel( cs->m_linear, -cs->m_linear, angular0, angular1, 
 			*linVelA+*dLinVelA, *angVelA+*dAngVelA, *linVelB+*dLinVelB, *angVelB+*dAngVelB ) + cs->m_b[ic];
 		rambdaDt *= cs->m_jacCoeffInv[ic];
 
+		
 		{
 			float prevSum = cs->m_appliedRambdaDt[ic];
 			float updated = prevSum;
@@ -551,11 +553,13 @@ void solveContact(__global Constraint4* cs,
 			cs->m_appliedRambdaDt[ic] = updated;
 		}
 
+			
 		float4 linImp0 = invMassA*linear*rambdaDt;
 		float4 linImp1 = invMassB*(-linear)*rambdaDt;
 		float4 angImp0 = mtMul1(invInertiaA, angular0)*rambdaDt;
 		float4 angImp1 = mtMul1(invInertiaB, angular1)*rambdaDt;
 
+		
 		if (invMassA)
 		{
 			*dLinVelA += linImp0;

opencl/gpu_rigidbody/kernels/solverUtils.h

@@ -464,7 +464,7 @@ static const char* solverUtilsCL= \
 "\n"
 "void setLinearAndAngular( float4 n, float4 r0, float4 r1, float4* linear, float4* angular0, float4* angular1)\n"
 "{\n"
-"	*linear = -n;\n"
+"	*linear = make_float4(-n.xyz,0.f);\n"
 "	*angular0 = -cross3(r0, n);\n"
 "	*angular1 = cross3(r1, n);\n"
 "}\n"
@@ -539,10 +539,12 @@ static const char* solverUtilsCL= \
 "		setLinearAndAngular( -cs->m_linear, r0, r1, &linear, &angular0, &angular1 );\n"
 "	\n"
 "\n"
+"\n"
 "		float rambdaDt = calcRelVel( cs->m_linear, -cs->m_linear, angular0, angular1, \n"
 "			*linVelA+*dLinVelA, *angVelA+*dAngVelA, *linVelB+*dLinVelB, *angVelB+*dAngVelB ) + cs->m_b[ic];\n"
 "		rambdaDt *= cs->m_jacCoeffInv[ic];\n"
 "\n"
+"		\n"
 "		{\n"
 "			float prevSum = cs->m_appliedRambdaDt[ic];\n"
 "			float updated = prevSum;\n"
@@ -552,12 +554,14 @@ static const char* solverUtilsCL= \
 "			rambdaDt = updated - prevSum;\n"
 "			cs->m_appliedRambdaDt[ic] = updated;\n"
 "		}\n"
+"\n"
 "			\n"
 "		float4 linImp0 = invMassA*linear*rambdaDt;\n"
 "		float4 linImp1 = invMassB*(-linear)*rambdaDt;\n"
 "		float4 angImp0 = mtMul1(invInertiaA, angular0)*rambdaDt;\n"
 "		float4 angImp1 = mtMul1(invInertiaB, angular1)*rambdaDt;\n"
 "\n"
+"		\n"
 "		if (invMassA)\n"
 "		{\n"
 "			*dLinVelA += linImp0;\n"

opencl/gpu_sat/host/ConvexHullContact.cpp

@@ -36,6 +36,7 @@ typedef btAlignedObjectArray<btVector3> btVertexArray;
 #include "../kernels/satKernels.h"
 #include "../kernels/satClipHullContacts.h"
 #include "../kernels/bvhTraversal.h"
+#include "../kernels/primitiveContacts.h"
 
 #include "BulletGeometry/btAabbUtil2.h"
 
@@ -113,15 +114,24 @@ m_totalContactsOut(m_context, m_queue)
 	 if (1)
 	{
 		const char* srcBvh = bvhTraversalKernelCL;
-		//cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl");
+		cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,srcBvh,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl");
-		cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl", true);
+		//cl_program bvhTraversalProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,0,&errNum,"","opencl/gpu_sat/kernels/bvhTraversal.cl", true);
 		btAssert(errNum==CL_SUCCESS);
 
-		m_bvhTraversalKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcBvh, "bvhTraversalKernel",&errNum,bvhTraversalProg,"-g");
+		m_bvhTraversalKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,srcBvh, "bvhTraversalKernel",&errNum,bvhTraversalProg,"");
 		btAssert(errNum==CL_SUCCESS);
 
 	}
         
+	 {
+		 const char* primitiveContactsSrc = primitiveContactsKernelsCL;
+		cl_program primitiveContactsProg = btOpenCLUtils::compileCLProgramFromString(m_context,m_device,primitiveContactsSrc,&errNum,"","opencl/gpu_sat/kernels/primitiveContacts.cl");
+		btAssert(errNum==CL_SUCCESS);
+
+		m_primitiveContactsKernel = btOpenCLUtils::compileCLKernelFromString(m_context, m_device,primitiveContactsSrc, "primitiveContactsKernel",&errNum,primitiveContactsProg,"");
+		btAssert(errNum==CL_SUCCESS);
+		 
+	 }
 	
 
 }
@@ -148,6 +158,8 @@ GpuSatCollision::~GpuSatCollision()
         clReleaseKernel(m_clipFacesAndContactReductionKernel);
     if (m_newContactReductionKernel)
         clReleaseKernel(m_newContactReductionKernel);
+	if (m_primitiveContactsKernel)
+		clReleaseKernel(m_primitiveContactsKernel);
     
 	if (m_clipHullHullKernel)
 		clReleaseKernel(m_clipHullHullKernel);
@@ -176,6 +188,224 @@ struct MyTriangleCallback : public btNodeOverlapCallback
 	}
 };
 
+
+#define float4 btVector3
+#define make_float4(x,y,z,w) btVector4(x,y,z,w)
+
+float signedDistanceFromPointToPlane(const float4& point, const float4& planeEqn, float4* closestPointOnFace)
+{
+	float4 n = planeEqn;
+	n[3] = 0.f;
+	float dist = dot3F4(n, point) + planeEqn[3];
+	*closestPointOnFace = point - dist * n;
+	return dist;
+}
+
+
+
+inline bool IsPointInPolygon(const btVector3& p,
+							const btVector3& posConvex,
+							const btQuaternion& ornConvex,
+							const btGpuFace* face,
+							const btVector3* baseVertex,
+							const int* convexIndices,
+                            btVector3* out)
+{
+    btVector3 a;
+    btVector3 b;
+    btVector3 ab;
+    btVector3 ap;
+    btVector3 v;
+
+	btVector3 plane (face->m_plane[0],face->m_plane[1],face->m_plane[2]);
+	
+	if (face->m_numIndices<2)
+		return false;
+
+	btTransform tr;
+	tr.setIdentity();
+	tr.setOrigin(posConvex);
+	tr.setRotation(ornConvex);
+
+	float4 v0 = baseVertex[convexIndices[face->m_indexOffset + face->m_numIndices-1]];
+	btVector3 worldV0 = tr(v0);
+	b = worldV0;
+
+    for(unsigned i=0; i != face->m_numIndices; ++i)
+    {
+		a = b;
+		float4 vi = baseVertex[convexIndices[face->m_indexOffset + i]];
+		btVector3 worldVi = tr(vi);
+		b = worldVi;
+        ab = b-a;
+        ap = p-a;
+        v = ab.cross(plane);
+
+        if (btDot(ap, v) > 0.f)
+        {
+            btScalar ab_m2 = btDot(ab, ab);
+            btScalar s = ab_m2 != btScalar(0.0) ? btDot(ab, ap) / ab_m2 : btScalar(0.0);
+            if (s <= btScalar(0.0))
+            {
+                *out = a;
+            }
+            else if (s >= btScalar(1.0)) 
+            {
+                *out = b;
+            }
+            else
+            {
+                out->setInterpolate3(a,b,s);
+            }
+            return false;
+        }
+    }
+    return true;
+}
+
+
+
+void	computeContactSphereConvex(int pairIndex,
+																int bodyIndexA, int bodyIndexB, 
+																int collidableIndexA, int collidableIndexB, 
+																const btRigidBodyCL* rigidBodies, 
+																const btCollidable* collidables,
+																const btConvexPolyhedronCL* convexShapes,
+																const btVector3* convexVertices,
+																const int* convexIndices,
+																const btGpuFace* faces,
+																btContact4* globalContactsOut,
+																int& nGlobalContactsOut,
+																int maxContactCapacity)
+{
+	
+	float radius = collidables[collidableIndexA].m_radius;
+	float4 spherePos1 = rigidBodies[bodyIndexA].m_pos;
+	btQuaternion sphereOrn = rigidBodies[bodyIndexA].m_quat;
+
+
+
+	float4 pos = rigidBodies[bodyIndexB].m_pos;
+	float4 spherePos = spherePos1-pos;
+	btQuaternion quat = rigidBodies[bodyIndexB].m_quat;
+
+	int collidableIndex = rigidBodies[bodyIndexB].m_collidableIdx;
+	int shapeIndex = collidables[collidableIndex].m_shapeIndex;
+	int numFaces = convexShapes[shapeIndex].m_numFaces;
+	float4 closestPnt = make_float4(0, 0, 0, 0);
+	float4 hitNormalWorld = make_float4(0, 0, 0, 0);
+	float minDist = -1000000.f; // TODO: What is the largest/smallest float?
+	bool bCollide = true;
+	int region = -1;
+	for ( int f = 0; f < numFaces; f++ )
+	{
+		btGpuFace face = faces[convexShapes[shapeIndex].m_faceOffset+f];
+		float4 planeEqn;
+		float4 localPlaneNormal = make_float4(face.m_plane.x(),face.m_plane.y(),face.m_plane.z(),0.f);
+		float4 n1 = quatRotate(quat,localPlaneNormal);
+		planeEqn = n1;
+		planeEqn[3] = face.m_plane[3];
+
+		float4 pntReturn;
+		float dist = signedDistanceFromPointToPlane(spherePos, planeEqn, &pntReturn);
+
+		if ( dist > radius)
+		{
+			bCollide = false;
+			break;
+		}
+
+		if ( dist > 0 )
+		{
+			//might hit an edge or vertex
+			btVector3 out;
+			bool isInPoly = IsPointInPolygon(spherePos,
+					pos,
+					quat,
+					&face,
+					&convexVertices[convexShapes[shapeIndex].m_vertexOffset],
+					convexIndices,
+                    &out);
+			if (isInPoly)
+			{
+				if (dist>minDist)
+				{
+					minDist = dist;
+					closestPnt = pntReturn;
+					hitNormalWorld = planeEqn;
+					region=1;
+				}
+			} else
+			{
+				btVector3 tmp = spherePos-out;
+				btScalar l2 = tmp.length2();
+				if (l2<radius*radius)
+				{
+					dist  = btSqrt(l2);
+					if (dist>minDist)
+					{
+						minDist = dist;
+						closestPnt = out;
+						hitNormalWorld = tmp/dist;
+						region=2;
+					}
+					
+				} else
+				{
+					bCollide = false;
+					break;
+				}
+			}
+		}
+		else
+		{
+			if ( dist > minDist )
+			{
+				minDist = dist;
+				closestPnt = pntReturn;
+				hitNormalWorld = planeEqn;
+				region=3;
+			}
+		}
+	}
+	
+	
+	if (bCollide && minDist > -100)
+	{
+		float4 normalOnSurfaceB1 = -hitNormalWorld;
+		float4 pOnB1 = closestPnt+pos;
+		//printf("dist ,%f,",minDist);
+		float actualDepth = minDist-radius;
+		//printf("actualDepth = ,%f,", actualDepth);
+		//printf("normalOnSurfaceB1 = ,%f,%f,%f,", normalOnSurfaceB1.getX(),normalOnSurfaceB1.getY(),normalOnSurfaceB1.getZ());
+		//printf("region=,%d,\n", region);
+		pOnB1[3] = actualDepth;
+
+		int dstIdx;
+//    dstIdx = nGlobalContactsOut++;//AppendInc( nGlobalContactsOut, dstIdx );
+		
+		if (nGlobalContactsOut < maxContactCapacity)
+		{
+			dstIdx=nGlobalContactsOut;
+			nGlobalContactsOut++;
+
+			btContact4* c = &globalContactsOut[dstIdx];
+			c->m_worldNormal = normalOnSurfaceB1;
+			c->setFrictionCoeff(0.7);
+			c->setRestituitionCoeff(0.f);
+
+			c->m_batchIdx = pairIndex;
+			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
+			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
+			c->m_worldPos[0] = pOnB1;
+			int numPoints = 1;
+			c->m_worldNormal[3] = numPoints;
+		}//if (dstIdx < numPairs)
+	}//if (hasCollision)
+	
+}
+
+
 void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btInt2>* pairs, int nPairs,
 			const btOpenCLArray<btRigidBodyCL>* bodyBuf,
 			btOpenCLArray<btContact4>* contactOut, int& nContacts,
@@ -206,6 +436,113 @@ void GpuSatCollision::computeConvexConvexContactsGPUSAT( const btOpenCLArray<btI
 	if (!nPairs)
 		return;
 
+
+//#define CHECK_ON_HOST
+#ifdef CHECK_ON_HOST
+	btAlignedObjectArray<btYetAnotherAabb> hostAabbs;
+	clAabbsWS.copyToHost(hostAabbs);
+	btAlignedObjectArray<btInt2> hostPairs;
+	pairs->copyToHost(hostPairs);
+
+	btAlignedObjectArray<btRigidBodyCL> hostBodyBuf;
+	bodyBuf->copyToHost(hostBodyBuf);
+
+	
+
+	btAlignedObjectArray<btConvexPolyhedronCL> hostConvexData;
+	convexData.copyToHost(hostConvexData);
+
+	btAlignedObjectArray<btVector3> hostVertices;
+	gpuVertices.copyToHost(hostVertices);
+
+	btAlignedObjectArray<btVector3> hostUniqueEdges;
+	gpuUniqueEdges.copyToHost(hostUniqueEdges);
+	btAlignedObjectArray<btGpuFace> hostFaces;
+	gpuFaces.copyToHost(hostFaces);
+	btAlignedObjectArray<int> hostIndices;
+	gpuIndices.copyToHost(hostIndices);
+	btAlignedObjectArray<btCollidable> hostCollidables;
+	gpuCollidables.copyToHost(hostCollidables);
+	
+	btAlignedObjectArray<btGpuChildShape> cpuChildShapes;
+	gpuChildShapes.copyToHost(cpuChildShapes);
+	
+
+	btAlignedObjectArray<btInt4> hostTriangleConvexPairs;
+
+	btAlignedObjectArray<btContact4> hostContacts;
+	if (nContacts)
+	{
+		contactOut->copyToHost(hostContacts);
+	}
+
+	hostContacts.resize(nPairs);
+
+	for (int i=0;i<nPairs;i++)
+	{
+		int bodyIndexA = hostPairs[i].x;
+		int bodyIndexB = hostPairs[i].y;
+		int collidableIndexA = hostBodyBuf[bodyIndexA].m_collidableIdx;
+		int collidableIndexB = hostBodyBuf[bodyIndexB].m_collidableIdx;
+
+		if (hostCollidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&
+			hostCollidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)
+		{
+			printf("sphere-convex\n");
+		}
+
+		if (hostCollidables[collidableIndexA].m_shapeType == SHAPE_CONVEX_HULL &&
+			hostCollidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
+		{
+			computeContactSphereConvex(i,bodyIndexB,bodyIndexA,collidableIndexB,collidableIndexA,&hostBodyBuf[0],
+				&hostCollidables[0],&hostConvexData[0],&hostVertices[0],&hostIndices[0],&hostFaces[0],&hostContacts[0],nContacts,nPairs);
+			//printf("convex-sphere\n");
+			
+		}
+
+		
+	}
+
+	if (nContacts)
+		{
+			hostContacts.resize(nContacts);
+			contactOut->copyFromHost(hostContacts);
+		}
+
+	
+#else
+
+	{
+		if (nPairs)
+		{
+			m_totalContactsOut.copyFromHostPointer(&nContacts,1,0,true);
+
+			BT_PROFILE("primitiveContactsKernel");
+			btBufferInfoCL bInfo[] = {
+				btBufferInfoCL( pairs->getBufferCL(), true ), 
+				btBufferInfoCL( bodyBuf->getBufferCL(),true), 
+				btBufferInfoCL( gpuCollidables.getBufferCL(),true), 
+				btBufferInfoCL( convexData.getBufferCL(),true),
+				btBufferInfoCL( gpuVertices.getBufferCL(),true),
+				btBufferInfoCL( gpuUniqueEdges.getBufferCL(),true),
+				btBufferInfoCL( gpuFaces.getBufferCL(),true),
+				btBufferInfoCL( gpuIndices.getBufferCL(),true),
+				btBufferInfoCL( contactOut->getBufferCL()),
+				btBufferInfoCL( m_totalContactsOut.getBufferCL())	
+			};
+			
+			btLauncherCL launcher(m_queue, m_primitiveContactsKernel);
+			launcher.setBuffers( bInfo, sizeof(bInfo)/sizeof(btBufferInfoCL) );
+			launcher.setConst( nPairs  );
+			int num = nPairs;
+			launcher.launch1D( num);
+			clFinish(m_queue);
+		
+			nContacts = m_totalContactsOut.at(0);
+			contactOut->resize(nContacts);
+		}
+	}
+#endif//CHECK_ON_HOST
 	BT_PROFILE("computeConvexConvexContactsGPUSAT");
    // printf("nContacts = %d\n",nContacts);
     

opencl/gpu_sat/host/ConvexHullContact.h

@@ -52,6 +52,7 @@ struct GpuSatCollision
     cl_kernel               m_newContactReductionKernel;
 
 	cl_kernel				m_bvhTraversalKernel;
+	cl_kernel				m_primitiveContactsKernel;
     
 
 	btOpenCLArray<int>		m_totalContactsOut;

opencl/gpu_sat/kernels/primitiveContacts.cl

@@ -0,0 +1,667 @@
+#define TRIANGLE_NUM_CONVEX_FACES 5
+
+#define SHAPE_CONVEX_HULL 3
+#define SHAPE_PLANE 4
+#define SHAPE_CONCAVE_TRIMESH 5
+#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6
+#define SHAPE_SPHERE 7
+
+
+#pragma OPENCL EXTENSION cl_amd_printf : enable
+#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
+
+#ifdef cl_ext_atomic_counters_32
+#pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable
+#else
+#define counter32_t volatile __global int*
+#endif
+
+#define GET_GROUP_IDX get_group_id(0)
+#define GET_LOCAL_IDX get_local_id(0)
+#define GET_GLOBAL_IDX get_global_id(0)
+#define GET_GROUP_SIZE get_local_size(0)
+#define GET_NUM_GROUPS get_num_groups(0)
+#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)
+#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)
+#define AtomInc(x) atom_inc(&(x))
+#define AtomInc1(x, out) out = atom_inc(&(x))
+#define AppendInc(x, out) out = atomic_inc(x)
+#define AtomAdd(x, value) atom_add(&(x), value)
+#define AtomCmpxhg(x, cmp, value) atom_cmpxchg( &(x), cmp, value )
+#define AtomXhg(x, value) atom_xchg ( &(x), value )
+
+#define max2 max
+#define min2 min
+
+typedef unsigned int u32;
+
+
+
+typedef struct
+{
+	float4 m_worldPos[4];
+	float4 m_worldNormal;	//	w: m_nPoints
+	u32 m_coeffs;
+	u32 m_batchIdx;
+
+	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
+	int m_bodyBPtrAndSignBit;
+} Contact4;
+
+
+///keep this in sync with btCollidable.h
+typedef struct
+{
+	int m_numChildShapes;
+	float m_radius;
+	int m_shapeType;
+	int m_shapeIndex;
+	
+} btCollidableGpu;
+
+typedef struct
+{
+	float4	m_childPosition;
+	float4	m_childOrientation;
+	int m_shapeIndex;
+	int m_unused0;
+	int m_unused1;
+	int m_unused2;
+} btGpuChildShape;
+
+#define GET_NPOINTS(x) (x).m_worldNormal.w
+
+typedef struct
+{
+	float4 m_pos;
+	float4 m_quat;
+	float4 m_linVel;
+	float4 m_angVel;
+
+	u32 m_collidableIdx;	
+	float m_invMass;
+	float m_restituitionCoeff;
+	float m_frictionCoeff;
+} BodyData;
+
+
+typedef struct  
+{
+	float4		m_localCenter;
+	float4		m_extents;
+	float4		mC;
+	float4		mE;
+	
+	float			m_radius;
+	int	m_faceOffset;
+	int m_numFaces;
+	int	m_numVertices;
+	
+	int m_vertexOffset;
+	int	m_uniqueEdgesOffset;
+	int	m_numUniqueEdges;
+	int m_unused;
+
+} ConvexPolyhedronCL;
+
+typedef struct
+{
+	float4 m_plane;
+	int m_indexOffset;
+	int m_numIndices;
+} btGpuFace;
+
+#define SELECT_UINT4( b, a, condition ) select( b,a,condition )
+
+#define make_float4 (float4)
+#define make_float2 (float2)
+#define make_uint4 (uint4)
+#define make_int4 (int4)
+#define make_uint2 (uint2)
+#define make_int2 (int2)
+
+
+__inline
+float fastDiv(float numerator, float denominator)
+{
+	return native_divide(numerator, denominator);	
+//	return numerator/denominator;	
+}
+
+__inline
+float4 fastDiv4(float4 numerator, float4 denominator)
+{
+	return native_divide(numerator, denominator);	
+}
+
+
+__inline
+float4 cross3(float4 a, float4 b)
+{
+	return cross(a,b);
+}
+
+//#define dot3F4 dot
+
+__inline
+float dot3F4(float4 a, float4 b)
+{
+	float4 a1 = make_float4(a.xyz,0.f);
+	float4 b1 = make_float4(b.xyz,0.f);
+	return dot(a1, b1);
+}
+
+__inline
+float4 fastNormalize4(float4 v)
+{
+	return fast_normalize(v);
+}
+
+
+///////////////////////////////////////
+//	Quaternion
+///////////////////////////////////////
+
+typedef float4 Quaternion;
+
+__inline
+Quaternion qtMul(Quaternion a, Quaternion b);
+
+__inline
+Quaternion qtNormalize(Quaternion in);
+
+__inline
+float4 qtRotate(Quaternion q, float4 vec);
+
+__inline
+Quaternion qtInvert(Quaternion q);
+
+
+
+
+__inline
+Quaternion qtMul(Quaternion a, Quaternion b)
+{
+	Quaternion ans;
+	ans = cross3( a, b );
+	ans += a.w*b+b.w*a;
+//	ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);
+	ans.w = a.w*b.w - dot3F4(a, b);
+	return ans;
+}
+
+__inline
+Quaternion qtNormalize(Quaternion in)
+{
+	return fastNormalize4(in);
+//	in /= length( in );
+//	return in;
+}
+__inline
+float4 qtRotate(Quaternion q, float4 vec)
+{
+	Quaternion qInv = qtInvert( q );
+	float4 vcpy = vec;
+	vcpy.w = 0.f;
+	float4 out = qtMul(qtMul(q,vcpy),qInv);
+	return out;
+}
+
+__inline
+Quaternion qtInvert(Quaternion q)
+{
+	return (Quaternion)(-q.xyz, q.w);
+}
+
+__inline
+float4 qtInvRotate(const Quaternion q, float4 vec)
+{
+	return qtRotate( qtInvert( q ), vec );
+}
+
+__inline
+float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)
+{
+	return qtRotate( *orientation, *p ) + (*translation);
+}
+
+void	trInverse(float4 translationIn, Quaternion orientationIn,
+		float4* translationOut, Quaternion* orientationOut)
+{
+	*orientationOut = qtInvert(orientationIn);
+	*translationOut = qtRotate(*orientationOut, -translationIn);
+}
+
+void	trMul(float4 translationA, Quaternion orientationA,
+						float4 translationB, Quaternion orientationB,
+		float4* translationOut, Quaternion* orientationOut)
+{
+	*orientationOut = qtMul(orientationA,orientationB);
+	*translationOut = transform(&translationB,&translationA,&orientationA);
+}
+
+
+
+__inline
+float4 normalize3(const float4 a)
+{
+	float4 n = make_float4(a.x, a.y, a.z, 0.f);
+	return fastNormalize4( n );
+}
+
+
+__inline float4 lerp3(const float4 a,const float4 b, float  t)
+{
+	return make_float4(	a.x + (b.x - a.x) * t,
+						a.y + (b.y - a.y) * t,
+						a.z + (b.z - a.z) * t,
+						0.f);
+}
+
+
+float signedDistanceFromPointToPlane(float4 point, float4 planeEqn, float4* closestPointOnFace)
+{
+	float4 n = (float4)(planeEqn.x, planeEqn.y, planeEqn.z, 0);
+	float dist = dot3F4(n, point) + planeEqn.w;
+	*closestPointOnFace = point - dist * n;
+	return dist;
+}
+
+
+
+inline bool IsPointInPolygon(float4 p, 
+							float4	posConvex,
+							float4	ornConvex,
+							const btGpuFace* face,
+							__global const float4* baseVertex,
+							__global const  int* convexIndices,
+							float4* out)
+{
+    float4 a;
+    float4 b;
+    float4 ab;
+    float4 ap;
+    float4 v;
+
+	float4 plane = make_float4(face->m_plane.x,face->m_plane.y,face->m_plane.z,0.f);
+	
+	if (face->m_numIndices<2)
+		return false;
+
+	
+	float4 v0 = baseVertex[convexIndices[face->m_indexOffset + face->m_numIndices-1]];
+	float4 worldV0 = transform(&v0, &posConvex, &ornConvex);
+	
+	b = worldV0;
+
+    for(unsigned i=0; i != face->m_numIndices; ++i)
+    {
+		a = b;
+		float4 vi = baseVertex[convexIndices[face->m_indexOffset + i]];
+		float4 worldVi = transform(&vi, &posConvex, &ornConvex);
+		b = worldVi;
+        ab = b-a;
+        ap = p-a;
+        v = cross3(ab,plane);
+
+        if (dot(ap, v) > 0.f)
+        {
+            float ab_m2 = dot(ab, ab);
+            float rt = ab_m2 != 0.f ? dot(ab, ap) / ab_m2 : 0.f;
+            if (rt <= 0.f)
+            {
+                *out = a;
+            }
+            else if (rt >= 1.f) 
+            {
+                *out = b;
+            }
+            else
+            {
+            	float s = 1.f - rt;
+							out[0].x = s * a.x + rt * b.x;
+							out[0].y = s * a.y + rt * b.y;
+							out[0].z = s * a.z + rt * b.z;
+            }
+            return false;
+        }
+    }
+    return true;
+}
+
+
+
+
+void	computeContactSphereConvex(int pairIndex,
+																int bodyIndexA, int bodyIndexB, 
+																int collidableIndexA, int collidableIndexB, 
+																__global const BodyData* rigidBodies, 
+																__global const btCollidableGpu* collidables,
+																__global const ConvexPolyhedronCL* convexShapes,
+																__global const float4* convexVertices,
+																__global const int* convexIndices,
+																__global const btGpuFace* faces,
+																__global Contact4* restrict globalContactsOut,
+																counter32_t nGlobalContactsOut,
+																int numPairs)
+{
+
+	float radius = collidables[collidableIndexA].m_radius;
+	float4 spherePos1 = rigidBodies[bodyIndexA].m_pos;
+	float4 sphereOrn = rigidBodies[bodyIndexA].m_quat;
+
+
+
+	float4 pos = rigidBodies[bodyIndexB].m_pos;
+	float4 quat = rigidBodies[bodyIndexB].m_quat;
+
+	float4 spherePos = spherePos1 - pos;
+
+	int collidableIndex = rigidBodies[bodyIndexB].m_collidableIdx;
+	int shapeIndex = collidables[collidableIndex].m_shapeIndex;
+	int numFaces = convexShapes[shapeIndex].m_numFaces;
+	float4 closestPnt = (float4)(0, 0, 0, 0);
+	float4 hitNormalWorld = (float4)(0, 0, 0, 0);
+	float minDist = -1000000.f;
+	bool bCollide = true;
+
+	for ( int f = 0; f < numFaces; f++ )
+	{
+		btGpuFace face = faces[convexShapes[shapeIndex].m_faceOffset+f];
+
+		// set up a plane equation 
+		float4 planeEqn;
+		float4 n1 = qtRotate(quat, (float4)(face.m_plane.xyz, 0));
+		planeEqn = n1;
+		planeEqn.w = face.m_plane.w;
+		
+	
+		// compute a signed distance from the vertex in cloth to the face of rigidbody.
+		float4 pntReturn;
+		float dist = signedDistanceFromPointToPlane(spherePos, planeEqn, &pntReturn);
+
+		// If the distance is positive, the plane is a separating plane. 
+		if ( dist > radius )
+		{
+			bCollide = false;
+			break;
+		}
+
+
+		if (dist>0)
+		{
+			//might hit an edge or vertex
+			float4 out;
+			bool isInPoly = IsPointInPolygon(spherePos,
+					pos,
+					quat,
+					&face,
+					&convexVertices[convexShapes[shapeIndex].m_vertexOffset],
+					convexIndices,
+           &out);
+			if (isInPoly)
+			{
+				if (dist>minDist)
+				{
+					minDist = dist;
+					closestPnt = pntReturn;
+					hitNormalWorld = planeEqn;
+					
+				}
+			} else
+			{
+				float4 tmp = spherePos-out;
+				float l2 = dot(tmp,tmp);
+				if (l2<radius*radius)
+				{
+					dist  = sqrt(l2);
+					if (dist>minDist)
+					{
+						minDist = dist;
+						closestPnt = out;
+						hitNormalWorld = tmp/dist;
+						
+					}
+					
+				} else
+				{
+					bCollide = false;
+					break;
+				}
+			}
+		} else
+		{
+			if ( dist > minDist )
+			{
+				minDist = dist;
+				closestPnt = pntReturn;
+				hitNormalWorld.xyz = planeEqn.xyz;
+			}
+		}
+		
+	}
+
+	
+
+	if (bCollide)
+	{
+		float4 normalOnSurfaceB1 = -hitNormalWorld;
+		float4 pOnB1 = closestPnt+pos;
+		float actualDepth = minDist-radius;
+		pOnB1.w = actualDepth;
+
+		int dstIdx;
+    AppendInc( nGlobalContactsOut, dstIdx );
+		
+		if (dstIdx < numPairs)
+		{
+			__global Contact4* c = &globalContactsOut[dstIdx];
+			c->m_worldNormal = normalOnSurfaceB1;
+			c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
+			c->m_batchIdx = pairIndex;
+			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
+			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
+			c->m_worldPos[0] = pOnB1;
+			GET_NPOINTS(*c) = 1;
+		}//if (dstIdx < numPairs)
+	}//if (hasCollision)
+
+}
+							
+
+									
+void	computeContactPlaneConvex(int pairIndex,
+																int bodyIndexA, int bodyIndexB, 
+																int collidableIndexA, int collidableIndexB, 
+																__global const BodyData* rigidBodies, 
+																__global const btCollidableGpu* collidables,
+																__global const btGpuFace* faces,
+																__global Contact4* restrict globalContactsOut,
+																counter32_t nGlobalContactsOut,
+																int numPairs)
+{
+	float4 planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;
+	float radius = collidables[collidableIndexB].m_radius;
+	float4 posA1 = rigidBodies[bodyIndexA].m_pos;
+	float4 ornA1 = rigidBodies[bodyIndexA].m_quat;
+	float4 posB1 = rigidBodies[bodyIndexB].m_pos;
+	float4 ornB1 = rigidBodies[bodyIndexB].m_quat;
+	
+	bool hasCollision = false;
+	float4 planeNormal1 = make_float4(planeEq.x,planeEq.y,planeEq.z,0.f);
+	float planeConstant = planeEq.w;
+	float4 convexInPlaneTransPos1; Quaternion convexInPlaneTransOrn1;
+	{
+		float4 invPosA;Quaternion invOrnA;
+		trInverse(posA1,ornA1,&invPosA,&invOrnA);
+		trMul(invPosA,invOrnA,posB1,ornB1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);
+	}
+	float4 planeInConvexPos1;	Quaternion planeInConvexOrn1;
+	{
+		float4 invPosB;Quaternion invOrnB;
+		trInverse(posB1,ornB1,&invPosB,&invOrnB);
+		trMul(invPosB,invOrnB,posA1,ornA1,&planeInConvexPos1,&planeInConvexOrn1);	
+	}
+	float4 vtx1 = qtRotate(planeInConvexOrn1,-planeNormal1)*radius;
+	float4 vtxInPlane1 = transform(&vtx1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);
+	float distance = dot3F4(planeNormal1,vtxInPlane1) - planeConstant;
+	hasCollision = distance < 0.f;//m_manifoldPtr->getContactBreakingThreshold();
+	if (hasCollision)
+	{
+		float4 vtxInPlaneProjected1 = vtxInPlane1 -   distance*planeNormal1;
+		float4 vtxInPlaneWorld1 = transform(&vtxInPlaneProjected1,&posA1,&ornA1);
+		float4 normalOnSurfaceB1 = qtRotate(ornA1,planeNormal1);
+		float4 pOnB1 = vtxInPlaneWorld1+normalOnSurfaceB1*distance;
+		pOnB1.w = distance;
+
+		int dstIdx;
+    AppendInc( nGlobalContactsOut, dstIdx );
+		
+		if (dstIdx < numPairs)
+		{
+			__global Contact4* c = &globalContactsOut[dstIdx];
+			c->m_worldNormal = normalOnSurfaceB1;
+			c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
+			c->m_batchIdx = pairIndex;
+			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
+			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
+			c->m_worldPos[0] = pOnB1;
+			GET_NPOINTS(*c) = 1;
+		}//if (dstIdx < numPairs)
+	}//if (hasCollision)
+}
+
+
+
+
+__kernel void   primitiveContactsKernel( __global const int2* pairs, 
+																					__global const BodyData* rigidBodies, 
+																					__global const btCollidableGpu* collidables,
+																					__global const ConvexPolyhedronCL* convexShapes, 
+																					__global const float4* vertices,
+																					__global const float4* uniqueEdges,
+																					__global const btGpuFace* faces,
+																					__global const int* indices,
+																					__global Contact4* restrict globalContactsOut,
+																					counter32_t nGlobalContactsOut,
+																					int numPairs)
+{
+
+	int i = get_global_id(0);
+	int pairIndex = i;
+	
+	float4 worldVertsB1[64];
+	float4 worldVertsB2[64];
+	int capacityWorldVerts = 64;	
+
+	float4 localContactsOut[64];
+	int localContactCapacity=64;
+	
+	float minDist = -1e30f;
+	float maxDist = 0.02f;
+
+	if (i<numPairs)
+	{
+
+		int bodyIndexA = pairs[i].x;
+		int bodyIndexB = pairs[i].y;
+			
+		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
+		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
+	
+
+		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&
+			collidables[collidableIndexB].m_shapeType == SHAPE_PLANE)
+		{
+
+
+			computeContactPlaneConvex( pairIndex, bodyIndexB,bodyIndexA,  collidableIndexB,collidableIndexA, 
+																rigidBodies,collidables,faces,	globalContactsOut, nGlobalContactsOut,numPairs);
+			return;
+		}
+
+		if (collidables[collidableIndexA].m_shapeType == SHAPE_PLANE &&
+			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
+		{
+
+
+			computeContactPlaneConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, 
+																rigidBodies,collidables,faces, globalContactsOut, nGlobalContactsOut,numPairs);
+			return;
+			
+		}
+	
+		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&
+			collidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)
+		{
+		
+			computeContactSphereConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, 
+																rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,numPairs);
+			return;
+		}
+
+		if (collidables[collidableIndexA].m_shapeType == SHAPE_CONVEX_HULL &&
+			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
+		{
+		
+			computeContactSphereConvex(pairIndex, bodyIndexB, bodyIndexA, collidableIndexB, collidableIndexA, 
+																rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,numPairs);
+			return;
+		}
+	
+	
+	
+	
+	
+		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&
+			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
+		{
+			//sphere-sphere
+			float radiusA = collidables[collidableIndexA].m_radius;
+			float radiusB = collidables[collidableIndexB].m_radius;
+			float4 posA = rigidBodies[bodyIndexA].m_pos;
+			float4 posB = rigidBodies[bodyIndexB].m_pos;
+
+			float4 diff = posA-posB;
+			float len = length(diff);
+			
+			///iff distance positive, don't generate a new contact
+			if ( len <= (radiusA+radiusB))
+			{
+				///distance (negative means penetration)
+				float dist = len - (radiusA+radiusB);
+				float4 normalOnSurfaceB = make_float4(1.f,0.f,0.f,0.f);
+				if (len > 0.00001)
+				{
+					normalOnSurfaceB = diff / len;
+				}
+				float4 contactPosB = posB + normalOnSurfaceB*radiusB;
+				contactPosB.w = dist;
+								
+				int dstIdx;
+        AppendInc( nGlobalContactsOut, dstIdx );
+				
+				if (dstIdx < numPairs)
+				{
+					__global Contact4* c = &globalContactsOut[dstIdx];
+					c->m_worldNormal = -normalOnSurfaceB;
+					c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
+					c->m_batchIdx = pairIndex;
+					int bodyA = pairs[pairIndex].x;
+					int bodyB = pairs[pairIndex].y;
+					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
+					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
+					c->m_worldPos[0] = contactPosB;
+					GET_NPOINTS(*c) = 1;
+				}//if (dstIdx < numPairs)
+			}//if ( len <= (radiusA+radiusB))
+
+			return;
+		}//SHAPE_SPHERE SHAPE_SPHERE
+
+	}//	if (i<numPairs)
+
+}

opencl/gpu_sat/kernels/primitiveContacts.h

@@ -0,0 +1,671 @@
+//this file is autogenerated using stringify.bat (premake --stringify) in the build folder of this project
+static const char* primitiveContactsKernelsCL= \
+"#define TRIANGLE_NUM_CONVEX_FACES 5\n"
+"\n"
+"#define SHAPE_CONVEX_HULL 3\n"
+"#define SHAPE_PLANE 4\n"
+"#define SHAPE_CONCAVE_TRIMESH 5\n"
+"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
+"#define SHAPE_SPHERE 7\n"
+"\n"
+"\n"
+"#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
+"#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
+"#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
+"#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable\n"
+"#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable\n"
+"\n"
+"#ifdef cl_ext_atomic_counters_32\n"
+"#pragma OPENCL EXTENSION cl_ext_atomic_counters_32 : enable\n"
+"#else\n"
+"#define counter32_t volatile __global int*\n"
+"#endif\n"
+"\n"
+"#define GET_GROUP_IDX get_group_id(0)\n"
+"#define GET_LOCAL_IDX get_local_id(0)\n"
+"#define GET_GLOBAL_IDX get_global_id(0)\n"
+"#define GET_GROUP_SIZE get_local_size(0)\n"
+"#define GET_NUM_GROUPS get_num_groups(0)\n"
+"#define GROUP_LDS_BARRIER barrier(CLK_LOCAL_MEM_FENCE)\n"
+"#define GROUP_MEM_FENCE mem_fence(CLK_LOCAL_MEM_FENCE)\n"
+"#define AtomInc(x) atom_inc(&(x))\n"
+"#define AtomInc1(x, out) out = atom_inc(&(x))\n"
+"#define AppendInc(x, out) out = atomic_inc(x)\n"
+"#define AtomAdd(x, value) atom_add(&(x), value)\n"
+"#define AtomCmpxhg(x, cmp, value) atom_cmpxchg( &(x), cmp, value )\n"
+"#define AtomXhg(x, value) atom_xchg ( &(x), value )\n"
+"\n"
+"#define max2 max\n"
+"#define min2 min\n"
+"\n"
+"typedef unsigned int u32;\n"
+"\n"
+"\n"
+"\n"
+"typedef struct\n"
+"{\n"
+"	float4 m_worldPos[4];\n"
+"	float4 m_worldNormal;	//	w: m_nPoints\n"
+"	u32 m_coeffs;\n"
+"	u32 m_batchIdx;\n"
+"\n"
+"	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
+"	int m_bodyBPtrAndSignBit;\n"
+"} Contact4;\n"
+"\n"
+"\n"
+"///keep this in sync with btCollidable.h\n"
+"typedef struct\n"
+"{\n"
+"	int m_numChildShapes;\n"
+"	float m_radius;\n"
+"	int m_shapeType;\n"
+"	int m_shapeIndex;\n"
+"	\n"
+"} btCollidableGpu;\n"
+"\n"
+"typedef struct\n"
+"{\n"
+"	float4	m_childPosition;\n"
+"	float4	m_childOrientation;\n"
+"	int m_shapeIndex;\n"
+"	int m_unused0;\n"
+"	int m_unused1;\n"
+"	int m_unused2;\n"
+"} btGpuChildShape;\n"
+"\n"
+"#define GET_NPOINTS(x) (x).m_worldNormal.w\n"
+"\n"
+"typedef struct\n"
+"{\n"
+"	float4 m_pos;\n"
+"	float4 m_quat;\n"
+"	float4 m_linVel;\n"
+"	float4 m_angVel;\n"
+"\n"
+"	u32 m_collidableIdx;	\n"
+"	float m_invMass;\n"
+"	float m_restituitionCoeff;\n"
+"	float m_frictionCoeff;\n"
+"} BodyData;\n"
+"\n"
+"\n"
+"typedef struct  \n"
+"{\n"
+"	float4		m_localCenter;\n"
+"	float4		m_extents;\n"
+"	float4		mC;\n"
+"	float4		mE;\n"
+"	\n"
+"	float			m_radius;\n"
+"	int	m_faceOffset;\n"
+"	int m_numFaces;\n"
+"	int	m_numVertices;\n"
+"	\n"
+"	int m_vertexOffset;\n"
+"	int	m_uniqueEdgesOffset;\n"
+"	int	m_numUniqueEdges;\n"
+"	int m_unused;\n"
+"\n"
+"} ConvexPolyhedronCL;\n"
+"\n"
+"typedef struct\n"
+"{\n"
+"	float4 m_plane;\n"
+"	int m_indexOffset;\n"
+"	int m_numIndices;\n"
+"} btGpuFace;\n"
+"\n"
+"#define SELECT_UINT4( b, a, condition ) select( b,a,condition )\n"
+"\n"
+"#define make_float4 (float4)\n"
+"#define make_float2 (float2)\n"
+"#define make_uint4 (uint4)\n"
+"#define make_int4 (int4)\n"
+"#define make_uint2 (uint2)\n"
+"#define make_int2 (int2)\n"
+"\n"
+"\n"
+"__inline\n"
+"float fastDiv(float numerator, float denominator)\n"
+"{\n"
+"	return native_divide(numerator, denominator);	\n"
+"//	return numerator/denominator;	\n"
+"}\n"
+"\n"
+"__inline\n"
+"float4 fastDiv4(float4 numerator, float4 denominator)\n"
+"{\n"
+"	return native_divide(numerator, denominator);	\n"
+"}\n"
+"\n"
+"\n"
+"__inline\n"
+"float4 cross3(float4 a, float4 b)\n"
+"{\n"
+"	return cross(a,b);\n"
+"}\n"
+"\n"
+"//#define dot3F4 dot\n"
+"\n"
+"__inline\n"
+"float dot3F4(float4 a, float4 b)\n"
+"{\n"
+"	float4 a1 = make_float4(a.xyz,0.f);\n"
+"	float4 b1 = make_float4(b.xyz,0.f);\n"
+"	return dot(a1, b1);\n"
+"}\n"
+"\n"
+"__inline\n"
+"float4 fastNormalize4(float4 v)\n"
+"{\n"
+"	return fast_normalize(v);\n"
+"}\n"
+"\n"
+"\n"
+"///////////////////////////////////////\n"
+"//	Quaternion\n"
+"///////////////////////////////////////\n"
+"\n"
+"typedef float4 Quaternion;\n"
+"\n"
+"__inline\n"
+"Quaternion qtMul(Quaternion a, Quaternion b);\n"
+"\n"
+"__inline\n"
+"Quaternion qtNormalize(Quaternion in);\n"
+"\n"
+"__inline\n"
+"float4 qtRotate(Quaternion q, float4 vec);\n"
+"\n"
+"__inline\n"
+"Quaternion qtInvert(Quaternion q);\n"
+"\n"
+"\n"
+"\n"
+"\n"
+"__inline\n"
+"Quaternion qtMul(Quaternion a, Quaternion b)\n"
+"{\n"
+"	Quaternion ans;\n"
+"	ans = cross3( a, b );\n"
+"	ans += a.w*b+b.w*a;\n"
+"//	ans.w = a.w*b.w - (a.x*b.x+a.y*b.y+a.z*b.z);\n"
+"	ans.w = a.w*b.w - dot3F4(a, b);\n"
+"	return ans;\n"
+"}\n"
+"\n"
+"__inline\n"
+"Quaternion qtNormalize(Quaternion in)\n"
+"{\n"
+"	return fastNormalize4(in);\n"
+"//	in /= length( in );\n"
+"//	return in;\n"
+"}\n"
+"__inline\n"
+"float4 qtRotate(Quaternion q, float4 vec)\n"
+"{\n"
+"	Quaternion qInv = qtInvert( q );\n"
+"	float4 vcpy = vec;\n"
+"	vcpy.w = 0.f;\n"
+"	float4 out = qtMul(qtMul(q,vcpy),qInv);\n"
+"	return out;\n"
+"}\n"
+"\n"
+"__inline\n"
+"Quaternion qtInvert(Quaternion q)\n"
+"{\n"
+"	return (Quaternion)(-q.xyz, q.w);\n"
+"}\n"
+"\n"
+"__inline\n"
+"float4 qtInvRotate(const Quaternion q, float4 vec)\n"
+"{\n"
+"	return qtRotate( qtInvert( q ), vec );\n"
+"}\n"
+"\n"
+"__inline\n"
+"float4 transform(const float4* p, const float4* translation, const Quaternion* orientation)\n"
+"{\n"
+"	return qtRotate( *orientation, *p ) + (*translation);\n"
+"}\n"
+"\n"
+"void	trInverse(float4 translationIn, Quaternion orientationIn,\n"
+"		float4* translationOut, Quaternion* orientationOut)\n"
+"{\n"
+"	*orientationOut = qtInvert(orientationIn);\n"
+"	*translationOut = qtRotate(*orientationOut, -translationIn);\n"
+"}\n"
+"\n"
+"void	trMul(float4 translationA, Quaternion orientationA,\n"
+"						float4 translationB, Quaternion orientationB,\n"
+"		float4* translationOut, Quaternion* orientationOut)\n"
+"{\n"
+"	*orientationOut = qtMul(orientationA,orientationB);\n"
+"	*translationOut = transform(&translationB,&translationA,&orientationA);\n"
+"}\n"
+"\n"
+"\n"
+"\n"
+"__inline\n"
+"float4 normalize3(const float4 a)\n"
+"{\n"
+"	float4 n = make_float4(a.x, a.y, a.z, 0.f);\n"
+"	return fastNormalize4( n );\n"
+"}\n"
+"\n"
+"\n"
+"__inline float4 lerp3(const float4 a,const float4 b, float  t)\n"
+"{\n"
+"	return make_float4(	a.x + (b.x - a.x) * t,\n"
+"						a.y + (b.y - a.y) * t,\n"
+"						a.z + (b.z - a.z) * t,\n"
+"						0.f);\n"
+"}\n"
+"\n"
+"\n"
+"float signedDistanceFromPointToPlane(float4 point, float4 planeEqn, float4* closestPointOnFace)\n"
+"{\n"
+"	float4 n = (float4)(planeEqn.x, planeEqn.y, planeEqn.z, 0);\n"
+"	float dist = dot3F4(n, point) + planeEqn.w;\n"
+"	*closestPointOnFace = point - dist * n;\n"
+"	return dist;\n"
+"}\n"
+"\n"
+"\n"
+"\n"
+"inline bool IsPointInPolygon(float4 p, \n"
+"							float4	posConvex,\n"
+"							float4	ornConvex,\n"
+"							const btGpuFace* face,\n"
+"							__global const float4* baseVertex,\n"
+"							__global const  int* convexIndices,\n"
+"							float4* out)\n"
+"{\n"
+"    float4 a;\n"
+"    float4 b;\n"
+"    float4 ab;\n"
+"    float4 ap;\n"
+"    float4 v;\n"
+"\n"
+"	float4 plane = make_float4(face->m_plane.x,face->m_plane.y,face->m_plane.z,0.f);\n"
+"	\n"
+"	if (face->m_numIndices<2)\n"
+"		return false;\n"
+"\n"
+"	\n"
+"	float4 v0 = baseVertex[convexIndices[face->m_indexOffset + face->m_numIndices-1]];\n"
+"	float4 worldV0 = transform(&v0, &posConvex, &ornConvex);\n"
+"	\n"
+"	b = worldV0;\n"
+"\n"
+"    for(unsigned i=0; i != face->m_numIndices; ++i)\n"
+"    {\n"
+"		a = b;\n"
+"		float4 vi = baseVertex[convexIndices[face->m_indexOffset + i]];\n"
+"		float4 worldVi = transform(&vi, &posConvex, &ornConvex);\n"
+"		b = worldVi;\n"
+"        ab = b-a;\n"
+"        ap = p-a;\n"
+"        v = cross3(ab,plane);\n"
+"\n"
+"        if (dot(ap, v) > 0.f)\n"
+"        {\n"
+"            float ab_m2 = dot(ab, ab);\n"
+"            float rt = ab_m2 != 0.f ? dot(ab, ap) / ab_m2 : 0.f;\n"
+"            if (rt <= 0.f)\n"
+"            {\n"
+"                *out = a;\n"
+"            }\n"
+"            else if (rt >= 1.f) \n"
+"            {\n"
+"                *out = b;\n"
+"            }\n"
+"            else\n"
+"            {\n"
+"            	float s = 1.f - rt;\n"
+"							out[0].x = s * a.x + rt * b.x;\n"
+"							out[0].y = s * a.y + rt * b.y;\n"
+"							out[0].z = s * a.z + rt * b.z;\n"
+"            }\n"
+"            return false;\n"
+"        }\n"
+"    }\n"
+"    return true;\n"
+"}\n"
+"\n"
+"\n"
+"\n"
+"\n"
+"void	computeContactSphereConvex(int pairIndex,\n"
+"																int bodyIndexA, int bodyIndexB, \n"
+"																int collidableIndexA, int collidableIndexB, \n"
+"																__global const BodyData* rigidBodies, \n"
+"																__global const btCollidableGpu* collidables,\n"
+"																__global const ConvexPolyhedronCL* convexShapes,\n"
+"																__global const float4* convexVertices,\n"
+"																__global const int* convexIndices,\n"
+"																__global const btGpuFace* faces,\n"
+"																__global Contact4* restrict globalContactsOut,\n"
+"																counter32_t nGlobalContactsOut,\n"
+"																int numPairs)\n"
+"{\n"
+"\n"
+"	float radius = collidables[collidableIndexA].m_radius;\n"
+"	float4 spherePos1 = rigidBodies[bodyIndexA].m_pos;\n"
+"	float4 sphereOrn = rigidBodies[bodyIndexA].m_quat;\n"
+"\n"
+"\n"
+"\n"
+"	float4 pos = rigidBodies[bodyIndexB].m_pos;\n"
+"	float4 quat = rigidBodies[bodyIndexB].m_quat;\n"
+"\n"
+"	float4 spherePos = spherePos1 - pos;\n"
+"\n"
+"	int collidableIndex = rigidBodies[bodyIndexB].m_collidableIdx;\n"
+"	int shapeIndex = collidables[collidableIndex].m_shapeIndex;\n"
+"	int numFaces = convexShapes[shapeIndex].m_numFaces;\n"
+"	float4 closestPnt = (float4)(0, 0, 0, 0);\n"
+"	float4 hitNormalWorld = (float4)(0, 0, 0, 0);\n"
+"	float minDist = -1000000.f;\n"
+"	bool bCollide = true;\n"
+"\n"
+"	for ( int f = 0; f < numFaces; f++ )\n"
+"	{\n"
+"		btGpuFace face = faces[convexShapes[shapeIndex].m_faceOffset+f];\n"
+"\n"
+"		// set up a plane equation \n"
+"		float4 planeEqn;\n"
+"		float4 n1 = qtRotate(quat, (float4)(face.m_plane.xyz, 0));\n"
+"		planeEqn = n1;\n"
+"		planeEqn.w = face.m_plane.w;\n"
+"		\n"
+"	\n"
+"		// compute a signed distance from the vertex in cloth to the face of rigidbody.\n"
+"		float4 pntReturn;\n"
+"		float dist = signedDistanceFromPointToPlane(spherePos, planeEqn, &pntReturn);\n"
+"\n"
+"		// If the distance is positive, the plane is a separating plane. \n"
+"		if ( dist > radius )\n"
+"		{\n"
+"			bCollide = false;\n"
+"			break;\n"
+"		}\n"
+"\n"
+"\n"
+"		if (dist>0)\n"
+"		{\n"
+"			//might hit an edge or vertex\n"
+"			float4 out;\n"
+"			bool isInPoly = IsPointInPolygon(spherePos,\n"
+"					pos,\n"
+"					quat,\n"
+"					&face,\n"
+"					&convexVertices[convexShapes[shapeIndex].m_vertexOffset],\n"
+"					convexIndices,\n"
+"           &out);\n"
+"			if (isInPoly)\n"
+"			{\n"
+"				if (dist>minDist)\n"
+"				{\n"
+"					minDist = dist;\n"
+"					closestPnt = pntReturn;\n"
+"					hitNormalWorld = planeEqn;\n"
+"					\n"
+"				}\n"
+"			} else\n"
+"			{\n"
+"				float4 tmp = spherePos-out;\n"
+"				float l2 = dot(tmp,tmp);\n"
+"				if (l2<radius*radius)\n"
+"				{\n"
+"					dist  = sqrt(l2);\n"
+"					if (dist>minDist)\n"
+"					{\n"
+"						minDist = dist;\n"
+"						closestPnt = out;\n"
+"						hitNormalWorld = tmp/dist;\n"
+"						\n"
+"					}\n"
+"					\n"
+"				} else\n"
+"				{\n"
+"					bCollide = false;\n"
+"					break;\n"
+"				}\n"
+"			}\n"
+"		} else\n"
+"		{\n"
+"			if ( dist > minDist )\n"
+"			{\n"
+"				minDist = dist;\n"
+"				closestPnt = pntReturn;\n"
+"				hitNormalWorld.xyz = planeEqn.xyz;\n"
+"			}\n"
+"		}\n"
+"		\n"
+"	}\n"
+"\n"
+"	\n"
+"\n"
+"	if (bCollide)\n"
+"	{\n"
+"		float4 normalOnSurfaceB1 = -hitNormalWorld;\n"
+"		float4 pOnB1 = closestPnt+pos;\n"
+"		float actualDepth = minDist-radius;\n"
+"		pOnB1.w = actualDepth;\n"
+"\n"
+"		int dstIdx;\n"
+"    AppendInc( nGlobalContactsOut, dstIdx );\n"
+"		\n"
+"		if (dstIdx < numPairs)\n"
+"		{\n"
+"			__global Contact4* c = &globalContactsOut[dstIdx];\n"
+"			c->m_worldNormal = normalOnSurfaceB1;\n"
+"			c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
+"			c->m_batchIdx = pairIndex;\n"
+"			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
+"			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
+"			c->m_worldPos[0] = pOnB1;\n"
+"			GET_NPOINTS(*c) = 1;\n"
+"		}//if (dstIdx < numPairs)\n"
+"	}//if (hasCollision)\n"
+"\n"
+"}\n"
+"							\n"
+"\n"
+"									\n"
+"void	computeContactPlaneConvex(int pairIndex,\n"
+"																int bodyIndexA, int bodyIndexB, \n"
+"																int collidableIndexA, int collidableIndexB, \n"
+"																__global const BodyData* rigidBodies, \n"
+"																__global const btCollidableGpu* collidables,\n"
+"																__global const btGpuFace* faces,\n"
+"																__global Contact4* restrict globalContactsOut,\n"
+"																counter32_t nGlobalContactsOut,\n"
+"																int numPairs)\n"
+"{\n"
+"	float4 planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;\n"
+"	float radius = collidables[collidableIndexB].m_radius;\n"
+"	float4 posA1 = rigidBodies[bodyIndexA].m_pos;\n"
+"	float4 ornA1 = rigidBodies[bodyIndexA].m_quat;\n"
+"	float4 posB1 = rigidBodies[bodyIndexB].m_pos;\n"
+"	float4 ornB1 = rigidBodies[bodyIndexB].m_quat;\n"
+"	\n"
+"	bool hasCollision = false;\n"
+"	float4 planeNormal1 = make_float4(planeEq.x,planeEq.y,planeEq.z,0.f);\n"
+"	float planeConstant = planeEq.w;\n"
+"	float4 convexInPlaneTransPos1; Quaternion convexInPlaneTransOrn1;\n"
+"	{\n"
+"		float4 invPosA;Quaternion invOrnA;\n"
+"		trInverse(posA1,ornA1,&invPosA,&invOrnA);\n"
+"		trMul(invPosA,invOrnA,posB1,ornB1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);\n"
+"	}\n"
+"	float4 planeInConvexPos1;	Quaternion planeInConvexOrn1;\n"
+"	{\n"
+"		float4 invPosB;Quaternion invOrnB;\n"
+"		trInverse(posB1,ornB1,&invPosB,&invOrnB);\n"
+"		trMul(invPosB,invOrnB,posA1,ornA1,&planeInConvexPos1,&planeInConvexOrn1);	\n"
+"	}\n"
+"	float4 vtx1 = qtRotate(planeInConvexOrn1,-planeNormal1)*radius;\n"
+"	float4 vtxInPlane1 = transform(&vtx1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);\n"
+"	float distance = dot3F4(planeNormal1,vtxInPlane1) - planeConstant;\n"
+"	hasCollision = distance < 0.f;//m_manifoldPtr->getContactBreakingThreshold();\n"
+"	if (hasCollision)\n"
+"	{\n"
+"		float4 vtxInPlaneProjected1 = vtxInPlane1 -   distance*planeNormal1;\n"
+"		float4 vtxInPlaneWorld1 = transform(&vtxInPlaneProjected1,&posA1,&ornA1);\n"
+"		float4 normalOnSurfaceB1 = qtRotate(ornA1,planeNormal1);\n"
+"		float4 pOnB1 = vtxInPlaneWorld1+normalOnSurfaceB1*distance;\n"
+"		pOnB1.w = distance;\n"
+"\n"
+"		int dstIdx;\n"
+"    AppendInc( nGlobalContactsOut, dstIdx );\n"
+"		\n"
+"		if (dstIdx < numPairs)\n"
+"		{\n"
+"			__global Contact4* c = &globalContactsOut[dstIdx];\n"
+"			c->m_worldNormal = normalOnSurfaceB1;\n"
+"			c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
+"			c->m_batchIdx = pairIndex;\n"
+"			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
+"			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
+"			c->m_worldPos[0] = pOnB1;\n"
+"			GET_NPOINTS(*c) = 1;\n"
+"		}//if (dstIdx < numPairs)\n"
+"	}//if (hasCollision)\n"
+"}\n"
+"\n"
+"\n"
+"\n"
+"\n"
+"__kernel void   primitiveContactsKernel( __global const int2* pairs, \n"
+"																					__global const BodyData* rigidBodies, \n"
+"																					__global const btCollidableGpu* collidables,\n"
+"																					__global const ConvexPolyhedronCL* convexShapes, \n"
+"																					__global const float4* vertices,\n"
+"																					__global const float4* uniqueEdges,\n"
+"																					__global const btGpuFace* faces,\n"
+"																					__global const int* indices,\n"
+"																					__global Contact4* restrict globalContactsOut,\n"
+"																					counter32_t nGlobalContactsOut,\n"
+"																					int numPairs)\n"
+"{\n"
+"\n"
+"	int i = get_global_id(0);\n"
+"	int pairIndex = i;\n"
+"	\n"
+"	float4 worldVertsB1[64];\n"
+"	float4 worldVertsB2[64];\n"
+"	int capacityWorldVerts = 64;	\n"
+"\n"
+"	float4 localContactsOut[64];\n"
+"	int localContactCapacity=64;\n"
+"	\n"
+"	float minDist = -1e30f;\n"
+"	float maxDist = 0.02f;\n"
+"\n"
+"	if (i<numPairs)\n"
+"	{\n"
+"\n"
+"		int bodyIndexA = pairs[i].x;\n"
+"		int bodyIndexB = pairs[i].y;\n"
+"			\n"
+"		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
+"		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
+"	\n"
+"\n"
+"		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
+"			collidables[collidableIndexB].m_shapeType == SHAPE_PLANE)\n"
+"		{\n"
+"\n"
+"\n"
+"			computeContactPlaneConvex( pairIndex, bodyIndexB,bodyIndexA,  collidableIndexB,collidableIndexA, \n"
+"																rigidBodies,collidables,faces,	globalContactsOut, nGlobalContactsOut,numPairs);\n"
+"			return;\n"
+"		}\n"
+"\n"
+"		if (collidables[collidableIndexA].m_shapeType == SHAPE_PLANE &&\n"
+"			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
+"		{\n"
+"\n"
+"\n"
+"			computeContactPlaneConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, \n"
+"																rigidBodies,collidables,faces, globalContactsOut, nGlobalContactsOut,numPairs);\n"
+"			return;\n"
+"			\n"
+"		}\n"
+"	\n"
+"		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
+"			collidables[collidableIndexB].m_shapeType == SHAPE_CONVEX_HULL)\n"
+"		{\n"
+"		\n"
+"			computeContactSphereConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, \n"
+"																rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,numPairs);\n"
+"			return;\n"
+"		}\n"
+"\n"
+"		if (collidables[collidableIndexA].m_shapeType == SHAPE_CONVEX_HULL &&\n"
+"			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
+"		{\n"
+"		\n"
+"			computeContactSphereConvex(pairIndex, bodyIndexB, bodyIndexA, collidableIndexB, collidableIndexA, \n"
+"																rigidBodies,collidables,convexShapes,vertices,indices,faces, globalContactsOut, nGlobalContactsOut,numPairs);\n"
+"			return;\n"
+"		}\n"
+"	\n"
+"	\n"
+"	\n"
+"	\n"
+"	\n"
+"		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
+"			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
+"		{\n"
+"			//sphere-sphere\n"
+"			float radiusA = collidables[collidableIndexA].m_radius;\n"
+"			float radiusB = collidables[collidableIndexB].m_radius;\n"
+"			float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
+"			float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
+"\n"
+"			float4 diff = posA-posB;\n"
+"			float len = length(diff);\n"
+"			\n"
+"			///iff distance positive, don't generate a new contact\n"
+"			if ( len <= (radiusA+radiusB))\n"
+"			{\n"
+"				///distance (negative means penetration)\n"
+"				float dist = len - (radiusA+radiusB);\n"
+"				float4 normalOnSurfaceB = make_float4(1.f,0.f,0.f,0.f);\n"
+"				if (len > 0.00001)\n"
+"				{\n"
+"					normalOnSurfaceB = diff / len;\n"
+"				}\n"
+"				float4 contactPosB = posB + normalOnSurfaceB*radiusB;\n"
+"				contactPosB.w = dist;\n"
+"								\n"
+"				int dstIdx;\n"
+"        AppendInc( nGlobalContactsOut, dstIdx );\n"
+"				\n"
+"				if (dstIdx < numPairs)\n"
+"				{\n"
+"					__global Contact4* c = &globalContactsOut[dstIdx];\n"
+"					c->m_worldNormal = -normalOnSurfaceB;\n"
+"					c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
+"					c->m_batchIdx = pairIndex;\n"
+"					int bodyA = pairs[pairIndex].x;\n"
+"					int bodyB = pairs[pairIndex].y;\n"
+"					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
+"					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
+"					c->m_worldPos[0] = contactPosB;\n"
+"					GET_NPOINTS(*c) = 1;\n"
+"				}//if (dstIdx < numPairs)\n"
+"			}//if ( len <= (radiusA+radiusB))\n"
+"\n"
+"			return;\n"
+"		}//SHAPE_SPHERE SHAPE_SPHERE\n"
+"\n"
+"	}//	if (i<numPairs)\n"
+"\n"
+"}\n"
+"\n"
+;

opencl/gpu_sat/kernels/satClipHullContacts.cl

@@ -1,9 +1,13 @@
 
 #define TRIANGLE_NUM_CONVEX_FACES 5
 
+#define SHAPE_CONVEX_HULL 3
 #define SHAPE_PLANE 4
+#define SHAPE_CONCAVE_TRIMESH 5
+#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6
 #define SHAPE_SPHERE 7
 
+
 #pragma OPENCL EXTENSION cl_amd_printf : enable
 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
@@ -41,12 +45,8 @@ typedef struct
 {
 	float4 m_worldPos[4];
 	float4 m_worldNormal;	//	w: m_nPoints
-//	float m_restituitionCoeff;
-//	float m_frictionCoeff;
 	u32 m_coeffs;
 	u32 m_batchIdx;
-//	int m_nPoints;
-//	int m_padding0;
 
 	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr
 	int m_bodyBPtrAndSignBit;
@@ -872,11 +872,6 @@ int extractManifoldSequential(const float4* p, int nPoints, float4 nearNormal, i
 		contactIdx[2] = idx[2];
 		contactIdx[3] = idx[3];
 
-//		if( max00.y < 0.0f )
-//			contactIdx[0] = (int)max00.x;
-
-		//does this sort happen on GPU too?
-		//std::sort( contactIdx, contactIdx+4 );
 
 		return 4;
 	}
@@ -908,7 +903,7 @@ __kernel void   extractManifoldAndAddContactKernel(__global const int2* pairs,
 		{
 			localPoints[i] = pointsIn[i];
 		}
-//		int contactIdx[4] = {-1,-1,-1,-1};
+
 		int contactIdx[4];// = {-1,-1,-1,-1};
 		contactIdx[0] = -1;
 		contactIdx[1] = -1;
@@ -954,66 +949,7 @@ void	trMul(float4 translationA, Quaternion orientationA,
 	*translationOut = transform(&translationB,&translationA,&orientationA);
 }
 
-void	computeContactPlaneConvex(int pairIndex,
-																int bodyIndexA, int bodyIndexB, 
-																int collidableIndexA, int collidableIndexB, 
-																__global const BodyData* rigidBodies, 
-																__global const btCollidableGpu* collidables,
-																__global const btGpuFace* faces,
-																__global Contact4* restrict globalContactsOut,
-																counter32_t nGlobalContactsOut,
-																int numPairs)
-{
-	float4 planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;
-	float radius = collidables[collidableIndexB].m_radius;
-	float4 posA1 = rigidBodies[bodyIndexA].m_pos;
-	float4 ornA1 = rigidBodies[bodyIndexA].m_quat;
-	float4 posB1 = rigidBodies[bodyIndexB].m_pos;
-	float4 ornB1 = rigidBodies[bodyIndexB].m_quat;
 
-	bool hasCollision = false;
-	float4 planeNormal1 = make_float4(planeEq.x,planeEq.y,planeEq.z,0.f);
-	float planeConstant = planeEq.w;
-	float4 convexInPlaneTransPos1; Quaternion convexInPlaneTransOrn1;
-	{
-		float4 invPosA;Quaternion invOrnA;
-		trInverse(posA1,ornA1,&invPosA,&invOrnA);
-		trMul(invPosA,invOrnA,posB1,ornB1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);
-	}
-	float4 planeInConvexPos1;	Quaternion planeInConvexOrn1;
-	{
-		float4 invPosB;Quaternion invOrnB;
-		trInverse(posB1,ornB1,&invPosB,&invOrnB);
-		trMul(invPosB,invOrnB,posA1,ornA1,&planeInConvexPos1,&planeInConvexOrn1);	
-	}
-	float4 vtx1 = qtRotate(planeInConvexOrn1,-planeNormal1)*radius;
-	float4 vtxInPlane1 = transform(&vtx1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);
-	float distance = dot3F4(planeNormal1,vtxInPlane1) - planeConstant;
-	hasCollision = distance < 0.f;//m_manifoldPtr->getContactBreakingThreshold();
-	if (hasCollision)
-	{
-		float4 vtxInPlaneProjected1 = vtxInPlane1 -   distance*planeNormal1;
-		float4 vtxInPlaneWorld1 = transform(&vtxInPlaneProjected1,&posA1,&ornA1);
-		float4 normalOnSurfaceB1 = qtRotate(ornA1,planeNormal1);
-		float4 pOnB1 = vtxInPlaneWorld1+normalOnSurfaceB1*distance;
-		pOnB1.w = distance;
-
-		int dstIdx;
-    AppendInc( nGlobalContactsOut, dstIdx );
-		
-		if (dstIdx < numPairs)
-		{
-			__global Contact4* c = &globalContactsOut[dstIdx];
-			c->m_worldNormal = normalOnSurfaceB1;
-			c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
-			c->m_batchIdx = pairIndex;
-			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;
-			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;
-			c->m_worldPos[0] = pOnB1;
-			GET_NPOINTS(*c) = 1;
-		}//if (dstIdx < numPairs)
-	}//if (hasCollision)
-}
 
 
 __kernel void   clipHullHullKernel( __global const int2* pairs, 
@@ -1053,74 +989,6 @@ __kernel void   clipHullHullKernel( __global const int2* pairs,
 		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;
 		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;
 
-
-		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&
-			collidables[collidableIndexB].m_shapeType == SHAPE_PLANE)
-		{
-
-
-			computeContactPlaneConvex( pairIndex, bodyIndexB,bodyIndexA,  collidableIndexB,collidableIndexA, 
-																rigidBodies,collidables,faces,	globalContactsOut, nGlobalContactsOut,numPairs);
-			return;
-		}
-
-		if (collidables[collidableIndexA].m_shapeType == SHAPE_PLANE &&
-			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
-		{
-
-
-			computeContactPlaneConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, 
-																rigidBodies,collidables,faces, globalContactsOut, nGlobalContactsOut,numPairs);
-			return;
-			
-		}
-	
-		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&
-			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)
-		{
-			//sphere-sphere
-			float radiusA = collidables[collidableIndexA].m_radius;
-			float radiusB = collidables[collidableIndexB].m_radius;
-			float4 posA = rigidBodies[bodyIndexA].m_pos;
-			float4 posB = rigidBodies[bodyIndexB].m_pos;
-
-			float4 diff = posA-posB;
-			float len = length(diff);
-			
-			///iff distance positive, don't generate a new contact
-			if ( len <= (radiusA+radiusB))
-			{
-				///distance (negative means penetration)
-				float dist = len - (radiusA+radiusB);
-				float4 normalOnSurfaceB = make_float4(1.f,0.f,0.f,0.f);
-				if (len > 0.00001)
-				{
-					normalOnSurfaceB = diff / len;
-				}
-				float4 contactPosB = posB + normalOnSurfaceB*radiusB;
-				contactPosB.w = dist;
-								
-				int dstIdx;
-        AppendInc( nGlobalContactsOut, dstIdx );
-				
-				if (dstIdx < numPairs)
-				{
-					__global Contact4* c = &globalContactsOut[dstIdx];
-					c->m_worldNormal = -normalOnSurfaceB;
-					c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);
-					c->m_batchIdx = pairIndex;
-					int bodyA = pairs[pairIndex].x;
-					int bodyB = pairs[pairIndex].y;
-					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;
-					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;
-					c->m_worldPos[0] = contactPosB;
-					GET_NPOINTS(*c) = 1;
-				}//if (dstIdx < numPairs)
-			}//if ( len <= (radiusA+radiusB))
-
-			return;
-		}//SHAPE_SPHERE SHAPE_SPHERE
-
 		if (hasSeparatingAxis[i])
 		{
 
@@ -1262,10 +1130,6 @@ __kernel void   clipCompoundsHullHullKernel( __global const int4* gpuCompoundPai
 			int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;
 			int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;
 		
-			
-
-
-		
 			int numLocalContactsOut = clipHullAgainstHull(gpuCompoundSepNormalsOut[i],
 														&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],
 														posA,ornA,

opencl/gpu_sat/kernels/satClipHullContacts.h

@@ -3,9 +3,13 @@ static const char* satClipKernelsCL= \
 "\n"
 "#define TRIANGLE_NUM_CONVEX_FACES 5\n"
 "\n"
+"#define SHAPE_CONVEX_HULL 3\n"
 "#define SHAPE_PLANE 4\n"
+"#define SHAPE_CONCAVE_TRIMESH 5\n"
+"#define SHAPE_COMPOUND_OF_CONVEX_HULLS 6\n"
 "#define SHAPE_SPHERE 7\n"
 "\n"
+"\n"
 "#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
 "#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n"
 "#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n"
@@ -43,12 +47,8 @@ static const char* satClipKernelsCL= \
 "{\n"
 "	float4 m_worldPos[4];\n"
 "	float4 m_worldNormal;	//	w: m_nPoints\n"
-"//	float m_restituitionCoeff;\n"
-"//	float m_frictionCoeff;\n"
 "	u32 m_coeffs;\n"
 "	u32 m_batchIdx;\n"
-"//	int m_nPoints;\n"
-"//	int m_padding0;\n"
 "\n"
 "	int m_bodyAPtrAndSignBit;//x:m_bodyAPtr, y:m_bodyBPtr\n"
 "	int m_bodyBPtrAndSignBit;\n"
@@ -874,11 +874,6 @@ static const char* satClipKernelsCL= \
 "		contactIdx[2] = idx[2];\n"
 "		contactIdx[3] = idx[3];\n"
 "\n"
-"//		if( max00.y < 0.0f )\n"
-"//			contactIdx[0] = (int)max00.x;\n"
-"\n"
-"		//does this sort happen on GPU too?\n"
-"		//std::sort( contactIdx, contactIdx+4 );\n"
 "\n"
 "		return 4;\n"
 "	}\n"
@@ -910,7 +905,7 @@ static const char* satClipKernelsCL= \
 "		{\n"
 "			localPoints[i] = pointsIn[i];\n"
 "		}\n"
-"//		int contactIdx[4] = {-1,-1,-1,-1};\n"
+"\n"
 "		int contactIdx[4];// = {-1,-1,-1,-1};\n"
 "		contactIdx[0] = -1;\n"
 "		contactIdx[1] = -1;\n"
@@ -956,66 +951,7 @@ static const char* satClipKernelsCL= \
 "	*translationOut = transform(&translationB,&translationA,&orientationA);\n"
 "}\n"
 "\n"
-"void	computeContactPlaneConvex(int pairIndex,\n"
-"																int bodyIndexA, int bodyIndexB, \n"
-"																int collidableIndexA, int collidableIndexB, \n"
-"																__global const BodyData* rigidBodies, \n"
-"																__global const btCollidableGpu* collidables,\n"
-"																__global const btGpuFace* faces,\n"
-"																__global Contact4* restrict globalContactsOut,\n"
-"																counter32_t nGlobalContactsOut,\n"
-"																int numPairs)\n"
-"{\n"
-"	float4 planeEq = faces[collidables[collidableIndexA].m_shapeIndex].m_plane;\n"
-"	float radius = collidables[collidableIndexB].m_radius;\n"
-"	float4 posA1 = rigidBodies[bodyIndexA].m_pos;\n"
-"	float4 ornA1 = rigidBodies[bodyIndexA].m_quat;\n"
-"	float4 posB1 = rigidBodies[bodyIndexB].m_pos;\n"
-"	float4 ornB1 = rigidBodies[bodyIndexB].m_quat;\n"
-"	\n"
-"	bool hasCollision = false;\n"
-"	float4 planeNormal1 = make_float4(planeEq.x,planeEq.y,planeEq.z,0.f);\n"
-"	float planeConstant = planeEq.w;\n"
-"	float4 convexInPlaneTransPos1; Quaternion convexInPlaneTransOrn1;\n"
-"	{\n"
-"		float4 invPosA;Quaternion invOrnA;\n"
-"		trInverse(posA1,ornA1,&invPosA,&invOrnA);\n"
-"		trMul(invPosA,invOrnA,posB1,ornB1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);\n"
-"	}\n"
-"	float4 planeInConvexPos1;	Quaternion planeInConvexOrn1;\n"
-"	{\n"
-"		float4 invPosB;Quaternion invOrnB;\n"
-"		trInverse(posB1,ornB1,&invPosB,&invOrnB);\n"
-"		trMul(invPosB,invOrnB,posA1,ornA1,&planeInConvexPos1,&planeInConvexOrn1);	\n"
-"	}\n"
-"	float4 vtx1 = qtRotate(planeInConvexOrn1,-planeNormal1)*radius;\n"
-"	float4 vtxInPlane1 = transform(&vtx1,&convexInPlaneTransPos1,&convexInPlaneTransOrn1);\n"
-"	float distance = dot3F4(planeNormal1,vtxInPlane1) - planeConstant;\n"
-"	hasCollision = distance < 0.f;//m_manifoldPtr->getContactBreakingThreshold();\n"
-"	if (hasCollision)\n"
-"	{\n"
-"		float4 vtxInPlaneProjected1 = vtxInPlane1 -   distance*planeNormal1;\n"
-"		float4 vtxInPlaneWorld1 = transform(&vtxInPlaneProjected1,&posA1,&ornA1);\n"
-"		float4 normalOnSurfaceB1 = qtRotate(ornA1,planeNormal1);\n"
-"		float4 pOnB1 = vtxInPlaneWorld1+normalOnSurfaceB1*distance;\n"
-"		pOnB1.w = distance;\n"
 "\n"
-"		int dstIdx;\n"
-"    AppendInc( nGlobalContactsOut, dstIdx );\n"
-"		\n"
-"		if (dstIdx < numPairs)\n"
-"		{\n"
-"			__global Contact4* c = &globalContactsOut[dstIdx];\n"
-"			c->m_worldNormal = normalOnSurfaceB1;\n"
-"			c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
-"			c->m_batchIdx = pairIndex;\n"
-"			c->m_bodyAPtrAndSignBit = rigidBodies[bodyIndexA].m_invMass==0?-bodyIndexA:bodyIndexA;\n"
-"			c->m_bodyBPtrAndSignBit = rigidBodies[bodyIndexB].m_invMass==0?-bodyIndexB:bodyIndexB;\n"
-"			c->m_worldPos[0] = pOnB1;\n"
-"			GET_NPOINTS(*c) = 1;\n"
-"		}//if (dstIdx < numPairs)\n"
-"	}//if (hasCollision)\n"
-"}\n"
 "\n"
 "\n"
 "__kernel void   clipHullHullKernel( __global const int2* pairs, \n"
@@ -1054,74 +990,6 @@ static const char* satClipKernelsCL= \
 "			\n"
 "		int collidableIndexA = rigidBodies[bodyIndexA].m_collidableIdx;\n"
 "		int collidableIndexB = rigidBodies[bodyIndexB].m_collidableIdx;\n"
-"	\n"
-"\n"
-"		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
-"			collidables[collidableIndexB].m_shapeType == SHAPE_PLANE)\n"
-"		{\n"
-"\n"
-"\n"
-"			computeContactPlaneConvex( pairIndex, bodyIndexB,bodyIndexA,  collidableIndexB,collidableIndexA, \n"
-"																rigidBodies,collidables,faces,	globalContactsOut, nGlobalContactsOut,numPairs);\n"
-"			return;\n"
-"		}\n"
-"\n"
-"		if (collidables[collidableIndexA].m_shapeType == SHAPE_PLANE &&\n"
-"			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
-"		{\n"
-"\n"
-"\n"
-"			computeContactPlaneConvex(pairIndex, bodyIndexA, bodyIndexB, collidableIndexA, collidableIndexB, \n"
-"																rigidBodies,collidables,faces, globalContactsOut, nGlobalContactsOut,numPairs);\n"
-"			return;\n"
-"			\n"
-"		}\n"
-"	\n"
-"		if (collidables[collidableIndexA].m_shapeType == SHAPE_SPHERE &&\n"
-"			collidables[collidableIndexB].m_shapeType == SHAPE_SPHERE)\n"
-"		{\n"
-"			//sphere-sphere\n"
-"			float radiusA = collidables[collidableIndexA].m_radius;\n"
-"			float radiusB = collidables[collidableIndexB].m_radius;\n"
-"			float4 posA = rigidBodies[bodyIndexA].m_pos;\n"
-"			float4 posB = rigidBodies[bodyIndexB].m_pos;\n"
-"\n"
-"			float4 diff = posA-posB;\n"
-"			float len = length(diff);\n"
-"			\n"
-"			///iff distance positive, don't generate a new contact\n"
-"			if ( len <= (radiusA+radiusB))\n"
-"			{\n"
-"				///distance (negative means penetration)\n"
-"				float dist = len - (radiusA+radiusB);\n"
-"				float4 normalOnSurfaceB = make_float4(1.f,0.f,0.f,0.f);\n"
-"				if (len > 0.00001)\n"
-"				{\n"
-"					normalOnSurfaceB = diff / len;\n"
-"				}\n"
-"				float4 contactPosB = posB + normalOnSurfaceB*radiusB;\n"
-"				contactPosB.w = dist;\n"
-"								\n"
-"				int dstIdx;\n"
-"        AppendInc( nGlobalContactsOut, dstIdx );\n"
-"				\n"
-"				if (dstIdx < numPairs)\n"
-"				{\n"
-"					__global Contact4* c = &globalContactsOut[dstIdx];\n"
-"					c->m_worldNormal = -normalOnSurfaceB;\n"
-"					c->m_coeffs = (u32)(0.f*0xffff) | ((u32)(0.7f*0xffff)<<16);\n"
-"					c->m_batchIdx = pairIndex;\n"
-"					int bodyA = pairs[pairIndex].x;\n"
-"					int bodyB = pairs[pairIndex].y;\n"
-"					c->m_bodyAPtrAndSignBit = rigidBodies[bodyA].m_invMass==0?-bodyA:bodyA;\n"
-"					c->m_bodyBPtrAndSignBit = rigidBodies[bodyB].m_invMass==0?-bodyB:bodyB;\n"
-"					c->m_worldPos[0] = contactPosB;\n"
-"					GET_NPOINTS(*c) = 1;\n"
-"				}//if (dstIdx < numPairs)\n"
-"			}//if ( len <= (radiusA+radiusB))\n"
-"\n"
-"			return;\n"
-"		}//SHAPE_SPHERE SHAPE_SPHERE\n"
 "\n"
 "		if (hasSeparatingAxis[i])\n"
 "		{\n"
@@ -1263,10 +1131,6 @@ static const char* satClipKernelsCL= \
 "			\n"
 "			int shapeIndexA = collidables[collidableIndexA].m_shapeIndex;\n"
 "			int shapeIndexB = collidables[collidableIndexB].m_shapeIndex;\n"
-"\n"
-"			\n"
-"\n"
-"\n"
 "		\n"
 "			int numLocalContactsOut = clipHullAgainstHull(gpuCompoundSepNormalsOut[i],\n"
 "														&convexShapes[shapeIndexA], &convexShapes[shapeIndexB],\n"