//Creating map in host code pre-kernel:
    cl_int status;
    cl_image_format map_form = {CL_LUMINANCE, CL_UNORM_INT8};
    clCreateImage2D(context, CL_MEM_WRITE_ONLY, map_form,
        width, height, 0, NULL, &status);


//Doing one worker per 4-pixel block so we can vectorize the transformations,
//meaning dimensions of map must be divisible by 4.  Kernel is dispatched in
//two dimensions.
__kernel __attribute__((vec_type_hint(float4)))
void genMap( __global float min,
                      __global float max,
                      __global float *heights,
                      __global image2d_t map) {
    //Worker ID on X and Y
    int x = get_global_id(0);
    int y = get_global_id(1);
    //Job size in X and Y
    int xs = get_global_size(0);
    //int ys = get_global_size(1);  //Not needed

    //TBD: is there any performance benefit to this explicit vectorization?
    //..maybe the compiler will know better what benefits it.
    float4 norm = vload4(&heights[y*xs + x];
    //Convert so 0 <= norm <= 1, treat it like luminance
    norm -= (float4)min;
    norm /= (float4)(max - min);

    write_imagef(map, (int2)(x+0,y), norm.x);
    write_imagef(map, (int2)(x+1,y), norm.y);
    write_imagef(map, (int2)(x+2,y), norm.z);
    write_imagef(map, (int2)(x+3,y), norm.w);
}