Code Device Kernel Sharing

NYC_Trucks · May 25, 2025, 9:59pm

I’ve been scouring the forum and beyond for a hub where people swap and dissect their Code Device host‐scripts and OpenCL kernels—everything from custom noise generators to mask-building utilities—but I’m coming up empty. Is there a corner of the community I’ve missed, or is kernel‐sharing just an under-tapped ritual? Either way, I’d love to kick off a discussion: what holds us back from pooling our best Code‐Device tricks, and how could we build a more vibrant exchange?

Tomatiy · May 25, 2025, 11:32pm

WM community generally not huge due to steep learning curve

and Code Device was presented in October of last year… So new part of community related to kernels coding just not formed yet (i think we need years and years to form such)

i know only couple of them:
@sijmen_v_b
@NuoCheng
@Kr0wn

HYLK · May 27, 2025, 6:45am

The only corner I can think of you may have missed, is the Terraformers Discussion group on Discord. That said, it may be great to have a category on this forum specifically for the Code Device, @Stephen, do you think that is possible/suitable?

Tomatiy · May 27, 2025, 11:04am

Development section already exists
Just extend it description with not only “plugins and macros” but also code device

Stephen · May 29, 2025, 10:35pm

I was definitely intending that the Plugin category would be the place to talk about the Code device, but you’re right, this needs to be marked better. It might be easier to simply archive the old category and create a new one.

NYC_Trucks · June 2, 2025, 3:04pm

–[[
Perlin Noise Implementation
Based on Ken Perlin’s improved noise reference implementation.
MIT License

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.

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.
]]

– Basic Perlin terrain generator using fixed internal permutation table in kernel
local scale = wm.parm(“scale”) or 8.0
local zslice = wm.parm(“zslice”) or 0.0

– Get the device-space extent and create the output buffer
local space = wm.devicespace()
local result = wm.createBuffer()

– Run the kernel with the correct arguments
wm.runKernel(“perlin_gen”, space, result, scale, zslice)

– Output to port 0
wm.output(result, 0)

/*
Perlin Noise Implementation
Based on Ken Perlin’s improved noise reference implementation.
MIT License
Copyright (c) 2025 V. Wade Construction Corp.

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.

*/

/*
Basic Perlin terrain generator using fixed internal permutation table in kernel
*/

inline float fade(float t) {
return t * t * t * (t * (t * 6.0f - 15.0f) + 10.0f);
}

inline float lerp(float a, float b, float t) {
return a + t * (b - a);
}

inline float grad(int hash, float x, float y, float z) {
switch(hash & 15) {
case 0: return x + y;
case 1: return -x + y;
case 2: return x - y;
case 3: return -x - y;
case 4: return x + z;
case 5: return -x + z;
case 6: return x - z;
case 7: return -x - z;
case 8: return y + z;
case 9: return -y + z;
case 10: return y - z;
case 11: return -y - z;
case 12: return x + y;
case 13: return -x + y;
case 14: return -y + z;
case 15: return -y - z;
default: return 0.0f;
}
}

__constant int perm[512] = {
151,160,137,91,90,15,131,13,201,95,96,53,194,233,7,225,
140,36,103,30,69,142,8,99,37,240,21,10,23,190,6,148,247,
120,234,75,0,26,197,62,94,252,219,203,117,35,11,32,57,
177,33,88,237,149,56,87,174,20,125,136,171,168,68,175,
74,165,71,134,139,48,27,166,77,146,158,231,83,111,229,
122,60,211,133,230,220,105,92,41,55,46,245,40,244,102,
143,54,65,25,63,161,1,216,80,73,209,76,132,187,208,89,
18,169,200,196,135,130,116,188,159,86,164,100,109,198,
173,186,3,64,52,217,226,250,124,123,5,202,38,147,118,
126,255,82,85,212,207,206,59,227,47,16,58,17,182,189,
28,42,223,183,170,213,119,248,152,2,44,154,163,70,221,
153,101,155,167,43,172,9,129,22,39,253,19,98,108,110,
79,113,224,232,178,185,112,104,218,246,97,228,251,34,
242,193,238,210,144,12,191,179,162,241,81,51,145,235,
249,14,239,107,49,192,214,31,181,199,106,157,184,84,
204,176,115,121,50,45,127,4,150,254,138,236,205,93,
222,114,67,29,24,72,243,141,128,195,78,66,215,61,156,
180,
// Repeat for wraparound
151,160,137,91,90,15,131,13,201,95,96,53,194,233,7,225,
140,36,103,30,69,142,8,99,37,240,21,10,23,190,6,148,247,
120,234,75,0,26,197,62,94,252,219,203,117,35,11,32,57,
177,33,88,237,149,56,87,174,20,125,136,171,168,68,175,
74,165,71,134,139,48,27,166,77,146,158,231,83,111,229,
122,60,211,133,230,220,105,92,41,55,46,245,40,244,102,
143,54,65,25,63,161,1,216,80,73,209,76,132,187,208,89,
18,169,200,196,135,130,116,188,159,86,164,100,109,198,
173,186,3,64,52,217,226,250,124,123,5,202,38,147,118,
126,255,82,85,212,207,206,59,227,47,16,58,17,182,189,
28,42,223,183,170,213,119,248,152,2,44,154,163,70,221,
153,101,155,167,43,172,9,129,22,39,253,19,98,108,110,
79,113,224,232,178,185,112,104,218,246,97,228,251,34,
242,193,238,210,144,12,191,179,162,241,81,51,145,235,
249,14,239,107,49,192,214,31,181,199,106,157,184,84,
204,176,115,121,50,45,127,4,150,254,138,236,205,93,
222,114,67,29,24,72,243,141,128,195,78,66,215,61,156,
180
};

float perlin3(float x, float y, float z) {
int X = ((int)floor(x)) & 255;
int Y = ((int)floor(y)) & 255;
int Z = ((int)floor(z)) & 255;

x -= floor(x); y -= floor(y); z -= floor(z);
float u = fade(x), v = fade(y), w = fade(z);

int A  = perm[X]   + Y,  AA = perm[A]   + Z,  AB = perm[A+1]   + Z;
int B  = perm[X+1] + Y,  BA = perm[B]   + Z,  BB = perm[B+1]   + Z;

return lerp(
    lerp(
        lerp(grad(perm[AA],   x,   y,   z),   grad(perm[BA],   x-1,   y,   z),   u),
        lerp(grad(perm[AB],   x,   y-1,   z), grad(perm[BB],   x-1,   y-1,   z), u),
        v),
    lerp(
        lerp(grad(perm[AA+1], x,   y,   z-1), grad(perm[BA+1], x-1,   y,   z-1), u),
        lerp(grad(perm[AB+1], x,   y-1, z-1), grad(perm[BB+1], x-1, y-1, z-1), u),
        v),
    w);

}

__kernel void perlin_gen(__read_only image2d_t coords, __global float* out, float scale, float z) {
int2 gid = (int2)(get_global_id(0), get_global_id(1));
float2 uv = read_imagef(coords, gid).xy;

float val = perlin3(uv.x * scale, uv.y * scale, z);
out[gid.y * get_global_size(0) + gid.x] = clamp(val * 0.5f + 0.5f, 0.0f, 1.0f);

}

NYC_Trucks · June 2, 2025, 3:10pm

Overview
This Perlin noise-based kernel represents my very first successful terrain-generating “Code Device” for World Machine. After digging through countless online references to understand how to translate gradient-noise algorithms into a GPU-friendly form, I ported Ken Perlin’s improved noise implementation into OpenCL so that it would compile and execute within World Machine’s sandboxed Code Device environment. Since then, I’ve used this device as a foundational noise source to seed procedural terrain, generate masks, and drive texture placement across a variety of map- and world-design workflows.

Technical Details

Lua Host Script
- Retrieves two user parameters—scale (frequency multiplier) and zslice (third-dimension offset for 3D noise).
- Calls wm.devicespace() to obtain the per-pixel coordinates image.
- Allocates a float buffer via wm.createBuffer() to store the computed noise values.
- Invokes wm.runKernel("perlin_gen", space, result, scale, zslice) in the exact argument order that World Machine expects.
- Outputs the resulting float buffer to port 0 (wm.output(result, 0)), which can then be connected to subsequent devices or exported as a heightfield.
OpenCL Kernel (perlin_gen)
- Reads a single pixel’s device-space XY coordinates from the built-in coords image2d_t.
- Scales those coordinates by the user-specified scale and adds zslice to compute a 3D point (x, y, z).
- Calls perlin3(x, y, z), which:
  - Computes integer lattice cell indices (X, Y, Z).
  - Applies the fade() function to each fractional component.
  - Looks up four 1D gradient contributions per lattice corner via a 512-element perm[]‐based hash.
  - Performs nested lerp() operations to blend those gradients in x, y, and z.
- Normalizes the output to [0, 1] via clamp(val * 0.5f + 0.5f, 0.0f, 1.0f).
- Stores the result into the global float buffer at index gid.y * width + gid.x, where width is get_global_size(0)—ensuring correct row-major ordering.

Usage & Applications

Seeding Procedural Terrain: Connect this device’s output directly to a World Machine heightfield output, then feed it into downstream mask-generators, erosion devices, or texture nodes. The noise forms a base “mountain” shape or subtle undulating ground.
Texture & Mask Generation: Because the output is a smooth, tileable grayscale field, you can sample it at different scales (e.g., lower frequency for large landforms, higher frequency for fine detail) and use it to drive rock/grass blend masks, snow accumulation masks, or placement density maps for trees and props.
Parametric Control: Two parameters—scale and zslice—let you quickly dial in different noise frequencies and animate or layer multiple noise passes by varying zslice over time. For example, stacking several Perlin passes, each at a different zslice, can produce multifractal or FBM (fractal Brownian motion) effects without needing an external noise library.

Why Share on the World Machine Forum?

Proven Foundation
This kernel has been in active use in my own map-design pipelines for months. It compiles reliably, produces visually clean noise, and respects tile boundaries (no seams) when used with WM’s default coordinate system.
Community Reference
Many users struggle to find a fully working Perlin implementation that “just works” in World Machine’s Code Device. By sharing the complete host+kernel with proper parameter hookup and licensing, anyone else can pick it up, tweak scale/zslice, or combine it with their own downstream nodes.
Extensible Example
Because it’s fully commented and uses only inline functions plus a built-in constant array, it’s easy for others to experiment—swap in your own permutation table, replace the gradient logic with 4D OpenSimplex2S, or layer multiple noise octaves for ridged/multifractal terrain.

Final Notes

Performance: For most mid-resolution builds (512×512 to 2048×2048), this kernel runs in under a second on a mid-range GPU. If you need larger tile sizes, consider tiling at lower frequency or implementing an FBM wrapper to reduce kernel invocations.
Further Improvements: You can extend this basic Perlin device by adding:
- Multiple Octaves: Wrap consecutive calls to perlin3() at different scale values, accumulate results for fractal noise.
- Domain Warping: Feed one Perlin pass into the input of another to create “turbulent” or “ridged” terrain patterns.
- Speed Optimizations: Precompute and upload a 256×256 gradient texture instead of the 512-entry lookup, letting the GPU’s texture cache accelerate gradient reads.
Acknowledgments:
- All credit to Ken Perlin for the improved noise algorithm.
- Thanks to the World Machine community for scattered code snippets that helped clarify argument ordering and buffer conventions.

Feel free to copy, adapt, and build upon this device—just keep the MIT license header intact so future users know they have full permission to use or modify it. Good luck with your own procedural worlds!

Stephen · June 4, 2025, 9:36pm

Just a note that an opencl perlin noise implementation is one of the included sample worlds with World Machine (Device Examples → Code → OpenCL Perlin Noise)

NYC_Trucks · June 5, 2025, 4:09am

Yes, thank you. I am now aware. I developed the kernel prior to discovering the example. I am glad I at least got to experience the blissfulness of that ignorance. I really just wanted to see if anyone else had created or was trying to create any useful solutions with the Code Device.