Add a typedef to 32 bit unsigned int as "Color"

src/CUDA/3d_points_lines.cu

@@ -10,7 +10,7 @@
 #include "common_graphics.cuh" // Contains fill_circle
 
 __global__ void render_points_kernel(
-    unsigned int* pixels, int width, int height,
+    Color* pixels, int width, int height,
     float geom_mean_size, float points_opacity, float points_radius_multiplier,
     Cuda::Point* points, int num_points,
     Cuda::quat camera_direction, Cuda::vec3 camera_pos, float fov)
@@ -31,7 +31,7 @@ __global__ void render_points_kernel(
 }
 
 __global__ void render_lines_kernel(
-    unsigned int* pixels, int width, int height,
+    Color* pixels, int width, int height,
     float geom_mean_size, int thickness, float lines_opacity,
     Cuda::Line* lines, int num_lines,
     Cuda::quat camera_direction, Cuda::vec3 camera_pos, float fov)
@@ -59,14 +59,14 @@ __global__ void render_lines_kernel(
 }
 
 extern "C" void render_points_on_gpu(
-    unsigned int* h_pixels, int width, int height,
+    Color* h_pixels, int width, int height,
     float geom_mean_size, float points_opacity, float points_radius_multiplier,
     Cuda::Point* h_points, int num_points,
     Cuda::quat camera_direction, Cuda::vec3 camera_pos, float fov)
 {
-    unsigned int* d_pixels = nullptr;
+    Color* d_pixels = nullptr;
     Cuda::Point*        d_points = nullptr;
-    size_t pix_sz = width * height * sizeof(unsigned int);
+    size_t pix_sz = width * height * sizeof(Color);
     size_t pt_sz  = num_points * sizeof(Cuda::Point);
 
     cudaMalloc((void**)&d_pixels, pix_sz);
@@ -90,14 +90,14 @@ extern "C" void render_points_on_gpu(
 }
 
 extern "C" void render_lines_on_gpu(
-    unsigned int* h_pixels, int width, int height,
+    Color* h_pixels, int width, int height,
     float geom_mean_size, int thickness, float lines_opacity,
     Cuda::Line* h_lines, int num_lines,
     Cuda::quat camera_direction, Cuda::vec3 camera_pos, float fov)
 {
-    unsigned int* d_pixels = nullptr;
+    Color* d_pixels = nullptr;
     Cuda::Line*         d_lines  = nullptr;
-    size_t pix_sz = width * height * sizeof(unsigned int);
+    size_t pix_sz = width * height * sizeof(Color);
     size_t ln_sz  = num_lines * sizeof(Cuda::Line);
 
     cudaMalloc((void**)&d_pixels, pix_sz);

src/CUDA/background_color.cu

@@ -6,9 +6,9 @@
 #include "color.cuh"
 
 __global__
-void alpha_overlay_kernel(unsigned int* pixels,
+void alpha_overlay_kernel(Color* pixels,
                           int width, int height,
-                          int bg)
+                          Color bg)
 {
     // 2D coordinates
     int x = blockIdx.x * blockDim.x + threadIdx.x;
@@ -18,15 +18,15 @@ void alpha_overlay_kernel(unsigned int* pixels,
 
     int idx = y * width + x;
 
-    unsigned int pix = pixels[idx];
+    Color pix = pixels[idx];
 
     pixels[idx] = d_color_combine(bg, pix);
 }
 
 extern "C"
-void alpha_overlay_cuda(unsigned int* src_host,
+void alpha_overlay_cuda(Color* src_host,
                         int width, int height,
-                        unsigned int bg_color)
+                        Color bg_color)
 {
     if (!src_host || width <= 0 || height <= 0) return;
 
@@ -54,4 +54,3 @@ void alpha_overlay_cuda(unsigned int* src_host,
     // Free device memory
     cudaFree(d_src);
 }
-

src/CUDA/beaver_grid.cu

@@ -37,7 +37,7 @@ __device__ void decode_turing_machine_index(int x, int y, int grid_w, int grid_h
     }
 }
 
-__global__ void beaver_grid_kernel(int num_states, int num_symbols, unsigned int* pixels, int w, int h, int grid_w, int grid_h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, int max_steps) {
+__global__ void beaver_grid_kernel(int num_states, int num_symbols, Color* pixels, int w, int h, int grid_w, int grid_h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, int max_steps) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     int idy = blockIdx.y * blockDim.y + threadIdx.y;
 
@@ -91,8 +91,8 @@ __global__ void beaver_grid_kernel(int num_states, int num_symbols, unsigned int
     pixels[pixel_index] = halted ? d_rainbow(atan_steps) : 0xff000000;
 }
 
-extern "C" void beaver_grid_cuda(int num_states, int num_symbols, unsigned int* pixels, int w, int h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, int max_steps) {
-    unsigned int* d_pixels;
+extern "C" void beaver_grid_cuda(int num_states, int num_symbols, Color* pixels, int w, int h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, int max_steps) {
+    Color* d_pixels;
     int w_base = (num_states + 1);
     int h_base = 2 * num_symbols;
 // given by (2m(n+1))^(mn)
@@ -103,7 +103,7 @@ extern "C" void beaver_grid_cuda(int num_states, int num_symbols, unsigned int*
         grid_w *= w_base;
         grid_h *= h_base;
     }
-    size_t size = w * h * sizeof(unsigned int);
+    size_t size = w * h * sizeof(Color);
     cudaMalloc(&d_pixels, size);
     dim3 blockSize(16, 16);
     dim3 gridSize((w + blockSize.x - 1) / blockSize.x, (h + blockSize.y - 1) / blockSize.y);

src/CUDA/beaver_grid_TNF.cu

@@ -37,7 +37,7 @@ __device__ void get_child(TuringMachine& tm, int action_index, const Cuda::vec2&
     tm.num_states += (int)(tm.next_state[action_index] == tm.num_states-1);
 }
 
-__global__ void beaver_TNF_kernel(unsigned int* pixels, int w, int h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, bool corners_in_01, float border_thickness, TuringMachine TNF_root, int max_steps) {
+__global__ void beaver_TNF_kernel(Color* pixels, int w, int h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, bool corners_in_01, float border_thickness, TuringMachine TNF_root, int max_steps) {
     // convert the corners so that instead of storing the screen corners in the CoordinateScene's coordinate system, they store the TNF grid corners in the screen's coordinate system
     Cuda::vec2 grid_lx_ty = -lx_ty / (rx_by - lx_ty);
     rx_by = (Cuda::vec2(1,1) - lx_ty) / (rx_by - lx_ty);
@@ -134,9 +134,9 @@ __global__ void beaver_TNF_kernel(unsigned int* pixels, int w, int h, Cuda::vec2
     pixels[pixel_index] = d_rainbow(atanf(depth*depth / 200.) / 1.57079632679f);
 }
 
-extern "C" void beaver_grid_TNF_cuda(unsigned int* pixels, int w, int h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, int max_steps) {
-    unsigned int* d_pixels;
-    size_t size = w * h * sizeof(unsigned int);
+extern "C" void beaver_grid_TNF_cuda(Color* pixels, int w, int h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, int max_steps) {
+    Color* d_pixels;
+    size_t size = w * h * sizeof(Color);
     cudaMalloc(&d_pixels, size);
     dim3 blockSize(16, 16);
     dim3 gridSize((w + blockSize.x - 1) / blockSize.x, (h + blockSize.y - 1) / blockSize.y);

src/CUDA/beaver_grid_TNF_3d.cu

@@ -266,7 +266,7 @@ __device__ void loop5(Cuda::vec3& raypos, Cuda::vec3 raydir, Cuda::vec4& raycol,
 
 
 
-__global__ void beaver_raytrace_kernel(unsigned int* pixels, int w, int h, Cuda::vec3 raypos, Cuda::quat camera, float fov, int max_steps, float shell_border, float core_border, Cuda::vec3 scale, Cuda::vec3 target) {
+__global__ void beaver_raytrace_kernel(Color* pixels, int w, int h, Cuda::vec3 raypos, Cuda::quat camera, float fov, int max_steps, float shell_border, float core_border, Cuda::vec3 scale, Cuda::vec3 target) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     int idy = blockIdx.y * blockDim.y + threadIdx.y;
     if (idx >= w || idy >= h) {
@@ -318,21 +318,21 @@ __global__ void beaver_raytrace_kernel(unsigned int* pixels, int w, int h, Cuda:
             t++;
         }
     }
-    pixels[pixel_index] = ((unsigned int)(floor(raycol.x * 255)) << 24) | ((unsigned int)(floor(raycol.y * 255)) << 16) | ((unsigned int)(floor(raycol.z * 255)) << 8) | (unsigned int)(floor(raycol.w * 255));
+    pixels[pixel_index] = (static_cast<Color>(floor(raycol.x * 255)) << 24) | (static_cast<Color>(floor(raycol.y * 255)) << 16) | (static_cast<Color>(floor(raycol.z * 255)) << 8) | static_cast<Color>(floor(raycol.w * 255));
 }
 
 extern "C" void beaver_grid_TNF_3D_cuda(
-    unsigned int* pixels,
+    Color* pixels,
     int w, int h,
     Cuda::vec3 pos, Cuda::quat camera, float fov,
     Cuda::vec3 target, Cuda::vec3 up, float use_quat_camera,
     float shell_border, float core_border,
     Cuda::vec3 scale,
     int max_steps
 ){
-    unsigned int* d_pixels;
+    Color* d_pixels;
 
-    cudaMalloc(&d_pixels, w * h * sizeof(unsigned int));
+    cudaMalloc(&d_pixels, w * h * sizeof(Color));
     dim3 threads(16, 16);
     dim3 block((w + threads.x - 1) / threads.x, (h + threads.y - 1) / threads.y);
 
@@ -344,7 +344,7 @@ extern "C" void beaver_grid_TNF_3D_cuda(
 
     beaver_raytrace_kernel<<<block, threads>>>(d_pixels, w, h, pos, camera, fov, max_steps, shell_border, core_border, scale, target);
 
-    cudaMemcpy(pixels, d_pixels, w * h * sizeof(unsigned int), cudaMemcpyDeviceToHost);
+    cudaMemcpy(pixels, d_pixels, w * h * sizeof(Color), cudaMemcpyDeviceToHost);
 
     cudaFree(d_pixels);
-}
+}

src/CUDA/color.cuh

@@ -1,32 +1,33 @@
 #pragma once
 
 #include <thrust/complex.h>
+#include "../Host_Device_Shared/Color.h"
 #include "../Host_Device_Shared/helpers.h"
 
-__device__ __forceinline__ int d_argb(int a, int r, int g, int b){return (int(Cuda::clamp(a,0,255))<<24)|
-                                                                         (int(Cuda::clamp(r,0,255))<<16)|
-                                                                         (int(Cuda::clamp(g,0,255))<<8 )|
-                                                                         (int(Cuda::clamp(b,0,255))    );}
+__device__ __forceinline__ Color d_argb(int a, int r, int g, int b){return (static_cast<Color>(Cuda::clamp(a,0,255))<<24)|
+                                                                            (static_cast<Color>(Cuda::clamp(r,0,255))<<16)|
+                                                                            (static_cast<Color>(Cuda::clamp(g,0,255))<<8 )|
+                                                                             static_cast<Color>(Cuda::clamp(b,0,255))    ;}
 
-__device__ __forceinline__ int d_geta(int color) { 
+__device__ __forceinline__ int d_geta(Color color) {
     return (color >> 24) & 0xFF; 
 }
-__device__ __forceinline__ int d_getr(int color) { 
+__device__ __forceinline__ int d_getr(Color color) {
     return (color >> 16) & 0xFF; 
 }
-__device__ __forceinline__ int d_getg(int color) { 
+__device__ __forceinline__ int d_getg(Color color) {
     return (color >> 8) & 0xFF; 
 }
-__device__ __forceinline__ int d_getb(int color) { 
+__device__ __forceinline__ int d_getb(Color color) {
     return color & 0xFF; 
 }
 
-__device__ __forceinline__ int d_rainbow(float x){return d_argb(255,
-                                            sin((x+1/3.)*M_PI*2)*128.+128,
-                                            sin((x+2/3.)*M_PI*2)*128.+128,
-                                            sin((x     )*M_PI*2)*128.+128);}
+__device__ __forceinline__ Color d_rainbow(float x){return d_argb(255,
+                                              sin((x+1/3.)*M_PI*2)*128.+128,
+                                              sin((x+2/3.)*M_PI*2)*128.+128,
+                                              sin((x     )*M_PI*2)*128.+128);}
 
-__device__ __forceinline__ int d_colorlerp(int c1, int c2, float t) {
+__device__ __forceinline__ Color d_colorlerp(Color c1, Color c2, float t) {
     int a1 = (c1 >> 24) & 0xFF; int r1 = (c1 >> 16) & 0xFF; int g1 = (c1 >> 8) & 0xFF; int b1 = c1 & 0xFF;
     int a2 = (c2 >> 24) & 0xFF; int r2 = (c2 >> 16) & 0xFF; int g2 = (c2 >> 8) & 0xFF; int b2 = c2 & 0xFF;
     int a = roundf((1 - t) * a1 + t * a2);
@@ -36,54 +37,54 @@ __device__ __forceinline__ int d_colorlerp(int c1, int c2, float t) {
     return (a << 24) | (r << 16) | (g << 8) | b;
 }
 
-__device__ __forceinline__ int d_color_combine(int base_color, int over_color, float overlay_opacity_multiplier = 1) {
+__device__ __forceinline__ Color d_color_combine(Color base_color, Color over_color, float overlay_opacity_multiplier = 1) {
     float base_opacity = d_geta(base_color) / 255.0f;
     float over_opacity = d_geta(over_color) / 255.0f * overlay_opacity_multiplier;
     float final_opacity = 1 - (1 - base_opacity) * (1 - over_opacity);
     if (final_opacity == 0) return 0x00000000;
     int final_alpha = roundf(final_opacity * 255.0f);
     float chroma_weight = over_opacity / final_opacity;
-    int final_rgb = d_colorlerp(base_color, over_color, chroma_weight) & 0x00ffffff;
-    return (final_alpha << 24) | (final_rgb);
+    Color final_rgb = d_colorlerp(base_color, over_color, chroma_weight) & 0x00ffffff;
+    return (static_cast<Color>(final_alpha) << 24) | final_rgb;
 }
 
-__device__ __forceinline__ void d_set_pixel(int x, int y, int col, unsigned int* pixels, int width, int height) {
+__device__ __forceinline__ void d_set_pixel(int x, int y, Color col, Color* pixels, int width, int height) {
     if (x < 0 || x >= width || y < 0 || y >= height) return;
     pixels[y * width + x] = col;
 }
 
-__device__ __forceinline__ void d_overlay_pixel(int x, int y, int col, float opacity, unsigned int* pixels, int width, int height) {
+__device__ __forceinline__ void d_overlay_pixel(int x, int y, Color col, float opacity, Color* pixels, int width, int height) {
     if (x < 0 || x >= width || y < 0 || y >= height) return;
     opacity = Cuda::clamp(opacity, 0.0f, 1.0f);
     int idx = y * width + x;
-    int base = pixels[idx];
-    int blended = d_color_combine(base, col, opacity);
+    Color base = pixels[idx];
+    Color blended = d_color_combine(base, col, opacity);
     pixels[idx] = blended;
 }
 
-__device__ __forceinline__ void d_naive_add_pixel(int x, int y, int col, float opacity, unsigned int* pixels, int width, int height) {
+__device__ __forceinline__ void d_naive_add_pixel(int x, int y, Color col, float opacity, Color* pixels, int width, int height) {
     if (x < 0 || x >= width || y < 0 || y >= height) return;
     opacity = Cuda::clamp(opacity, 0.0f, 1.0f);
     int idx = y * width + x;
-    int base = pixels[idx];
+    Color base = pixels[idx];
     // Naively add each channel multiplied by opacity, clamping to 255
     int a = d_geta(base) + int(d_geta(col) * opacity);
     int r = d_getr(base) + int(d_getr(col) * opacity);
     int g = d_getg(base) + int(d_getg(col) * opacity);
     int b = d_getb(base) + int(d_getb(col) * opacity);
-    int blended = d_argb(a, r, g, b);
+    Color blended = d_argb(a, r, g, b);
     pixels[idx] = blended;
 }
 
-__device__ __forceinline__ void d_atomic_overlay_pixel(int x, int y, int col, float opacity, unsigned int* pixels, int width, int height) {
+__device__ __forceinline__ void d_atomic_overlay_pixel(int x, int y, Color col, float opacity, Color* pixels, int width, int height) {
     if (x < 0 || x >= width || y < 0 || y >= height) return;
     opacity = Cuda::clamp(opacity, 0.0f, 1.0f);
     int idx = y * width + x;
 
-    unsigned int old_pixel = pixels[idx];
-    int base = old_pixel;
-    int blended = d_color_combine(base, col, opacity);
-    int new_pixel = blended;
+    Color old_pixel = pixels[idx];
+    Color base = old_pixel;
+    Color blended = d_color_combine(base, col, opacity);
+    Color new_pixel = blended;
     atomicCAS(&pixels[idx], old_pixel, new_pixel);
 }
 
@@ -94,7 +95,7 @@ __device__ __forceinline__ float d_linear_srgb_to_srgb(float x) {
     return 12.92 * x;
 }
 
-__device__ __forceinline__ int d_OKLABtoRGB(int alpha, float L, float a, float b)
+__device__ __forceinline__ Color d_OKLABtoRGB(int alpha, float L, float a, float b)
 {
     float l_ = L + 0.3963377774f * a + 0.2158037573f * b;
     float m_ = L - 0.1055613458f * a - 0.0638541728f * b;
@@ -112,15 +113,15 @@ __device__ __forceinline__ int d_OKLABtoRGB(int alpha, float L, float a, float b
     );
 }
 
-__device__ __forceinline__ int d_complex_to_srgb(const thrust::complex<float>& c, float ab_dilation, float dot_radius) {
+__device__ __forceinline__ Color d_complex_to_srgb(const thrust::complex<float>& c, float ab_dilation, float dot_radius) {
     float mag = abs(c);
     if(mag < 1e-7) return d_argb(255, 255, 255, 255); // white to dodge division by zero 
     thrust::complex<float> norm = (c * ab_dilation / mag + thrust::complex<float>(1,1)) * .5;
     float am = 2*atan(mag/dot_radius)/M_PI;
     return d_OKLABtoRGB(255, 1-.8*am, Cuda::lerp(-.233888, .276216, norm.real()), Cuda::lerp(-.311528, .198570, norm.imag()));
 }
 
-__device__ __forceinline__ int d_HSVtoRGB(float h, float s, float v, int alpha = 255) {
+__device__ __forceinline__ Color d_HSVtoRGB(float h, float s, float v, int alpha = 255) {
     float r_f, g_f, b_f;
 
     if (s == 0.0) {

src/CUDA/common_graphics.cuh

@@ -5,7 +5,7 @@
 namespace Cuda {
 
 // Fill a circle on a pixel buffer
-__device__ __forceinline__ void d_fill_circle(float cx, float cy, float r, int col, unsigned int* pixels, int width, int height, float opa=1.0f) {
+__device__ __forceinline__ void d_fill_circle(float cx, float cy, float r, Color col, Color* pixels, int width, int height, float opa=1.0f) {
     // breakout if outside of screen
     if (cx + r < 0 || cx - r >= width || cy + r < 0 || cy - r >= height)
         return;
@@ -19,7 +19,7 @@ __device__ __forceinline__ void d_fill_circle(float cx, float cy, float r, int c
     }
 }
 
-__device__ __forceinline__ void bresenham(int x1, int y1, int x2, int y2, int col, float opacity, int thickness, unsigned int* pixels, int width, int height) {
+__device__ __forceinline__ void bresenham(int x1, int y1, int x2, int y2, Color col, float opacity, int thickness, Color* pixels, int width, int height) {
     int dx = abs(x2 - x1), dy = abs(y2 - y1);
     if (dx > 10000 || dy > 10000) return;
     int sx = (x1 < x2) ? 1 : -1;

src/CUDA/conway.cu

@@ -122,7 +122,7 @@ extern "C" void iterate_conway(Bitboard* d_board, Bitboard* d_board_2, int w_bit
     cudaDeviceSynchronize();
 }
 
-__global__ void conway_draw_kernel(Bitboard* board, Bitboard* board_2, int w_bitboards, int h_bitboards, unsigned int* pixels, int pixels_w, int pixels_h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, float w_t)
+__global__ void conway_draw_kernel(Bitboard* board, Bitboard* board_2, int w_bitboards, int h_bitboards, Color* pixels, int pixels_w, int pixels_h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, float w_t)
 {
     int px = blockDim.x * blockIdx.x + threadIdx.x;
     int py = blockDim.y * blockIdx.y + threadIdx.y;
@@ -221,11 +221,11 @@ __global__ void conway_draw_kernel(Bitboard* board, Bitboard* board_2, int w_bit
     }
 }
 
-extern "C" void draw_conway(Bitboard* d_board, Bitboard* d_board_2, int w_bitboards, int h_bitboards, unsigned int* h_pixels, int pixels_w, int pixels_h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, float transition)
+extern "C" void draw_conway(Bitboard* d_board, Bitboard* d_board_2, int w_bitboards, int h_bitboards, Color* h_pixels, int pixels_w, int pixels_h, Cuda::vec2 lx_ty, Cuda::vec2 rx_by, float transition)
 {
-    size_t pix_sz = pixels_w * pixels_h * sizeof(unsigned int);
+    size_t pix_sz = pixels_w * pixels_h * sizeof(Color);
 
-    unsigned int* d_pixels;
+    Color* d_pixels;
 
     cudaMalloc((void**)&d_pixels, pix_sz);