Az eddigi példaprogramok, valljuk be, nem voltak túl látványosak. Joggal kérdezhetnék tőlem, hogy miféle demoscene ember az, aki csak ilyen unalmas példaprogramokat tud bemutatni. Teljes mértékben igazat kell adnom. Ezennek törlesztem adósságom, és egy olyan programot készítünk, ami OpenGL-t is használ. Méghozzá egy csomó pontot fogunk dobálni a képernyőn. Ez nem nagy kaland még a CPU-nak sem, de megbolondítjuk azzal, hogy minden pont ütközhet egymással, nem csak a fallal. Ez már elég érdekesen hangzik? Ha igen, akkor vágjunk bele.
#define GL_GLEXT_PROTOTYPES
#include <stdio.h>
#include <stdlib.h>
#include <GL/glut.h>
#include <GL/gl.h>
#include <cuda_gl_interop.h>
#define WIDTH 1024
#define HEIGHT 768
#define PARTICLENUM 500
#define COLLIDEDIST 0.005f
GLuint pointBuff;
struct cudaGraphicsResource *pointBuff_CUDA;
float4 *particles;
/* Calculate distance of the points */
__device__ float distance(float x1, float y1, float x2, float y2){
return( sqrtf( (x1 - x2) * (x1 - x2) + (y1 - y2) * (y1 - y2)) );
}
__device__ void collision(float *p1x, float *p1y, float *p2x, float *p2y, float angle){
float m1 = sqrtf(*p1x**p1x + *p1y**p1y);
float m2 = sqrtf(*p2x**p2x + *p2y**p2y);
float d1 = atan2f(*p1y, *p1x);
float d2 = atan2f(*p2y, *p2x);
float p1newx = m1 * cos(d1 - angle);
float p1newy = m1 * sin(d1 - angle);
float p2newx = m2 * cos(d2 - angle);
float p2newy = m2 * sin(d2 - angle);
*p1x = cos(angle) * p2newx + cos(angle + 3.14f / 2.0f) * p1newy;
*p1y = sin(angle) * p2newx + sin(angle + 3.14f / 2.0f) * p1newy;
*p2x = cos(angle) * p1newx + cos(angle + 3.14f / 2.0f) * p2newy;
*p2y = sin(angle) * p1newx + sin(angle + 3.14f / 2.0f) * p2newy;
}
__global__ void calcParticlePos(float4 *p){
int index = threadIdx.x;
int i;
/* Collision with the wall */
if(p[index].x > 1.0f){
p[index].x = 1.0f;
p[index].z *= -1.0f;
}
if(p[index].x < -1.0f){
p[index].x = -1.0f;
p[index].z *= -1.0f;
}
if(p[index].y > 1.0f){
p[index].y = 1.0f;
p[index].w *= -1.0f;
}
if(p[index].y < -1.0f){
p[index].y = -1.0f;
p[index].w *= -1.0f;
}
/* Collision with each other */
for(i = 0; i < PARTICLENUM; i++){
if(i == index) continue; /*Prevent self collision */
if(distance(p[i].x, p[i].y, p[index].x, p[index].y) < COLLIDEDIST){
float angle = atan2f(p[i].y - p[index].y, p[i].x - p[index].x);
collision(&p[i].z, &p[i].w, &p[index].z, &p[index].w, angle);
}
}
p[index].x = p[index].x + p[index].z;
p[index].y = p[index].y + p[index].w;
}
/* Handle the keyboard */
void keyboard( unsigned char key, int x, int y){
switch(key){
case 27:
free(particles);
exit(0);
break;
}
}
/* Show the show */
void display(){
float4 *ppos;
cudaGraphicsMapResources(1, &pointBuff_CUDA, 0);
cudaGraphicsResourceGetMappedPointer( (void**)&ppos, NULL, pointBuff_CUDA);
calcParticlePos<<< 1, PARTICLENUM>>>(ppos);
cudaGraphicsUnmapResources(1, &pointBuff_CUDA, 0);
glClear(GL_COLOR_BUFFER_BIT);
glBindBuffer(GL_ARRAY_BUFFER, pointBuff);
glVertexPointer(2, GL_FLOAT, 0, NULL);
glDrawArrays(GL_POINTS, 0, PARTICLENUM);
glutSwapBuffers();
glutPostRedisplay();
}
/* Init a draw buffer and link it to the CUDA system */
void initOpenGL(){
int i;
particles = (float4*)malloc(sizeof(float4) * PARTICLENUM);
for( i = 0; i < PARTICLENUM; i++){
/* Postitions */
particles[i].x = drand48() * 2.0f - 1.0f;
particles[i].y = drand48() * 2.0f - 1.0f;
/* Velocity */
particles[i].z = (drand48() - 0.5f) / 1000.0f;
particles[i].w = (drand48() - 0.5f) / 1000.0f;
}
glGenBuffers(1, &pointBuff);
glBindBuffer(GL_ARRAY_BUFFER, pointBuff);
glBufferData(GL_ARRAY_BUFFER, sizeof(float4) * PARTICLENUM,
particles,
GL_DYNAMIC_DRAW);
cudaGraphicsGLRegisterBuffer(&pointBuff_CUDA, pointBuff,
cudaGraphicsMapFlagsWriteDiscard);
glPointSize(5);
glEnableClientState(GL_VERTEX_ARRAY);
glColor3f(1.0, 0.0, 0.0);
}
int main(int argc, char **argv){
/* Init cuda with opengl */
cudaGLSetGLDevice(0);
/* opegl window creation process */
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutInitWindowSize(WIDTH, HEIGHT);
glutCreateWindow("Collision test in CUDA");
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
initOpenGL();
glutMainLoop();
return(0);
}
A program GLUT-ot használ és OpenGL 2.0-t. Először közöljük a rendszerrel, hogy OpenGL-t is kívánunk használni. Erre szolgál a cudaGLSetGLDevice. Ismételten figyelmeztetek mindenkit, hogy nem vizsgálom le a visszatérési értékeket! A példaprogramban egy vertex buffert (pointBuff) hozunk létre, ami tárolja az ütköző részecskék koordinátáit. A cudaGraphicsGLRegisterBuffer segítségével létrehozunk egy CUDA azonosítót, és összekapcsoljuk az OpenGL pufferét a CUDA-éval. Erre azért van szükség, mert mindkét erőforrás a grafikus kártyán van, nincs szükség a mozgatására. Az érdekes lépések a display függvényben vannak. Először létrehozunk egy mutatót a pufferünkre (cudaGraphicsResourceGetMappedPointer), majd ezt a mutatót használjuk a kernelben. Mikor végeztünk a munkával, a cudaGraphicsUnmapResource segítségével elérhetővé tesszük a puffert az OpenGL számára.
Az ütközés vizsgálat nem saját találmány, azt a Emanuelle Feronato weboldaláról lestem le. Annyi változtatást eszközöltem, hogy nálam egyenlő súlyúak a pontok, ezért az impulzusok számításánál az ütköző részecskék tömegének arányára nincs szükségem. Aki kedvet érez, végezze el a módosításokat.
Ha valaki kipróbálja a fenti kódot, nagyon meg fog lepődni. Néhány helyen rendben lezajlanak az ütközések, de sok helyen a golyók átsüvítenek egymáson! A programban hiba van. A következő részben a hibakeresést és felderítést fogom bemutatni.
