Using OpenCL to generate a terrain heightmap on the GPU, using the Diamond-Square algorithm.
Made for an IGAD class on GPU programming.
OpenCL code to generate the heightmap using the Diamond-Square algorithm, by modifying the y position values of a 2D grid of vertices.
This code is called for every pass of the Diamond-Square algorithm: starting with a rectangle covering the whole terrain, then subdividing that into four smaller rectangles, then subdividing each rectangle again, etc. until the rectangles too small (2*2 vertices).
#include "tools/random.h" //rng library
#include "opencl_shared.h" //define INTENSITY, SQUARE_SIZE
__constant const float Intensity = 1.0f * INTENSITY;
__constant const float CIntensity = 0.25f * INTENSITY;
__kernel void task( __global float* vertexBuffer , int rectSize, int randseed )
{
int idx = get_global_id( 0 ); //x index of the current rectangle
int idz = get_global_id( 1 ); //z index of the current rectangle
rngstate rng;
rngseed( &rng, randseed + idx + SQUARE_SIZE * idy, 1 );
int x1 = (idx * (rectSize-1)) % SQUARE_SIZE;
int x2 = (x1 + rectSize - 1) % SQUARE_SIZE;
int z1 = (idz * (rectSize-1)) % SQUARE_SIZE;
int z2 = (z1 + rectSize - 1) % SQUARE_SIZE;
int centerX = (x1 + x2) / 2;
int centerY = (z1 + z2) / 2;
float IntensityModifier = (float)rectSize/SQUARE_SIZE;
//Diamond step
float NorthWest = vertexBuffer[(y1*SQUARE_SIZE+x1)*3+1]; //nw
float NorthEast = vertexBuffer[(y2*SQUARE_SIZE+x2)*3+1]; //ne
float SouthEast = vertexBuffer[(y1*SQUARE_SIZE+x2)*3+1]; //se
float SouthWest = vertexBuffer[(y2*SQUARE_SIZE+x1)*3+1]; //sw
float Center = ((NorthWest + NorthEast + SouthEast + SouthWest)/4.0f) + (-0.5f + rngfloat(&rng)) * CIntensity * IntensityModifier;
//set height of center vertex
vertexBuffer[(centerZ*SQUARE_SIZE+centerX)*3+1] = Center;
//Square step
float North = ((NorthWest+NorthEast)/2.0f) + (-0.5f + rngfloat(&rng)) * Intensity * IntensityModifier;
float South = ((SouthWest+SouthEast)/2.0f) + (-0.5f + rngfloat(&rng)) * Intensity * IntensityModifier;
float West = ((NorthWest+SouthWest)/2.0f) + (-0.5f + rngfloat(&rng)) * Intensity * IntensityModifier;
float East = ((NorthEast+SouthEast)/2.0f) + (-0.5f + rngfloat(&rng)) * Intensity * IntensityModifier;
//set height of vertices
vertexBuffer[(z1*SQUARE_SIZE+centerX)*3+1] = North; //n
vertexBuffer[(z2*SQUARE_SIZE+centerX)*3+1] = South; //s
vertexBuffer[(centerZ*SQUARE_SIZE+x1)*3+1] = West; //w
vertexBuffer[(centerZ*SQUARE_SIZE+x2)*3+1] = East; //e
}