Last active
December 11, 2023 21:50
-
-
Save lucasdamianjohnson/d5e0dc45071979f5ab03334a1c54dc3a to your computer and use it in GitHub Desktop.
voxel world gen on gpu
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// NOTE!: vec3u is padded to by 4 bytes | |
@group(0) @binding(0) var<storage, read_write> worldData: array<vec2u>; | |
@group(0) @binding(1) var<uniform> processingOptions: vec3<f32>; | |
fn generate_column(column_position: vec3<f32>) { | |
//put code here | |
let minY: f32 = 10; | |
var chunk_type: f32 = 1; | |
if(column_position.x % 2 == 1 || column_position.z % 2 == 1) { | |
chunk_type = 0; | |
} | |
for (var x = column_position.x; x < column_position.x + voxel_world.column_size.x; x += 1.) { | |
for (var z = column_position.z; z < column_position.z + voxel_world.column_size.z; z += 1.) { | |
for (var y = 0.; y < voxel_world.column_size.y; y += 1.) { | |
if(y > 120) { | |
continue; | |
} | |
var place_voxel : bool = false; | |
var r = (perlinNoise3(vec3f(x / 20,y / 20,z / 20))) * .2; | |
if(y > 30) { | |
var n = (perlinNoise3(vec3f(x / 100,y / 100,z / 100))) * 50 + 129; | |
if((r > .1 && r < .8) && y <= n) { | |
place_voxel = true; | |
} | |
} | |
if(y <= 30){ | |
var n = (perlinNoise3(vec3f(x / 50,y / 100,z / 200))) * 50; | |
if((r > .1 && r < .8) || y <= n) { | |
place_voxel = true; | |
} | |
} | |
if(y == 0) { | |
place_voxel = true; | |
} | |
if(place_voxel && y <= 40) { | |
set_voxel( | |
vec3(x,y,z), | |
Voxel( | |
21, | |
0, | |
Light(0,0,0,0), | |
State(0,0,0) | |
) | |
); | |
} | |
if(place_voxel && y > 40) { | |
if(noise3(vec3f(x,y,z)) < .5) { | |
set_voxel( | |
vec3(x,y,z), | |
Voxel( | |
21, | |
0, | |
Light(0,0,0,0), | |
State(0,0,0) | |
) | |
); | |
} else { | |
set_voxel( | |
vec3(x,y,z), | |
Voxel( | |
25, | |
0, | |
Light(0,15,0,15), | |
State(0,0,0) | |
) | |
); | |
} | |
} | |
if( y < 40 && !place_voxel) { | |
set_voxel( | |
vec3(x,y,z), | |
Voxel( | |
32, | |
0, | |
Light(0,0,0,0), | |
State(0,0,0) | |
) | |
); | |
} | |
} | |
} | |
} | |
} | |
fn mod289(x: vec4f) -> vec4f { return x - floor(x * (1. / 289.)) * 289.; } | |
fn perm4(x: vec4f) -> vec4f { return mod289(((x * 34.) + 1.) * x); } | |
fn noise3(p: vec3f) -> f32 { | |
let a = floor(p); | |
var d: vec3f = p - a; | |
d = d * d * (3. - 2. * d); | |
let b = a.xxyy + vec4f(0., 1., 0., 1.); | |
let k1 = perm4(b.xyxy); | |
let k2 = perm4(k1.xyxy + b.zzww); | |
let c = k2 + a.zzzz; | |
let k3 = perm4(c); | |
let k4 = perm4(c + 1.); | |
let o1 = fract(k3 * (1. / 41.)); | |
let o2 = fract(k4 * (1. / 41.)); | |
let o3 = o2 * d.z + o1 * (1. - d.z); | |
let o4 = o3.yw * d.x + o3.xz * (1. - d.x); | |
return o4.y * d.y + o4.x * (1. - d.y); | |
} | |
// MIT License. © Stefan Gustavson, Munrocket | |
// | |
fn permute4(x: vec4f) -> vec4f { return ((x * 34. + 1.) * x) % vec4f(289.); } | |
fn taylorInvSqrt4(r: vec4f) -> vec4f { return 1.79284291400159 - 0.85373472095314 * r; } | |
fn fade3(t: vec3f) -> vec3f { return t * t * t * (t * (t * 6. - 15.) + 10.); } | |
fn perlinNoise3(P: vec3f) -> f32 { | |
var Pi0 : vec3f = floor(P); // Integer part for indexing | |
var Pi1 : vec3f = Pi0 + vec3f(1.); // Integer part + 1 | |
Pi0 = Pi0 % vec3f(289.); | |
Pi1 = Pi1 % vec3f(289.); | |
let Pf0 = fract(P); // Fractional part for interpolation | |
let Pf1 = Pf0 - vec3f(1.); // Fractional part - 1. | |
let ix = vec4f(Pi0.x, Pi1.x, Pi0.x, Pi1.x); | |
let iy = vec4f(Pi0.yy, Pi1.yy); | |
let iz0 = Pi0.zzzz; | |
let iz1 = Pi1.zzzz; | |
let ixy = permute4(permute4(ix) + iy); | |
let ixy0 = permute4(ixy + iz0); | |
let ixy1 = permute4(ixy + iz1); | |
var gx0: vec4f = ixy0 / 7.; | |
var gy0: vec4f = fract(floor(gx0) / 7.) - 0.5; | |
gx0 = fract(gx0); | |
var gz0: vec4f = vec4f(0.5) - abs(gx0) - abs(gy0); | |
var sz0: vec4f = step(gz0, vec4f(0.)); | |
gx0 = gx0 + sz0 * (step(vec4f(0.), gx0) - 0.5); | |
gy0 = gy0 + sz0 * (step(vec4f(0.), gy0) - 0.5); | |
var gx1: vec4f = ixy1 / 7.; | |
var gy1: vec4f = fract(floor(gx1) / 7.) - 0.5; | |
gx1 = fract(gx1); | |
var gz1: vec4f = vec4f(0.5) - abs(gx1) - abs(gy1); | |
var sz1: vec4f = step(gz1, vec4f(0.)); | |
gx1 = gx1 - sz1 * (step(vec4f(0.), gx1) - 0.5); | |
gy1 = gy1 - sz1 * (step(vec4f(0.), gy1) - 0.5); | |
var g000: vec3f = vec3f(gx0.x, gy0.x, gz0.x); | |
var g100: vec3f = vec3f(gx0.y, gy0.y, gz0.y); | |
var g010: vec3f = vec3f(gx0.z, gy0.z, gz0.z); | |
var g110: vec3f = vec3f(gx0.w, gy0.w, gz0.w); | |
var g001: vec3f = vec3f(gx1.x, gy1.x, gz1.x); | |
var g101: vec3f = vec3f(gx1.y, gy1.y, gz1.y); | |
var g011: vec3f = vec3f(gx1.z, gy1.z, gz1.z); | |
var g111: vec3f = vec3f(gx1.w, gy1.w, gz1.w); | |
let norm0 = taylorInvSqrt4( | |
vec4f(dot(g000, g000), dot(g010, g010), dot(g100, g100), dot(g110, g110))); | |
g000 = g000 * norm0.x; | |
g010 = g010 * norm0.y; | |
g100 = g100 * norm0.z; | |
g110 = g110 * norm0.w; | |
let norm1 = taylorInvSqrt4( | |
vec4f(dot(g001, g001), dot(g011, g011), dot(g101, g101), dot(g111, g111))); | |
g001 = g001 * norm1.x; | |
g011 = g011 * norm1.y; | |
g101 = g101 * norm1.z; | |
g111 = g111 * norm1.w; | |
let n000 = dot(g000, Pf0); | |
let n100 = dot(g100, vec3f(Pf1.x, Pf0.yz)); | |
let n010 = dot(g010, vec3f(Pf0.x, Pf1.y, Pf0.z)); | |
let n110 = dot(g110, vec3f(Pf1.xy, Pf0.z)); | |
let n001 = dot(g001, vec3f(Pf0.xy, Pf1.z)); | |
let n101 = dot(g101, vec3f(Pf1.x, Pf0.y, Pf1.z)); | |
let n011 = dot(g011, vec3f(Pf0.x, Pf1.yz)); | |
let n111 = dot(g111, Pf1); | |
var fade_xyz: vec3f = fade3(Pf0); | |
let temp = vec4f(f32(fade_xyz.z)); // simplify after chrome bug fix | |
let n_z = mix(vec4f(n000, n100, n010, n110), vec4f(n001, n101, n011, n111), temp); | |
let n_yz = mix(n_z.xy, n_z.zw, vec2f(f32(fade_xyz.y))); // simplify after chrome bug fix | |
let n_xyz = mix(n_yz.x, n_yz.y, fade_xyz.x); | |
return 2.2 * n_xyz; | |
} | |
//math | |
fn get_position_from_index(index: f32, bounds: vec3<f32>) -> vec3<f32> { | |
return vec3<f32>( | |
floor(index % bounds.y), | |
floor((index / bounds.y) % bounds.x), | |
floor(index / (bounds.x * bounds.z)) | |
); | |
} | |
fn get_index_from_position(position: vec3<f32>, bounds: vec3<f32>) -> f32 { | |
return position.z + position.x * bounds.z + position.y * bounds.z * bounds.x; | |
} | |
//https://stackoverflow.com/questions/12964279/whats-the-origin-of-this-glsl-rand-one-liner | |
fn rand(co : vec2<f32>) -> f32 { | |
return fract(sin(dot(co, vec2(12.9898, 78.233))) * 43758.5453); | |
} | |
//world | |
struct VoxelWorld { | |
process_size: vec3<f32>, | |
region_size: vec3<f32>, | |
column_size: vec3<f32>, | |
chunk_size: vec3<f32> | |
} | |
const voxel_world = VoxelWorld( | |
vec3<f32>(256,128,256), | |
vec3<f32>(256,128,256), | |
vec3<f32>(16,128,16), | |
vec3<f32>(16,16,16) | |
); | |
fn get_voxel_position_from_index(index: f32) -> vec3<f32> { | |
return get_position_from_index(index,voxel_world.process_size); | |
} | |
fn get_voxel_index_from_position(position: vec3<f32>) -> f32 { | |
return get_index_from_position(position, voxel_world.process_size); | |
} | |
fn is_in_bounds(position: vec3<f32>) -> bool { | |
if(position.x < 0 || position.x > 256){ return false;} | |
if(position.y < 0 || position.y > 128){ return false;} | |
if(position.z < 0 || position.z > 256){ return false;} | |
return true; | |
} | |
//light | |
struct Light { | |
sun: u32, | |
red: u32, | |
green: u32, | |
blue: u32, | |
} | |
const light_mask: u32 = 15; | |
fn decode_light(data: u32) -> Light { | |
return Light( | |
light_mask & data, | |
((light_mask << 4) & data) >> 4, | |
((light_mask << 8) & data) >> 8, | |
((light_mask << 12) & data) >> 12 | |
); | |
} | |
fn encode_light(light: Light) -> u32 { | |
var data: u32 = 0; | |
data = (data & ~light_mask) | light.sun; | |
data = (data & ~(light_mask << 4)) | (light.red << 4); | |
data = (data & ~(light_mask << 8)) | (light.green << 8); | |
data = (data & ~(light_mask << 12)) | (light.blue << 12); | |
return data; | |
} | |
//state | |
const level_mask: u32 = 15; | |
const level_state_mask: u32 = 48; | |
const shape_state_mask: u32 = 65472; | |
struct State { | |
level: u32, | |
level_state: u32, | |
shape: u32, | |
} | |
fn decode_state(data: u32) -> State { | |
return State( | |
data & level_mask, | |
(data & level_state_mask) >> 4, | |
(data & shape_state_mask) >> 6 | |
); | |
} | |
fn encode_state(state: State) -> u32 { | |
var data: u32 = 0; | |
data = (data & ~level_mask) | state.level; | |
data = (data & ~level_state_mask) | (state.level_state << 4); | |
return (data & ~shape_state_mask) | (state.shape << 6); | |
} | |
//voxels | |
struct Voxel { | |
id: u32, | |
secondary_id: u32, | |
light: Light, | |
state: State | |
} | |
struct VoxelSegments { | |
id: u32, | |
light: u32, | |
state: u32, | |
secondary_id: u32, | |
} | |
const segment_mask: u32 = 65535; | |
fn get_voxel_segments(data: vec2u) -> VoxelSegments { | |
let raw_seg1 = data.x; | |
let raw_seg2 = data.y; | |
return VoxelSegments( | |
data.x & segment_mask, | |
(data.x & (segment_mask << 16)) >> 16, | |
data.y & segment_mask, | |
(data.y & (segment_mask << 16)) >> 16, | |
); | |
} | |
fn encode_voxel_segments(segments: VoxelSegments) -> vec2u { | |
var raw = vec2u(0,0); | |
raw.x = segments.id; | |
raw.x = (raw.x & ~(segment_mask << 16)) | (segments.light << 16); | |
raw.y = segments.state ; | |
raw.y = (raw.y & ~(segment_mask << 16)) | (segments.secondary_id << 16); | |
return raw; | |
} | |
fn voxel_to_segments(voxel: Voxel) -> VoxelSegments { | |
return VoxelSegments( | |
voxel.id, | |
encode_light(voxel.light), | |
encode_state(voxel.state), | |
voxel.secondary_id | |
); | |
} | |
fn voxel_to_raw_data(voxel: Voxel) -> vec2u { | |
return encode_voxel_segments( | |
voxel_to_segments(voxel) | |
); | |
} | |
fn decode_voxel_segments(voxel_segments: VoxelSegments) -> Voxel { | |
//state | |
let state: u32 = 0; | |
//level | |
let level: u32 = 0; | |
let level_state: u32 = 0; | |
return Voxel( | |
voxel_segments.id, | |
voxel_segments.secondary_id, | |
decode_light(voxel_segments.light), | |
decode_state(voxel_segments.state) | |
); | |
} | |
fn get_voxel(position: vec3<f32>) -> Voxel { | |
let voxel_segments = get_voxel_segments(worldData[i32( | |
get_voxel_index_from_position(position) | |
)]); | |
return decode_voxel_segments(voxel_segments); | |
} | |
fn set_voxel(position: vec3<f32>, voxel: Voxel) { | |
let voxel_index = get_voxel_index_from_position(position); | |
worldData[i32(voxel_index)] = voxel_to_raw_data(voxel); | |
} | |
const sun_light_proagation_amount: u32 = 2; | |
const sun_light_start: f32 = 128; | |
fn can_add_light(voxel:Voxel)-> bool { | |
if(voxel.light.red > 0 || voxel.light.green > 0 || voxel.light.blue > 0) {return true;} | |
if(voxel.id <= 2) { return true;} | |
return false; | |
} | |
fn run_sun_light_fill(voxel_position: vec3<f32>) { | |
var main_voxel = get_voxel(voxel_position); | |
var position : vec3<f32> = vec3(voxel_position.x,voxel_position.y ,voxel_position.z); | |
if(voxel_position.y >= sun_light_start) { | |
main_voxel.light.sun = 15; | |
set_voxel(voxel_position,main_voxel); | |
for(var y = voxel_position.y - 1; y > 0; y -= 1.) { | |
position = vec3(voxel_position.x,y,voxel_position.z); | |
var other_voxel = get_voxel(position); | |
if(can_add_light(other_voxel)) { | |
other_voxel.light.sun = get_minus_one_for_sun_under_voxel(main_voxel,other_voxel); | |
set_voxel(position,other_voxel); | |
} else { | |
break; | |
} | |
} | |
} | |
} | |
const directions = array<vec3<f32>, 6>( | |
vec3<f32>(1,0,0) | |
,vec3<f32>(-1,0,0) | |
,vec3<f32>(0,0,1) | |
,vec3<f32>(0,0,-1) | |
,vec3<f32>(0,1,0) | |
,vec3<f32>(0,-1,0) | |
); | |
const rgb_light_proagation_amount = 1; | |
fn run_light_flood(voxel_position: vec3<f32>) -> f32 { | |
var main_voxel = get_voxel(voxel_position); | |
var position : vec3<f32> = vec3(voxel_position.x,voxel_position.y,voxel_position.z); | |
if(!can_add_light(main_voxel)) {return 0;} | |
for(var i = 0; i < 6; i++ ){ | |
var direction = directions[i]; | |
position = vec3( | |
voxel_position.x + direction.x, | |
voxel_position.y + direction.y, | |
voxel_position.z + direction.z | |
); | |
if(is_in_bounds(position)) { | |
var other_voxel = get_voxel(position); | |
if(main_voxel.light.sun + sun_light_proagation_amount < other_voxel.light.sun) { | |
var value = other_voxel.light.sun - sun_light_proagation_amount; | |
if(value < 0) { | |
value = 0; | |
} | |
if(value < main_voxel.light.sun) { | |
value = main_voxel.light.sun; | |
} | |
main_voxel.light.sun = value; | |
} | |
if(main_voxel.light.red + 2 <= other_voxel.light.red) { | |
var value = other_voxel.light.red - rgb_light_proagation_amount; | |
if(value < 0) { | |
value = 0; | |
} | |
if(value < main_voxel.light.red) { | |
value = main_voxel.light.red; | |
} | |
main_voxel.light.red = value; | |
} | |
if(main_voxel.light.green + 2 <= other_voxel.light.green) { | |
var value = other_voxel.light.green - rgb_light_proagation_amount; | |
if(value < 0) { | |
value = 0; | |
} | |
if(value < main_voxel.light.green) { | |
value = main_voxel.light.green; | |
} | |
main_voxel.light.green = value; | |
} | |
if(main_voxel.light.blue + 2 <= other_voxel.light.blue) { | |
var value = other_voxel.light.blue - rgb_light_proagation_amount; | |
if(value < 0) { | |
value = 0; | |
} | |
if(value < main_voxel.light.blue) { | |
value = main_voxel.light.blue; | |
} | |
main_voxel.light.blue = value; | |
} | |
} | |
} | |
set_voxel(voxel_position,main_voxel); | |
return 1; | |
} | |
fn light_flood(column_position: vec3<f32>) -> f32 { | |
for (var x = column_position.x; x < column_position.x + voxel_world.column_size.x; x += 1.) { | |
for (var z = column_position.z; z < column_position.z + voxel_world.column_size.z; z += 1.) { | |
for (var y = 0.; y < voxel_world.column_size.y; y += 1.) { | |
run_light_flood(vec3<f32>(x,y,z)); | |
} | |
} | |
} | |
return 1; | |
} | |
@compute @workgroup_size(1) | |
fn computeSomething( | |
@builtin(workgroup_id) workgroup_id : vec3<u32>, | |
@builtin(local_invocation_id) local_invocation_id : vec3<u32>, | |
@builtin(global_invocation_id) global_invocation_id : vec3<u32>, | |
@builtin(local_invocation_index) local_invocation_index: u32, | |
@builtin(num_workgroups) num_workgroups: vec3<u32> | |
) { | |
if(processingOptions.x == 0){ | |
let column_position = vec3<f32>( | |
f32(workgroup_id.x) * voxel_world.column_size.x, | |
f32(processingOptions.y), | |
f32(workgroup_id.y) * voxel_world.column_size.z, | |
); | |
generate_column(column_position); | |
} | |
if(processingOptions.x == 1){ | |
let voxel_position = vec3<f32>( | |
f32(workgroup_id.x), | |
f32(processingOptions.y), | |
f32(workgroup_id.y) | |
); | |
run_sun_light_fill(voxel_position); | |
} | |
if(processingOptions.x == 2){ | |
let voxel_position = vec3<f32>( | |
f32(workgroup_id.x), | |
f32(processingOptions.y), | |
f32(workgroup_id.y) | |
); | |
for(var i = 0; i < 128; i++) { | |
run_light_flood(vec3<f32>( | |
f32(workgroup_id.x), | |
f32(i), | |
f32(workgroup_id.y) | |
)); | |
} | |
// run_sun_rgb_flood(voxel_position); | |
} | |
} | |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment