OpenCL Tutorial – Masspoint Movement, Collision Hash Grids (Download)

Der Schwerpunkt des heutigen Tutorials liegt im Entwurf eines einfachen OpenCL-Frameworks sowie in der Implementierung einer einfachen OpenCL-basierten Bewegungssimulation für ein Asteroidenfeld:

  • Simulation der Asteroiden-Bewegung (jedem Asteroid ist eine einzelne Punktmasse zugeordnet).
  • Berechnung der Gitter-Indices und -Positionen in einem Kollisionsgitter (Collision Hash Grid) für eine effiziente Kollisionsprüfung.



Für die Initialisierung einer OpenCL-Anwendung ist die COpenCLDevice-Klasse verantwortlich:

class COpenCLDevice
{
public:

    cl_platform_id   Platform;
    cl_device_id     Device;
    cl_context       DeviceContext;
    cl_command_queue DeviceCommandQueue;

    bool use_GPU_Device;

    COpenCLDevice();
    ~COpenCLDevice();

    void Init_Device(bool useGPU);
    void Wait(void);
};


Zuständig für die Handhabung eines einzelnen OpenCL-Programms (Kernel) ist die COpenCLProgram-Klasse:

class COpenCLProgram
{
public:

    cl_program Program;
    cl_kernel  Kernel;

    cl_uint WorkGroupDim;
    size_t local_ws;
    size_t global_ws;
    size_t MaxBufferElements;

    COpenCLDevice* pUsedOpenCLDevice;

    COpenCLProgram();
    ~COpenCLProgram();

    void Set_OpenCLDevice(COpenCLDevice* pOpenCLDevice);

   
// Laden und Kompilieren eines OpenCl-Programms
    void Build_ProgramFromSourceFile(char *pKernelName, char *pFileName,
         cl_uint workGroupDim, bool fastMath, size_t maxBufferElements,
         size_t num_WorkItems_Per_WorkGroup = 0);

   
// Ausführen eines OpenCL-Programms:
    void Execute_Program(void);
    void Execute_Program_With_UserDefined_Num_WorkItems_Per_WorkGroup(void);

    void Execute_Program(cl_uint numEventsInWaitList, const cl_event
                         *pEventWaitList, cl_event *pEvent);
    void Execute_Program_With_UserDefined_Num_WorkItems_Per_WorkGroup(
         cl_uint numEventsInWaitList, const cl_event *pEventWaitList,
         cl_event *pEvent);

   
// Methoden für die Parameter-Übergabe an ein OpenCL-Programm (Kernel):
    void Set_KernelArg(cl_uint pos, cl_mem* pInput);
    void Set_KernelArg(cl_uint pos, int input);
    void Set_KernelArg(cl_uint pos, double input);
    void Set_KernelArg(cl_uint pos, float input);
    void Set_KernelArg(cl_uint pos, size_t input);
};


Die nachfolgenden Container-Klassen (Arrays von Vektoren und Skalaren) fungieren als Schittstelle zwischen einem OpenCL-Kernel und dem Hauptprogramm:

class COpenCL_FloatScalar_Container;
class COpenCL_2DFloatVector_Container;
class COpenCL_3DFloatVector_Container;
class COpenCL_4DFloatVector_Container;

class COpenCL_IntegerScalar_Container;
class COpenCL_2DIntegerVector_Container;
class COpenCL_3DIntegerVector_Container;
class COpenCL_4DIntegerVector_Container;
class COpenCL_NDIntegerVector_Container;


Die nachfolgenden Klassen ermöglichen den Zugriff auf einzelne Elemente der Container-Klassen aus dem Hauptprogramm heraus:

class COpenCL_FloatScalar;
class COpenCL_2DFloatVector;
class COpenCL_3DFloatVector;
class COpenCL_4DFloatVector;

class COpenCL_IntegerScalar;
class COpenCL_2D IntegerVector;
class COpenCL_3D IntegerVector;
class COpenCL_4D IntegerVector;
class COpenCL_NDIntegerVector;


Verantwortlich für die Durchführung der Bewegungssimulation ist der nachfolgende OpenCL-Kernel:

__kernel void MasspointMovement_GPU(__global float4* Position,
                                    __global float4* Velocity,
                                    __global int4* GridPosition,
                                    __global int4* GridIndex,
                                    const float TimeStep,
                                    const float InvGridSize,
                                    const float GridOffset,
                                    const float GridCenterX,
                                    const float GridCenterY,
                                    const float GridCenterZ,
                          const int HalfNumCollisionGridElementsXDirPlus1,
                          const int HalfNumCollisionGridElementsYDirPlus1,
                          const int HalfNumCollisionGridElementsZDirPlus1,
                          const int HalfNumCollisionGridElementsXDirMinus1,
                          const int HalfNumCollisionGridElementsYDirMinus1,
                          const int HalfNumCollisionGridElementsZDirMinus1)
{
    size_t id = get_global_id(0);

   
// Bewegungs-Simulation (kräftefrei):
    Position[id] += Velocity[id]*TimeStep;

   
// Gitterkoordinaten berechnen:
    GridPosition[id].x = (int)(GridOffset+
                              (Position[id].x-GridCenterX)*InvGridSize);
    GridPosition[id].y = (int)(GridOffset+
                              (Position[id].y-GridCenterY)*InvGridSize);
    GridPosition[id].z = (int)(GridOffset+
                              (Position[id].z-GridCenterZ)*InvGridSize);

   
// Hash-Grid-Index berechnen:
    GridIndex[id].x =
    GridPosition[id].x % HalfNumCollisionGridElementsXDirPlus1 +
    HalfNumCollisionGridElementsXDirMinus1;

    GridIndex[id].y =
    GridPosition[id].y % HalfNumCollisionGridElementsYDirPlus1 +
    HalfNumCollisionGridElementsYDirMinus1;

    GridIndex[id].z =
    GridPosition[id].z % HalfNumCollisionGridElementsZDirPlus1 +
    HalfNumCollisionGridElementsZDirMinus1;
}

Hinweis:
Für die Ausführung dieses Programmbeispiels muss der Treiber Ihrer Grafikkarte die OpenGL Version 3.3 unterstützen.

Visual C++ 2010: DemoWithOpenGL2010_Tut48