diff --git a/benchmark/README.md b/benchmark/README.md index 68fd5ae2e9..d929d8c908 100644 --- a/benchmark/README.md +++ b/benchmark/README.md @@ -1,6 +1,11 @@ -# Sigmoid Kernel Benchmarks +# Darktable Mojo Benchmarks -This directory contains benchmarks and parity checks for the darktable `sigmoid` module kernels, comparing the traditional **OpenCL (C)** implementation against the **Mojo GPU** implementation. +This directory contains performance benchmarks and validation tools for darktable module kernels ported to Mojo. + +## Subprojects + +- **[Sigmoid](./sigmoid)**: Benchmarks and parity checks for the sigmoid module. +- **[Blurs](./blurs)**: Benchmarks for the blurs convolution module. ## Prerequisites @@ -10,49 +15,8 @@ This directory contains benchmarks and parity checks for the darktable `sigmoid` ## Environment Setup -Initialize the environment using Pixi: +Initialize the environment using Pixi in the root `benchmark/` directory: ```bash pixi install -# or enter the shell -pixi shell -``` - -## Running Benchmarks - -### 1. OpenCL Baseline (C) -The C benchmark measures the performance of the original OpenCL kernels. - -```bash -make run ``` - -### 2. Mojo GPU Benchmark -The Mojo benchmark measures the performance of the ported kernels using Mojo's GPU abstraction. - -```bash -pixi run mojo sigmoid_benchmark_gpu.mojo -``` - -## Parity Validation - -To ensure the Mojo implementation produces numerically identical results to the OpenCL baseline, run the parity check script: - -```bash -pixi run python validate_parity.py -``` - -This script: -1. Compiles and runs the OpenCL parity check (`parity_check.c`). -2. Runs the Mojo GPU benchmark. -3. Compares sampled pixels across both implementations (RGB Ratio and Per-Channel modes). -4. Reports "PASS" if the results match within a tolerance of $10^{-6}$. - -## Files - -- `benchmark_sigmoid.c`: Main C-based OpenCL benchmark. -- `sigmoid_benchmark_gpu.mojo`: Mojo implementation and benchmark. -- `parity_check.c`: Minimal OpenCL runner for numerical validation. -- `validate_parity.py`: Automated comparison tool. -- `Makefile`: Build instructions for C/OpenCL binaries. -- `pixi.toml`: Project dependencies (Mojo, Python, etc.). diff --git a/benchmark/blurs/Makefile b/benchmark/blurs/Makefile new file mode 100644 index 0000000000..e16dc6b570 --- /dev/null +++ b/benchmark/blurs/Makefile @@ -0,0 +1,18 @@ +CC=gcc +CFLAGS=-O3 -Wall +LDFLAGS=-lOpenCL -lm + +TARGET=benchmark_blurs +SRC=benchmark_blurs.c + +all: $(TARGET) + +.PHONY: force +$(TARGET): $(SRC) force + $(CC) $(CFLAGS) $(SRC) -o $(TARGET) $(LDFLAGS) + +clean: + rm -f $(TARGET) + +run: all + ./$(TARGET) diff --git a/benchmark/blurs/README.md b/benchmark/blurs/README.md new file mode 100644 index 0000000000..efae830150 --- /dev/null +++ b/benchmark/blurs/README.md @@ -0,0 +1,30 @@ +# Blurs Kernel Benchmarks + +This directory contains benchmarks for the darktable `blurs` module kernels, comparing the **OpenCL (C)** implementation against the **Mojo GPU (Tiled)** implementation. + +## Running Benchmarks + +### 1. OpenCL Baseline (C) +```bash +make run +``` + +### 2. Mojo GPU Benchmark +Run from the `benchmark/blurs` directory. Use `-I` to point to the `mojo/` source root. + +```bash +mojo -I /path/to/darktable/mojo blurs_benchmark_gpu.mojo +``` + +Or if using pixi: +```bash +pixi run mojo -I /path/to/darktable/mojo blurs_benchmark_gpu.mojo +``` + +From the darktable project root: +```bash +cd benchmark/blurs && mojo -I ../../mojo blurs_benchmark_gpu.mojo +``` + +## Performance Notes +The Mojo implementation uses a tiled approach with SIMD vectorization to ensure coalesced memory access on the GPU. This is designed to be significantly faster than a naive implementation. diff --git a/benchmark/blurs/benchmark_blurs b/benchmark/blurs/benchmark_blurs new file mode 100755 index 0000000000..498cf22a4f Binary files /dev/null and b/benchmark/blurs/benchmark_blurs differ diff --git a/benchmark/blurs/benchmark_blurs.c b/benchmark/blurs/benchmark_blurs.c new file mode 100644 index 0000000000..38a80a7406 --- /dev/null +++ b/benchmark/blurs/benchmark_blurs.c @@ -0,0 +1,180 @@ +#define CL_TARGET_OPENCL_VERSION 120 +#include +#include +#include +#include +#include +#include + +#define CHECK_CL(cmd) \ + { \ + cl_int _cl_err = cmd; \ + if (_cl_err != CL_SUCCESS) { \ + fprintf(stderr, "OpenCL error %d at %s:%d\n", _cl_err, __FILE__, __LINE__); \ + exit(1); \ + } \ + } + +char* read_file(const char* filename) { + FILE* f = fopen(filename, "rb"); + if (!f) return NULL; + fseek(f, 0, SEEK_END); + long size = ftell(f); + fseek(f, 0, SEEK_SET); + char* buf = (char*)malloc(size + 1); + if (!buf) { + fclose(f); + return NULL; + } + if (fread(buf, 1, size, f) != (size_t)size) { + free(buf); + fclose(f); + return NULL; + } + buf[size] = '\0'; + fclose(f); + return buf; +} + +static double run_benchmark(cl_context context, cl_command_queue queue, cl_program program, + cl_mem d_in, int width, int height, int radius) { + int k_width = 2 * radius + 1; + int iterations = 100; + if (radius >= 12) iterations = 30; + else if (radius >= 5) iterations = 80; + + // Create kernel + output image (reusable) + cl_image_format format = { CL_RGBA, CL_FLOAT }; + cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, 0, 0, 0, 0, {NULL} }; + cl_int err; + cl_mem d_out = clCreateImage(context, CL_MEM_WRITE_ONLY, &format, &desc, NULL, &err); + CHECK_CL(err); + + // Create kernel image for this radius + size_t k_data_size = (size_t)k_width * k_width * sizeof(float); + float* h_kern = (float*)malloc(k_data_size); + for (int i = 0; i < k_width * k_width; i++) + h_kern[i] = 1.0f / (k_width * k_width); + + cl_image_format kern_format = { CL_R, CL_FLOAT }; + cl_image_desc kern_desc = { CL_MEM_OBJECT_IMAGE2D, k_width, k_width, 0, 0, 0, 0, 0, 0, {NULL} }; + cl_mem d_kern = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &kern_format, &kern_desc, h_kern, &err); + CHECK_CL(err); + + cl_kernel kernel = clCreateKernel(program, "convolve", &err); + CHECK_CL(err); + + CHECK_CL(clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_in)); + CHECK_CL(clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_kern)); + CHECK_CL(clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_out)); + CHECK_CL(clSetKernelArg(kernel, 3, sizeof(int), &width)); + CHECK_CL(clSetKernelArg(kernel, 4, sizeof(int), &height)); + CHECK_CL(clSetKernelArg(kernel, 5, sizeof(int), &radius)); + + printf("Benchmarking OpenCL convolve (Radius: %d, %dx%d, %d iters)...\n", radius, width, height, iterations); + fflush(stdout); + + cl_event event; + double total_time = 0; + for (int i = 0; i < iterations; i++) { + size_t global_work_size[2] = { (size_t)width, (size_t)height }; + CHECK_CL(clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global_work_size, NULL, 0, NULL, &event)); + clWaitForEvents(1, &event); + cl_ulong start, end; + clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(start), &start, NULL); + clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(end), &end, NULL); + total_time += (double)(end - start) / 1000000.0; + clReleaseEvent(event); + } + double avg = total_time / iterations; + printf(" Average Time: %.4f ms\n", avg); + + clReleaseMemObject(d_kern); + clReleaseMemObject(d_out); + clReleaseKernel(kernel); + free(h_kern); + return avg; +} + +int main() { + printf("Starting C benchmark...\n"); + fflush(stdout); + + cl_int err; + cl_platform_id platform; + cl_device_id device; + cl_context context; + cl_command_queue queue; + int width = 6016; + int height = 4016; + size_t img_size = (size_t)width * height * 4 * sizeof(float); + + printf("Allocating memory for %dx%d image...\n", width, height); + fflush(stdout); + float* h_data = (float*)malloc(img_size); + if (!h_data) { fprintf(stderr, "Failed to allocate h_data\n"); return 1; } + + printf("Initializing host data (%zu bytes)...\n", img_size); + fflush(stdout); + for (size_t i = 0; i < (size_t)width * height * 4; i++) h_data[i] = 0.5f; + printf("Host data initialized.\n"); + fflush(stdout); + + // OpenCL Initialization + printf("Initializing OpenCL...\n"); + fflush(stdout); + cl_uint num_platforms; + err = clGetPlatformIDs(0, NULL, &num_platforms); + if (err != CL_SUCCESS || num_platforms == 0) { + fprintf(stderr, "No OpenCL platforms found\n"); + return 1; + } + CHECK_CL(clGetPlatformIDs(1, &platform, NULL)); + CHECK_CL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL)); + + char device_name[128]; + clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_name), device_name, NULL); + printf("Using device: %s\n", device_name); + + context = clCreateContext(NULL, 1, &device, NULL, NULL, &err); + CHECK_CL(err); + queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &err); + CHECK_CL(err); + + // Load kernel source + const char* kernel_path = "../../data/kernels/blurs.cl"; + char* source = read_file(kernel_path); + if (!source) { fprintf(stderr, "Failed to load kernel\n"); return 1; } + + cl_program program = clCreateProgramWithSource(context, 1, (const char**)&source, NULL, &err); + CHECK_CL(err); + const char* options = "-I ../../data/kernels/"; + err = clBuildProgram(program, 1, &device, options, NULL, NULL); + if (err != CL_SUCCESS) { + char log[16384]; + clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(log), log, NULL); + fprintf(stderr, "Build error:\n%s\n", log); + return 1; + } + + // Create input image (reused across radii) + cl_image_format format = { CL_RGBA, CL_FLOAT }; + cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, 0, 0, 0, 0, {NULL} }; + cl_mem d_in = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + &format, &desc, h_data, &err); + CHECK_CL(err); + + int radii[] = {3, 8, 15}; + int num_radii = sizeof(radii) / sizeof(radii[0]); + for (int r = 0; r < num_radii; r++) { + run_benchmark(context, queue, program, d_in, width, height, radii[r]); + } + + free(source); free(h_data); + clReleaseMemObject(d_in); + clReleaseProgram(program); + clReleaseCommandQueue(queue); + clReleaseContext(context); + return 0; +} diff --git a/benchmark/blurs/benchmark_results.md b/benchmark/blurs/benchmark_results.md new file mode 100644 index 0000000000..2571b2f9ef --- /dev/null +++ b/benchmark/blurs/benchmark_results.md @@ -0,0 +1,59 @@ +# GPU Convolution Benchmark Results + +An execution-time comparison on a massive 24.1 Megapixel ($6016 \times 4016$, 4-channel float32) image, comparing: +1. **OpenCL C-Benchmark** (`benchmark_blurs.c` using C/OpenCL host) +2. **Mojo GPU Naive Baseline** (replicating the global memory `UnsafePointer` implementation from `iop/blurs` in Mojo) +3. **Mojo GPU Tiled Convolution** (using a single compiled kernel with dynamic runtime `radius` parameter, zero `UnsafePointer` usage) + +--- + +## Benchmark Configuration + +- **Image Resolution**: $6016 \times 4016$ (4 channels, RGBA) +- **Data Type**: `Float32` +- **Kernel Size**: $31 \times 31$ Box Blur (`RADIUS = 15`) +- **Kernel Initialization**: Normalized Box Blur (each element = $1 / 961$) +- **Target GPU**: AMD HIP/OpenCL GPU (`gfx1201`) +- **LDS Buffer Allocation Size**: Sized to support any dynamic runtime `radius` up to $25$ + +--- + +## Performance Summary + +| Implementation | Average Latency (ms) | Rel. Performance | +| :--- | :--- | :--- | +| **Mojo GPU Tiled (`TileTensor`)** | **38.5 ms** | **1.65x** (Fastest) 🏆 | +| **Mojo GPU Naive Baseline** | **63 ms** | **1.00x** | +| **OpenCL convolve (C-Benchmark)** | **69 ms** | **0.91x** | + +--- + +## Technical Analysis & Key Discoveries + +### 1. Zero-Specialization Dynamic Runtime Radius +We transitioned the GPU kernel from compile-time specialization of the radius to a pure, user-selected runtime parameter: +```mojo +def convolve_gpu_kernel[ + InLayout: TensorLayout, + OutLayout: TensorLayout, + max_radius: Int, # Sizing upper bound for LDS buffer allocation +]( + in_t: TileTensor[DTYPE, InLayout, MutAnyOrigin], + out_t: TileTensor[DTYPE, OutLayout, MutAnyOrigin], + k_val: Float32, # Precomputed box-blur weight (1 / kern_size) + radius: Int, # Dynamic runtime radius parameter + ... +) +``` +- **LDS Allocation Sized by `MAX_RADIUS = 25`**: A compile-time constant `MAX_RADIUS` sizes the shared memory array. This single compilation supports *any* runtime user-selected radius $r \le 25$ without generating multiple specializations. +- **Division-Free 2D Strided Loader**: Instead of mapping a flat thread index using slow runtime modulo and division operations (which is highly detrimental when radius is dynamic), we developed an elegant 2D strided loader: + ```mojo + for ly in range(ty, actual_halo_h, TILE_H): + for lx in range(tx, actual_halo_w, TILE_W): + ... + sh_tile[c, ly, lx] = val + ``` + This eliminates all division/modulo instructions from the loading stage and preserves perfectly coalesced, bank-conflict-free shared memory writes. + +### 2. Double-Occupancy via Alpha-Elimination +By omitting the Alpha channel from the shared tile layout (`row_major[RGB, HALO_H, HALO_W]`), we reduced the LDS footprint from **45.6 KB** to **34.2 KB**. This allows **2 active blocks per CU** on gfx1201's 64 KB LDS limit, doubling latency hiding and keeping execution throughput at an optimal level. diff --git a/benchmark/blurs/check_cl.c b/benchmark/blurs/check_cl.c new file mode 100644 index 0000000000..e31e82bda3 --- /dev/null +++ b/benchmark/blurs/check_cl.c @@ -0,0 +1,14 @@ +#include +#include + +int main() { + printf("Checking OpenCL platforms...\n"); + cl_uint num_platforms; + cl_int err = clGetPlatformIDs(0, NULL, &num_platforms); + if (err != CL_SUCCESS) { + printf("clGetPlatformIDs failed with %d\n", err); + return 1; + } + printf("Found %u platforms.\n", num_platforms); + return 0; +} diff --git a/benchmark/Makefile b/benchmark/sigmoid/Makefile similarity index 100% rename from benchmark/Makefile rename to benchmark/sigmoid/Makefile diff --git a/benchmark/sigmoid/README.md b/benchmark/sigmoid/README.md new file mode 100644 index 0000000000..68fd5ae2e9 --- /dev/null +++ b/benchmark/sigmoid/README.md @@ -0,0 +1,58 @@ +# Sigmoid Kernel Benchmarks + +This directory contains benchmarks and parity checks for the darktable `sigmoid` module kernels, comparing the traditional **OpenCL (C)** implementation against the **Mojo GPU** implementation. + +## Prerequisites + +- **Mojo**: Required for running `.mojo` benchmarks. +- **OpenCL**: Required for hardware acceleration on GPU (headers and libraries). +- **Pixi**: Used for environment and dependency management. + +## Environment Setup + +Initialize the environment using Pixi: + +```bash +pixi install +# or enter the shell +pixi shell +``` + +## Running Benchmarks + +### 1. OpenCL Baseline (C) +The C benchmark measures the performance of the original OpenCL kernels. + +```bash +make run +``` + +### 2. Mojo GPU Benchmark +The Mojo benchmark measures the performance of the ported kernels using Mojo's GPU abstraction. + +```bash +pixi run mojo sigmoid_benchmark_gpu.mojo +``` + +## Parity Validation + +To ensure the Mojo implementation produces numerically identical results to the OpenCL baseline, run the parity check script: + +```bash +pixi run python validate_parity.py +``` + +This script: +1. Compiles and runs the OpenCL parity check (`parity_check.c`). +2. Runs the Mojo GPU benchmark. +3. Compares sampled pixels across both implementations (RGB Ratio and Per-Channel modes). +4. Reports "PASS" if the results match within a tolerance of $10^{-6}$. + +## Files + +- `benchmark_sigmoid.c`: Main C-based OpenCL benchmark. +- `sigmoid_benchmark_gpu.mojo`: Mojo implementation and benchmark. +- `parity_check.c`: Minimal OpenCL runner for numerical validation. +- `validate_parity.py`: Automated comparison tool. +- `Makefile`: Build instructions for C/OpenCL binaries. +- `pixi.toml`: Project dependencies (Mojo, Python, etc.). diff --git a/benchmark/benchmark_sigmoid.c b/benchmark/sigmoid/benchmark_sigmoid.c similarity index 85% rename from benchmark/benchmark_sigmoid.c rename to benchmark/sigmoid/benchmark_sigmoid.c index 4b8f32ad6b..2e55b13d1f 100644 --- a/benchmark/benchmark_sigmoid.c +++ b/benchmark/sigmoid/benchmark_sigmoid.c @@ -148,14 +148,10 @@ int main() { // Create command queue with profiling enabled cl_queue_properties props[] = { CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0 }; queue = clCreateCommandQueueWithProperties(context, device, props, &err); - if (err != CL_SUCCESS) { - // Fallback for older OpenCL versions if needed, but we targeting 2.2 above - queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &err); - } CHECK_CL(err); // Load kernel source - const char* kernel_path = "../data/kernels/sigmoid.cl"; + const char* kernel_path = "../../data/kernels/sigmoid.cl"; printf("Reading kernel source from %s...\n", kernel_path); char* source = read_file(kernel_path); if (!source) { @@ -168,7 +164,7 @@ int main() { CHECK_CL(err); // Need to point to the directory containing common.h and colorspace.h - const char* options = "-I ../data/kernels/"; + const char* options = "-I ../../data/kernels/"; err = clBuildProgram(program, 1, &device, options, NULL, NULL); if (err != CL_SUCCESS) { char build_log[16384]; @@ -183,9 +179,6 @@ int main() { cl_kernel kernel_rgb_ratio = clCreateKernel(program, "sigmoid_loglogistic_rgb_ratio", &err); CHECK_CL(err); - CHECK_CL(clGetPlatformIDs(1, &platform, &num_platforms)); - CHECK_CL(clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL)); - cl_image_format format = { CL_RGBA, CL_FLOAT }; cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height, 0, 0, 0, 0, 0, 0, {0} }; @@ -233,10 +226,11 @@ int main() { CHECK_CL(clSetKernelArg(kernel_per_channel, 11, sizeof(cl_mem), &d_mat2)); CHECK_CL(clSetKernelArg(kernel_per_channel, 12, sizeof(cl_mem), &d_mat3)); - int warmup_iters = 100; - int iterations = 1000; + int warmup_iters = 20; + int iterations = 200; + int batch_size = 20; + int num_batches = iterations / batch_size; cl_event event; - double total_time = 0; printf(" Warmup (%d iterations)...\n", warmup_iters); for (int i = 0; i < warmup_iters; i++) { @@ -245,7 +239,8 @@ int main() { } clFinish(queue); - printf(" Benchmarking (%d iterations)...\n", iterations); + printf(" Benchmarking (%d iterations, batch_size=%d)...\n", iterations, batch_size); + double *times = (double*)malloc(iterations * sizeof(double)); for (int i = 0; i < iterations; i++) { size_t global_work_size[2] = { width, height }; CHECK_CL(clEnqueueNDRangeKernel(queue, kernel_per_channel, 2, NULL, global_work_size, NULL, 0, NULL, &event)); @@ -254,10 +249,29 @@ int main() { cl_ulong start, end; clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(start), &start, NULL); clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(end), &end, NULL); - total_time += (double)(end - start) / 1000000.0; // ns to ms + times[i] = (double)(end - start) / 1000000.0; // ns to ms clReleaseEvent(event); } - printf(" Average Time: %.4f ms\n", total_time / iterations); + + double *batch_means = (double*)malloc(num_batches * sizeof(double)); + double overall_sum = 0; + for (int b = 0; b < num_batches; b++) { + double batch_sum = 0; + for (int j = 0; j < batch_size; j++) + batch_sum += times[b * batch_size + j]; + batch_means[b] = batch_sum / batch_size; + overall_sum += batch_means[b]; + } + double mean = overall_sum / num_batches; + double sum_sq = 0; + for (int b = 0; b < num_batches; b++) { + double d = batch_means[b] - mean; + sum_sq += d * d; + } + double std = sqrt(sum_sq / num_batches); + free(batch_means); + free(times); + printf(" Mean: %.4f ms, Std: %.4f ms\n", mean, std); // Save result of the last iteration to a file float* h_out = (float*)malloc(img_size); @@ -289,8 +303,8 @@ int main() { } clFinish(queue); - total_time = 0; - printf(" Benchmarking (%d iterations)...\n", iterations); + printf(" Benchmarking (%d iterations, batch_size=%d)...\n", iterations, batch_size); + times = (double*)malloc(iterations * sizeof(double)); for (int i = 0; i < iterations; i++) { size_t global_work_size[2] = { width, height }; CHECK_CL(clEnqueueNDRangeKernel(queue, kernel_rgb_ratio, 2, NULL, global_work_size, NULL, 0, NULL, &event)); @@ -299,10 +313,29 @@ int main() { cl_ulong start, end; clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(start), &start, NULL); clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(end), &end, NULL); - total_time += (double)(end - start) / 1000000.0; // ns to ms + times[i] = (double)(end - start) / 1000000.0; // ns to ms clReleaseEvent(event); } - printf(" Average Time: %.4f ms\n", total_time / iterations); + + batch_means = (double*)malloc(num_batches * sizeof(double)); + overall_sum = 0; + for (int b = 0; b < num_batches; b++) { + double batch_sum = 0; + for (int j = 0; j < batch_size; j++) + batch_sum += times[b * batch_size + j]; + batch_means[b] = batch_sum / batch_size; + overall_sum += batch_means[b]; + } + mean = overall_sum / num_batches; + sum_sq = 0; + for (int b = 0; b < num_batches; b++) { + double d = batch_means[b] - mean; + sum_sq += d * d; + } + std = sqrt(sum_sq / num_batches); + free(batch_means); + free(times); + printf(" Mean: %.4f ms, Std: %.4f ms\n", mean, std); // Save result of the last iteration to a file h_out = (float*)malloc(img_size); diff --git a/benchmark/sigmoid/opencl_input.jpg b/benchmark/sigmoid/opencl_input.jpg new file mode 100644 index 0000000000..e5d0773daf Binary files /dev/null and b/benchmark/sigmoid/opencl_input.jpg differ diff --git a/benchmark/sigmoid/opencl_output_per_channel.jpg b/benchmark/sigmoid/opencl_output_per_channel.jpg new file mode 100644 index 0000000000..76d7b4902d Binary files /dev/null and b/benchmark/sigmoid/opencl_output_per_channel.jpg differ diff --git a/benchmark/sigmoid/opencl_output_rgb_ratio.jpg b/benchmark/sigmoid/opencl_output_rgb_ratio.jpg new file mode 100644 index 0000000000..08176fb2ca Binary files /dev/null and b/benchmark/sigmoid/opencl_output_rgb_ratio.jpg differ diff --git a/benchmark/parity_check.c b/benchmark/sigmoid/parity_check.c similarity index 100% rename from benchmark/parity_check.c rename to benchmark/sigmoid/parity_check.c diff --git a/benchmark/sigmoid/sigmoid_benchmark_gpu.mojo b/benchmark/sigmoid/sigmoid_benchmark_gpu.mojo new file mode 100644 index 0000000000..45059945ff --- /dev/null +++ b/benchmark/sigmoid/sigmoid_benchmark_gpu.mojo @@ -0,0 +1,246 @@ +from std.gpu.host import DeviceContext, DeviceBuffer, HostBuffer +from std.utils import Index +from std.math import sqrt +from std.benchmark import Bench, BenchConfig, Bencher, BenchId +from layout import Layout, LayoutTensor, UNKNOWN_VALUE +from layout.runtime_layout import RuntimeLayout, IndexList +from std.memory.unsafe_pointer import UnsafePointer +from iop.sigmoid.kernels import ( + apply_sigmoid_rgb_ratio, + apply_sigmoid_per_channel, +) +from iop.sigmoid.lib import ( + _launch_rgb_ratio_gpu, + _launch_per_channel_gpu, +) + +comptime WIDTH = 6016 +comptime HEIGHT = 4016 +comptime CHANNELS = 4 +comptime IMAGE_LAYOUT = Layout.row_major(UNKNOWN_VALUE, UNKNOWN_VALUE, CHANNELS) +comptime DTYPE = DType.float32 + + +def main() raises: + var total_floats = HEIGHT * WIDTH * CHANNELS + var ctx = DeviceContext() + print("Using GPU API:", ctx.api()) + + var input_buffer_host = ctx.enqueue_create_host_buffer[DTYPE](total_floats) + var input_buffer_device = ctx.enqueue_create_buffer[DTYPE](total_floats) + var output_buffer_device = ctx.enqueue_create_buffer[DTYPE](total_floats) + + var rt = RuntimeLayout[IMAGE_LAYOUT].row_major( + IndexList[3](HEIGHT, WIDTH, CHANNELS) + ) + + var input_image_host = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=Int(input_buffer_host.unsafe_ptr()) + ), + rt, + ) + var input_image_device = LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin]( + UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=Int(input_buffer_device.unsafe_ptr()) + ), + rt, + ) + var output_image_device = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=Int(output_buffer_device.unsafe_ptr()) + ), + rt, + ) + + for y in range(HEIGHT): + var h = Float32(y) / (HEIGHT - 1) * 6.0 + var segment = Int(h) + var f1 = h - Float32(segment) + var pr: Float32 + var pg: Float32 + var pb: Float32 + if segment == 0: + pr = 1.0; pg = f1; pb = 0.0 + elif segment == 1: + pr = 1.0 - f1; pg = 1.0; pb = 0.0 + elif segment == 2: + pr = 0.0; pg = 1.0; pb = f1 + elif segment == 3: + pr = 0.0; pg = 1.0 - f1; pb = 1.0 + elif segment == 4: + pr = f1; pg = 0.0; pb = 1.0 + elif segment == 5: + pr = 1.0; pg = 0.0; pb = 1.0 - f1 + else: + pr = 1.0; pg = 0.0; pb = 0.0 + + for x in range(WIDTH): + var r: Float32 + var g: Float32 + var b: Float32 + var mid_x = Float32(WIDTH) / 2.0 + if Float32(x) < mid_x: + var t = Float32(x) / mid_x + r = pr * t; g = pg * t; b = pb * t + else: + var t = (Float32(x) - mid_x) / (Float32(WIDTH) - 1.0 - mid_x) + r = pr * (1.0 - t) + t + g = pg * (1.0 - t) + t + b = pb * (1.0 - t) + t + input_image_host.store[width=4]( + Index(y, x, 0), SIMD[DType.float32, 4](r, g, b, 1.0) + ) + + ctx.enqueue_copy(input_buffer_device, input_buffer_host) + + var white_target = Float32(1.0) + var black_target = Float32(0.000152) + var paper_exp = Float32(0.5) + var film_fog = Float32(0.0) + var contrast_power = Float32(2.5) + var skew_power = Float32(1.0) + var hue_preservation = Float32(1.0) + + var identity = SIMD[DType.float32, 16](0) + identity[0] = 1; identity[5] = 1; identity[10] = 1; identity[15] = 1 + + _launch_rgb_ratio_gpu( + ctx, + input_image_device, + output_image_device, + white_target, + black_target, + paper_exp, + film_fog, + contrast_power, + skew_power, + WIDTH, + HEIGHT, + WIDTH * HEIGHT, + ) + ctx.synchronize() + + var output_buffer_host = ctx.enqueue_create_host_buffer[DTYPE](total_floats) + ctx.enqueue_copy(output_buffer_host, output_buffer_device) + ctx.synchronize() + var output_image_host = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=Int(output_buffer_host.unsafe_ptr()) + ), + rt, + ) + + for i in range(1, 4): + var t = Float32(i) * 0.2 + var x = Int(t * Float32(WIDTH)) + var y = Int(t * Float32(HEIGHT)) + var res = output_image_host.load[width=4](Index(y, x, 0)) + print( + "RGB Ratio -", Int(t * 100), "% pixel: [", + res[0], res[1], res[2], res[3], "]", + ) + + _launch_per_channel_gpu( + ctx, + input_image_device, + output_image_device, + white_target, + paper_exp, + film_fog, + contrast_power, + skew_power, + hue_preservation, + identity, + identity, + identity, + WIDTH, + HEIGHT, + WIDTH * HEIGHT, + ) + ctx.synchronize() + + ctx.enqueue_copy(output_buffer_host, output_buffer_device) + ctx.synchronize() + + for i in range(1, 4): + var t = Float32(i) * 0.2 + var x = Int(t * Float32(WIDTH)) + var y = Int(t * Float32(HEIGHT)) + var res = output_image_host.load[width=4](Index(y, x, 0)) + print( + "Per Channel -", Int(t * 100), "% pixel: [", + res[0], res[1], res[2], res[3], "]", + ) + + var bench = Bench(BenchConfig( + max_iters=200, num_warmup_iters=20, max_batch_size=20, + )) + + @parameter + def bench_rgb(mut b: Bencher) raises: + @parameter + def run_rgb(ctx: DeviceContext) raises: + _launch_rgb_ratio_gpu( + ctx, + input_image_device, + output_image_device, + white_target, + black_target, + paper_exp, + film_fog, + contrast_power, + skew_power, + WIDTH, + HEIGHT, + WIDTH * HEIGHT, + ) + + b.iter_custom[run_rgb](ctx) + ctx.synchronize() + + @parameter + def bench_per(mut b: Bencher) raises: + @parameter + def run_per(ctx: DeviceContext) raises: + _launch_per_channel_gpu( + ctx, + input_image_device, + output_image_device, + white_target, + paper_exp, + film_fog, + contrast_power, + skew_power, + hue_preservation, + identity, + identity, + identity, + WIDTH, + HEIGHT, + WIDTH * HEIGHT, + ) + + b.iter_custom[run_per](ctx) + ctx.synchronize() + + bench.bench_function[bench_rgb]( + BenchId("Mojo-Sigmoid-GPU-RGB-Ratio-V2"), + ) + bench.bench_function[bench_per]( + BenchId("Mojo-Sigmoid-GPU-Per-Channel-V2"), + ) + print(bench) + + for idx in range(len(bench.info_vec)): + ref info = bench.info_vec[idx] + var overall_mean = info.result.mean("ms") + var sum_sq = 0.0 + var v_n = Float64(len(info.result.runs)) + for b in range(len(info.result.runs)): + var batch_mean = info.result.runs[b].mean("ms") + var d = batch_mean - overall_mean + sum_sq += d * d + var variance = sum_sq / v_n + var std = sqrt(variance) + print(info.name, "- std:", std, "ms") diff --git a/benchmark/stb_image_write.h b/benchmark/sigmoid/stb_image_write.h similarity index 100% rename from benchmark/stb_image_write.h rename to benchmark/sigmoid/stb_image_write.h diff --git a/benchmark/validate_parity.py b/benchmark/sigmoid/validate_parity.py similarity index 100% rename from benchmark/validate_parity.py rename to benchmark/sigmoid/validate_parity.py diff --git a/benchmark/sigmoid_benchmark_gpu.mojo b/benchmark/sigmoid_benchmark_gpu.mojo deleted file mode 100644 index 0321f52346..0000000000 --- a/benchmark/sigmoid_benchmark_gpu.mojo +++ /dev/null @@ -1,601 +0,0 @@ -from gpu import block_dim, block_idx, thread_idx -from gpu.host import DeviceContext, DeviceBuffer, Dim, HostBuffer -from layout import Layout, LayoutTensor -from utils import Index, IndexList -from math import ceildiv, isnan, sqrt, pow, max, min -from benchmark import Bench, BenchConfig, Bencher, BenchId -from algorithm.functional import elementwise - -# Configuration -comptime WIDTH = 6016 -comptime HEIGHT = 4016 -comptime CHANNELS = 4 -comptime IMAGE_LAYOUT = Layout.row_major(HEIGHT, WIDTH, CHANNELS) -comptime DTYPE = DType.float32 - - -# Removed Pixel struct to use 3D LayoutTensor directly as requested - - -@always_inline -fn generalized_loglogistic_sigmoid_scalar( - value: Float32, - magnitude: Float32, - paper_exp: Float32, - film_fog: Float32, - film_power: Float32, - paper_power: Float32, -) -> Float32: - var clamped_value = max(value, Float32(0.0)) - var film_response = pow(film_fog + clamped_value, film_power) - var paper_response = magnitude * pow( - film_response / (paper_exp + film_response), paper_power - ) - if isnan(paper_response): - return magnitude - return paper_response - - -@always_inline -fn apply_sigmoid_rgb_ratio( - in_r: Float32, - in_g: Float32, - in_b: Float32, - in_a: Float32, - white_target: Float32, - black_target: Float32, - paper_exp: Float32, - film_fog: Float32, - film_power: Float32, - paper_power: Float32, -) -> SIMD[DType.float32, 4]: - # Desaturate negative values - var avg = max((in_r + in_g + in_b) / 3.0, Float32(0.0)) - var min_v = min(min(in_r, in_g), in_b) - var sat = Float32(1.0) - if min_v < 0.0: - sat = -avg / (min_v - avg) - - var p_r = avg + sat * (in_r - avg) - var p_g = avg + sat * (in_g - avg) - var p_b = avg + sat * (in_b - avg) - - var luma = (p_r + p_g + p_b) / 3.0 - var mapped_luma = generalized_loglogistic_sigmoid_scalar( - luma, white_target, paper_exp, film_fog, film_power, paper_power - ) - - if luma > 1e-9: - var scale = mapped_luma / luma - p_r *= scale - p_g *= scale - p_b *= scale - else: - p_r = mapped_luma - p_g = mapped_luma - p_b = mapped_luma - - var p_min = min(min(p_r, p_g), p_b) - var p_max = max(max(p_r, p_g), p_b) - var eps = Float32(1e-6) - var d_white = (white_target - mapped_luma) / (p_max - mapped_luma + eps) - var d_black = (black_target - mapped_luma) / (p_min - mapped_luma - eps) - var db_vs_chroma = min(d_white, d_black) - var cvm_border = (mapped_luma - p_min) / (mapped_luma + eps) - var p_chr_adj = 1.0 / (cvm_border * db_vs_chroma + eps) - var h_chr = ( - 2.0 * cvm_border / (1.0 - cvm_border * cvm_border + eps) - ) * p_chr_adj - var h_z = sqrt(h_chr * h_chr + 1.0) - var chroma_f = h_chr / (1.0 + h_z) * db_vs_chroma - - return SIMD[DType.float32, 4]( - mapped_luma + chroma_f * (p_r - mapped_luma), - mapped_luma + chroma_f * (p_g - mapped_luma), - mapped_luma + chroma_f * (p_b - mapped_luma), - in_a, - ) - - -@always_inline -fn apply_sigmoid_per_channel( - in_r: Float32, - in_g: Float32, - in_b: Float32, - in_a: Float32, - white_target: Float32, - paper_exp: Float32, - film_fog: Float32, - contrast_power: Float32, - skew_power: Float32, - hue_preservation: Float32, - pipe_to_base: SIMD[DType.float32, 16], - base_to_rendering: SIMD[DType.float32, 16], - rendering_to_pipe: SIMD[DType.float32, 16], -) -> SIMD[DType.float32, 4]: - # 1. Transform to base space - var i_r = ( - pipe_to_base[0] * in_r + pipe_to_base[1] * in_g + pipe_to_base[2] * in_b - ) - var i_g = ( - pipe_to_base[4] * in_r + pipe_to_base[5] * in_g + pipe_to_base[6] * in_b - ) - var i_b = ( - pipe_to_base[8] * in_r - + pipe_to_base[9] * in_g - + pipe_to_base[10] * in_b - ) - - # 2. Desaturate negative - var avg = max((i_r + i_g + i_b) / 3.0, Float32(0.0)) - var min_v = min(min(i_r, i_g), i_b) - var sat = Float32(1.0) - if min_v < 0.0: - sat = -avg / (min_v - avg) - i_r = avg + sat * (i_r - avg) - i_g = avg + sat * (i_g - avg) - i_b = avg + sat * (i_b - avg) - - # 3. Transform to rendering space - var r_r = ( - base_to_rendering[0] * i_r - + base_to_rendering[1] * i_g - + base_to_rendering[2] * i_b - ) - var r_g = ( - base_to_rendering[4] * i_r - + base_to_rendering[5] * i_g - + base_to_rendering[6] * i_b - ) - var r_b = ( - base_to_rendering[8] * i_r - + base_to_rendering[9] * i_g - + base_to_rendering[10] * i_b - ) - - # 4. Per-channel sigmoid curves - var pc_r = generalized_loglogistic_sigmoid_scalar( - r_r, white_target, paper_exp, film_fog, contrast_power, skew_power - ) - var pc_g = generalized_loglogistic_sigmoid_scalar( - r_g, white_target, paper_exp, film_fog, contrast_power, skew_power - ) - var pc_b = generalized_loglogistic_sigmoid_scalar( - r_b, white_target, paper_exp, film_fog, contrast_power, skew_power - ) - - # 5. Preserve hue & energy - var p_min: Float32 - var p_mid: Float32 - var p_max: Float32 - var pc_min: Float32 - var pc_mid: Float32 - var pc_max: Float32 - - if r_r >= r_g: - if r_g >= r_b: # R G B - p_max = r_r - p_mid = r_g - p_min = r_b - pc_max = pc_r - pc_mid = pc_g - pc_min = pc_b - elif r_b >= r_r: # B R G - p_max = r_b - p_mid = r_r - p_min = r_g - pc_max = pc_b - pc_mid = pc_r - pc_min = pc_g - else: # R B G - p_max = r_r - p_mid = r_b - p_min = r_g - pc_max = pc_r - pc_mid = pc_b - pc_min = pc_g - else: - if r_r >= r_b: # G R B - p_max = r_g - p_mid = r_r - p_min = r_b - pc_max = pc_g - pc_mid = pc_r - pc_min = pc_b - elif r_b >= r_g: # B G R - p_max = r_b - p_mid = r_g - p_min = r_r - pc_max = pc_b - pc_mid = pc_g - pc_min = pc_r - else: # G B R - p_max = r_g - p_mid = r_b - p_min = r_r - pc_max = pc_g - pc_mid = pc_b - pc_min = pc_r - - var chroma = p_max - p_min - var midscale = Float32(0.0) - if chroma != 0.0: - midscale = (p_mid - p_min) / chroma - - var f_hc = pc_min + (pc_max - pc_min) * midscale - var n_mid = (1.0 - hue_preservation) * pc_mid + hue_preservation * f_hc - - var blend = 2.0 * p_min / (p_min + p_mid + 1e-9) - var target = blend * (pc_r + pc_g + pc_b) + (1.0 - blend) * ( - pc_min + n_mid + pc_max - ) - - var res_min: Float32 - var res_mid: Float32 - var res_max: Float32 - if n_mid <= pc_mid: - res_mid = ( - (1.0 - hue_preservation) * pc_mid - + hue_preservation - * (midscale * pc_max + (1.0 - midscale) * (target - pc_max)) - ) / (1.0 + hue_preservation * (1.0 - midscale)) - res_min = target - pc_max - res_mid - res_max = pc_max - else: - res_mid = ( - (1.0 - hue_preservation) * pc_mid - + hue_preservation - * (pc_min * (1.0 - midscale) + midscale * (target - pc_min)) - ) / (1.0 + hue_preservation * midscale) - res_min = pc_min - res_max = target - pc_min - res_mid - - var res_r: Float32 - var res_g: Float32 - var res_b: Float32 - if r_r >= r_g: - if r_g >= r_b: # R G B - res_r = res_max - res_g = res_mid - res_b = res_min - elif r_b >= r_r: # B R G - res_b = res_max - res_r = res_mid - res_g = res_min - else: # R B G - res_r = res_max - res_b = res_mid - res_g = res_min - else: - if r_r >= r_b: # G R B - res_g = res_max - res_r = res_mid - res_b = res_min - elif r_b >= r_g: # B G R - res_b = res_max - res_g = res_mid - res_r = res_min - else: # G B R - res_g = res_max - res_b = res_mid - res_r = res_min - - # 6. Transform to pipe space - var out_r = ( - rendering_to_pipe[0] * res_r - + rendering_to_pipe[1] * res_g - + rendering_to_pipe[2] * res_b - ) - var out_g = ( - rendering_to_pipe[4] * res_r - + rendering_to_pipe[5] * res_g - + rendering_to_pipe[6] * res_b - ) - var out_b = ( - rendering_to_pipe[8] * res_r - + rendering_to_pipe[9] * res_g - + rendering_to_pipe[10] * res_b - ) - - return SIMD[DType.float32, 4](out_r, out_g, out_b, in_a) - - -fn run_sigmoid_rgb_ratio( - ctx: DeviceContext, - output: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], - input: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], - white_target: Float32, - black_target: Float32, - paper_exp: Float32, - film_fog: Float32, - film_power: Float32, - paper_power: Float32, - num_pixels: Int, -) raises: - @parameter - @always_inline - fn rgb_ratio_closure[ - width: Int, rank: Int, alignment: Int - ](indices: IndexList[rank]) capturing -> None: - var px_idx = indices[0] - var y = px_idx // WIDTH - var x = px_idx % WIDTH - var pix = input.load[width=4](Index(y, x, 0)) - var res = apply_sigmoid_rgb_ratio( - pix[0], - pix[1], - pix[2], - pix[3], - white_target, - black_target, - paper_exp, - film_fog, - film_power, - paper_power, - ) - output.store[width=4](Index(y, x, 0), res) - - elementwise[rgb_ratio_closure, 1, target="gpu"](num_pixels, ctx) - - -fn run_sigmoid_per_channel( - ctx: DeviceContext, - output: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], - input: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], - white_target: Float32, - paper_exp: Float32, - film_fog: Float32, - contrast_power: Float32, - skew_power: Float32, - hue_preservation: Float32, - pipe_to_base: SIMD[DType.float32, 16], - base_to_rendering: SIMD[DType.float32, 16], - rendering_to_pipe: SIMD[DType.float32, 16], - num_pixels: Int, -) raises: - @parameter - @always_inline - fn per_channel_closure[ - width: Int, rank: Int, alignment: Int - ](indices: IndexList[rank]) capturing -> None: - var px_idx = indices[0] - var y = px_idx // WIDTH - var x = px_idx % WIDTH - var pix = input.load[width=4](Index(y, x, 0)) - var res = apply_sigmoid_per_channel( - pix[0], - pix[1], - pix[2], - pix[3], - white_target, - paper_exp, - film_fog, - contrast_power, - skew_power, - hue_preservation, - pipe_to_base, - base_to_rendering, - rendering_to_pipe, - ) - output.store[width=4](Index(y, x, 0), res) - - elementwise[per_channel_closure, 1, target="gpu"](num_pixels, ctx) - - -fn main() raises: - var total_floats = HEIGHT * WIDTH * CHANNELS - var ctx = DeviceContext() - print("Using GPU API:", ctx.api()) - - var input_buffer_host = ctx.enqueue_create_host_buffer[DTYPE](total_floats) - var input_buffer_device = ctx.enqueue_create_buffer[DTYPE](total_floats) - var output_buffer_device = ctx.enqueue_create_buffer[DTYPE](total_floats) - - var input_image_host = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( - input_buffer_host - ) - var input_image_device = LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin]( - input_buffer_device - ) - var output_image_device = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( - output_buffer_device - ) - - for y in range(HEIGHT): - var h = Float32(y) / (HEIGHT - 1) * 6.0 - var segment = Int(h) - var f1 = h - Float32(segment) - var pr: Float32 - var pg: Float32 - var pb: Float32 - if segment == 0: - pr = 1.0 - pg = f1 - pb = 0.0 - elif segment == 1: - pr = 1.0 - f1 - pg = 1.0 - pb = 0.0 - elif segment == 2: - pr = 0.0 - pg = 1.0 - pb = f1 - elif segment == 3: - pr = 0.0 - pg = 1.0 - f1 - pb = 1.0 - elif segment == 4: - pr = f1 - pg = 0.0 - pb = 1.0 - elif segment == 5: - pr = 1.0 - pg = 0.0 - pb = 1.0 - f1 - else: - pr = 1.0 - pg = 0.0 - pb = 0.0 - - for x in range(WIDTH): - var r: Float32 - var g: Float32 - var b: Float32 - var mid_x = Float32(WIDTH) / 2.0 - if Float32(x) < mid_x: - var t = Float32(x) / mid_x - r = pr * t - g = pg * t - b = pb * t - else: - var t = (Float32(x) - mid_x) / (Float32(WIDTH) - 1.0 - mid_x) - r = pr * (1.0 - t) + t - g = pg * (1.0 - t) + t - b = pb * (1.0 - t) + t - input_image_host.store[width=4]( - Index(y, x, 0), SIMD[DType.float32, 4](r, g, b, 1.0) - ) - - ctx.enqueue_copy(input_buffer_device, input_buffer_host) - - var white_target = Float32(1.0) - var black_target = Float32(0.000152) - var paper_exp = Float32(0.5) - var film_fog = Float32(0.0) - var contrast_power = Float32(2.5) - var skew_power = Float32(1.0) - var hue_preservation = Float32(1.0) - - var identity = SIMD[DType.float32, 16](0) - identity[0] = 1 - identity[5] = 1 - identity[10] = 1 - identity[15] = 1 - - # Runs - run_sigmoid_rgb_ratio( - ctx, - output_image_device, - input_image_device, - white_target, - black_target, - paper_exp, - film_fog, - contrast_power, - skew_power, - WIDTH * HEIGHT, - ) - ctx.synchronize() - - var output_buffer_host = ctx.enqueue_create_host_buffer[DTYPE](total_floats) - ctx.enqueue_copy(output_buffer_host, output_buffer_device) - ctx.synchronize() - var output_image_host = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( - output_buffer_host - ) - - for i in range(1, 4): - var t = Float32(i) * 0.2 - var x = Int(t * Float32(WIDTH)) - var y = Int(t * Float32(HEIGHT)) - var res = output_image_host.load[width=4](Index(y, x, 0)) - print( - "RGB Ratio -", - Int(t * 100), - "% pixel: [", - res[0], - res[1], - res[2], - res[3], - "]", - ) - - run_sigmoid_per_channel( - ctx, - output_image_device, - input_image_device, - white_target, - paper_exp, - film_fog, - contrast_power, - skew_power, - hue_preservation, - identity, - identity, - identity, - WIDTH * HEIGHT, - ) - ctx.synchronize() - - ctx.enqueue_copy(output_buffer_host, output_buffer_device) - ctx.synchronize() - - for i in range(1, 4): - var t = Float32(i) * 0.2 - var x = Int(t * Float32(WIDTH)) - var y = Int(t * Float32(HEIGHT)) - var res = output_image_host.load[width=4](Index(y, x, 0)) - print( - "Per Channel -", - Int(t * 100), - "% pixel: [", - res[0], - res[1], - res[2], - res[3], - "]", - ) - - # Benchmarking - var bench = Bench(BenchConfig(max_iters=1000, num_warmup_iters=100)) - - @parameter - fn bench_rgb(mut b: Bencher) raises: - @parameter - fn run(ctx: DeviceContext) raises: - run_sigmoid_rgb_ratio( - ctx, - output_image_device, - input_image_device, - white_target, - black_target, - paper_exp, - film_fog, - contrast_power, - skew_power, - WIDTH * HEIGHT, - ) - - b.iter_custom[run](ctx) - ctx.synchronize() - - @parameter - fn bench_per(mut b: Bencher) raises: - @parameter - fn run(ctx: DeviceContext) raises: - run_sigmoid_per_channel( - ctx, - output_image_device, - input_image_device, - white_target, - paper_exp, - film_fog, - contrast_power, - skew_power, - hue_preservation, - identity, - identity, - identity, - WIDTH * HEIGHT, - ) - - b.iter_custom[run](ctx) - ctx.synchronize() - - bench.bench_function[bench_rgb]( - BenchId("Mojo-Sigmoid-GPU-RGB-Ratio-V2"), fixed_iterations=1000 - ) - bench.bench_function[bench_per]( - BenchId("Mojo-Sigmoid-GPU-Per-Channel-V2"), fixed_iterations=1000 - ) - print(bench) diff --git a/mojo/.gitattributes b/mojo/.gitattributes new file mode 100644 index 0000000000..997504b465 --- /dev/null +++ b/mojo/.gitattributes @@ -0,0 +1,2 @@ +# SCM syntax highlighting & preventing 3-way merges +pixi.lock merge=binary linguist-language=YAML linguist-generated=true -diff diff --git a/mojo/.gitignore b/mojo/.gitignore new file mode 100644 index 0000000000..ae849e65b8 --- /dev/null +++ b/mojo/.gitignore @@ -0,0 +1,3 @@ +# pixi environments +.pixi/* +!.pixi/config.toml diff --git a/mojo/Makefile b/mojo/Makefile new file mode 100644 index 0000000000..c8648a298b --- /dev/null +++ b/mojo/Makefile @@ -0,0 +1,17 @@ +MOJO = pixi run mojo +PLUGIN_DIR = /usr/lib/darktable/plugins + +.PHONY: all build install bench clean + +all: build + +build: libsigmoid_mojo.so + +libsigmoid_mojo.so: iop/sigmoid/lib.mojo iop/sigmoid/kernels.mojo + $(MOJO) build -I . iop/sigmoid/lib.mojo --emit shared-lib -o libsigmoid_mojo.so + +install: libsigmoid_mojo.so + sudo cp libsigmoid_mojo.so $(PLUGIN_DIR)/ + +clean: + rm -f libsigmoid_mojo.so diff --git a/mojo/build_sigmoid_iop.sh b/mojo/build_sigmoid_iop.sh new file mode 100755 index 0000000000..0c06170049 --- /dev/null +++ b/mojo/build_sigmoid_iop.sh @@ -0,0 +1,51 @@ +#!/bin/bash +# Build libsigmoid.so from the darktable source tree. +# This mimics what CMake does (run introspection, then compile + link). + +set -e + +SRC=/home/mc/code/darktable +BUILD=/tmp/sigmoid_build +DEST=$BUILD/libsigmoid.so + +mkdir -p "$BUILD" + +# Step 1: Run introspection (generates introspection_sigmoid.c) +echo "=== Step 1: Introspection ===" +perl "$SRC/tools/introspection/parser.pl" \ + "$SRC/src/" \ + "$SRC/src/iop/sigmoid.c" \ + "$BUILD/introspection_sigmoid.c" +echo " -> $BUILD/introspection_sigmoid.c ($(wc -l < "$BUILD/introspection_sigmoid.c") lines)" + +# Step 2: Collect flags +GTK_FLAGS=$(pkg-config --cflags gtk+-3.0 glib-2.0 librsvg-2.0 lcms2 json-glib-1.0) +DT_INCLUDE="-I$SRC/src -I$SRC/src/iop -I$SRC/build/bin -I/usr/lib/darktable" +DEFINES="-DHAVE_CONFIG_H -DHAVE_OPENCL -D_GNU_SOURCE -include common/module_api.h -include iop/iop_api.h" +CFLAGS="-O3 -march=native -fPIC -fvisibility=hidden -fopenmp $GTK_FLAGS $DT_INCLUDE $DEFINES" +LDFLAGS="-L/usr/lib/darktable -ldarktable -lm -lgomp -Wl,-rpath,/usr/lib/darktable" + +echo "" +echo "=== Step 2: Compile ===" +gcc $CFLAGS \ + -Wno-unused-function \ + -Wno-deprecated-declarations \ + -c "$BUILD/introspection_sigmoid.c" \ + -o "$BUILD/introspection_sigmoid.o" 2>&1 +echo " -> $BUILD/introspection_sigmoid.o" + +echo "" +echo "=== Step 3: Link ===" +gcc -shared -fPIC -fopenmp \ + "$BUILD/introspection_sigmoid.o" \ + $LDFLAGS \ + -o "$DEST" +echo " -> $DEST" + +echo "" +echo "=== Step 4: Verify ===" +nm -D "$DEST" | grep " T " | awk '{print $3}' + +echo "" +echo "Build complete: $DEST" +echo "Size: $(du -sh "$DEST" | cut -f1)" diff --git a/mojo/iop/__init__.mojo b/mojo/iop/__init__.mojo new file mode 100644 index 0000000000..0c393fc1b4 --- /dev/null +++ b/mojo/iop/__init__.mojo @@ -0,0 +1 @@ +# iop package diff --git a/mojo/iop/blurs/__init__.mojo b/mojo/iop/blurs/__init__.mojo new file mode 100644 index 0000000000..8abcdf557a --- /dev/null +++ b/mojo/iop/blurs/__init__.mojo @@ -0,0 +1 @@ +# iop.blurs package diff --git a/mojo/iop/blurs/kernels.mojo b/mojo/iop/blurs/kernels.mojo new file mode 100644 index 0000000000..0898c38d54 --- /dev/null +++ b/mojo/iop/blurs/kernels.mojo @@ -0,0 +1,59 @@ +from layout import Layout, LayoutTensor, UNKNOWN_VALUE +from std.math import clamp +from std.utils import Index +from std.memory.unsafe_pointer import UnsafePointer + +comptime CHANNELS = 4 +comptime IMAGE_LAYOUT = Layout.row_major(UNKNOWN_VALUE, UNKNOWN_VALUE, CHANNELS) +comptime DTYPE = DType.float32 + + +@always_inline +def apply_convolve_vector[W: Int]( + x: Int, + y: Int, + width: Int, + height: Int, + radius: Int, + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], + kern_ptr: UnsafePointer[Float32, ImmutAnyOrigin], +) -> SIMD[DType.float32, W * 4]: + var acc = SIMD[DType.float32, W * 4](0) + + for l in range(-radius, radius + 1): + var ii = clamp(y + l, 0, height - 1) + for m in range(-radius, radius + 1): + var ik = l + radius + var jk = m + radius + var k_width = 2 * radius + 1 + var k = kern_ptr[ik * k_width + jk] + + for i in range(W): + var jj = clamp(x + i + m, 0, width - 1) + var pix = in_t.load[width=4](Index(ii, jj, 0)) + + acc[i * 4] += k * pix[0] + acc[i * 4 + 1] += k * pix[1] + acc[i * 4 + 2] += k * pix[2] + acc[i * 4 + 3] += k * pix[3] + + for i in range(W): + var curr_x = x + i + if curr_x < width: + var orig = in_t.load[width=4](Index(y, curr_x, 0)) + acc[i * 4 + 3] = orig[3] + + return acc + + +@always_inline +def apply_convolve( + x: Int, + y: Int, + width: Int, + height: Int, + radius: Int, + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], + kern_ptr: UnsafePointer[Float32, ImmutAnyOrigin], +) -> SIMD[DType.float32, 4]: + return apply_convolve_vector[1](x, y, width, height, radius, in_t, kern_ptr) diff --git a/mojo/iop/blurs/lib.mojo b/mojo/iop/blurs/lib.mojo new file mode 100644 index 0000000000..c8e966add0 --- /dev/null +++ b/mojo/iop/blurs/lib.mojo @@ -0,0 +1,190 @@ +from std.gpu.host import DeviceContext +from layout import Layout, LayoutTensor, UNKNOWN_VALUE +from layout.runtime_layout import RuntimeLayout, IndexList +from layout.coord import Coord +from std.utils import Index +from std.algorithm.functional import elementwise, vectorize +from std.gpu.host.compile import get_gpu_target +from std.memory import alloc +from std.memory.unsafe_pointer import UnsafePointer +from std.sys import simd_width_of +from iop.blurs.kernels import apply_convolve, apply_convolve_vector + +comptime CHANNELS = 4 +comptime IMAGE_LAYOUT = Layout.row_major(UNKNOWN_VALUE, UNKNOWN_VALUE, CHANNELS) +comptime DTYPE = DType.float32 +comptime SIMD_WIDTH = simd_width_of[DTYPE, target=get_gpu_target()]() + + +struct MojoCtx: + var use_gpu: Int + var dctx_addr: Int + + def __init__(out self, use_gpu: Int, dctx_addr: Int): + self.use_gpu = use_gpu + self.dctx_addr = dctx_addr + + +def _launch_convolve_gpu( + dctx: DeviceContext, + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], + out_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + kern_ptr: UnsafePointer[Float32, ImmutAnyOrigin], + width: Int, + height: Int, + radius: Int, + num_pixels: Int, +) raises: + @parameter + @always_inline + def gpu_kernel[simd_width: Int, alignment: Int](coord: Coord) capturing -> None: + var px_idx = Int(coord[0].value()) + var y = px_idx // width + var x = px_idx % width + var res = apply_convolve_vector[1]( + x, y, width, height, radius, in_t, kern_ptr + ) + out_t.store[width=4](Index(y, x, 0), res) + + elementwise[gpu_kernel, 1, target="gpu"](num_pixels, dctx) + dctx.synchronize() + + +@export("blurs_mojo_init") +def blurs_mojo_init( + ctx_out: UnsafePointer[Int, MutAnyOrigin], use_gpu: Int32 +) abi("C") -> None: + var gpu = use_gpu != 0 + var dctx_addr = 0 + if gpu: + try: + var d_ptr = alloc[DeviceContext](1) + d_ptr.unsafe_write(DeviceContext()) + dctx_addr = Int(d_ptr) + print("Mojo Blurs: GPU Context Initialized") + except e: + print("Mojo Blurs: GPU Init Error:", String(e)) + gpu = False + else: + print("Mojo Blurs: CPU Context Initialized") + var p = alloc[MojoCtx](1) + p[0].use_gpu = 1 if gpu else 0 + p[0].dctx_addr = dctx_addr + ctx_out[0] = Int(p) + + +@export("blurs_mojo_destroy") +def blurs_mojo_destroy(ctx_addr: Int) abi("C") -> None: + var p = UnsafePointer[MojoCtx, MutAnyOrigin](unsafe_from_address=ctx_addr) + if p[0].dctx_addr != 0: + var dctx_ptr = UnsafePointer[DeviceContext, MutAnyOrigin]( + unsafe_from_address=p[0].dctx_addr + ) + dctx_ptr.unsafe_deinit_pointee() + dctx_ptr.free() + p.free() + + +@export("blurs_mojo_convolve") +def blurs_mojo_convolve( + ctx_addr: Int, + in_addr: Int, + kern_addr: Int, + out_addr: Int, + width: Int32, + height: Int32, + radius: Int32, +) abi("C") -> None: + var ctx_p = UnsafePointer[MojoCtx, MutAnyOrigin]( + unsafe_from_address=ctx_addr + ) + var use_gpu = ctx_p[0].use_gpu != 0 + var dctx_addr = ctx_p[0].dctx_addr + + var kern_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=kern_addr + ) + var out_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=out_addr + ) + + var h = Int(height) + var w = Int(width) + var r = Int(radius) + var num_pixels = h * w + var k_size = (2 * r + 1) * (2 * r + 1) + + var rt = RuntimeLayout[IMAGE_LAYOUT].row_major( + IndexList[3](h, w, CHANNELS) + ) + + if use_gpu and dctx_addr != 0: + try: + var in_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=in_addr + ) + var dctx = UnsafePointer[DeviceContext, MutAnyOrigin]( + unsafe_from_address=dctx_addr + )[0] + var dev_in = dctx.enqueue_create_buffer[DTYPE](num_pixels * 4) + var dev_kern = dctx.enqueue_create_buffer[DTYPE](k_size) + var dev_out = dctx.enqueue_create_buffer[DTYPE](num_pixels * 4) + + dctx.enqueue_copy(dev_in, in_p) + dctx.enqueue_copy(dev_kern, kern_p) + + var in_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin]( + UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=Int(dev_in.unsafe_ptr()) + ), + rt, + ) + var out_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=Int(dev_out.unsafe_ptr()) + ), + rt, + ) + var kern_ptr_gpu = UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=Int(dev_kern.unsafe_ptr()) + ) + + _launch_convolve_gpu( + dctx, in_t, out_t, kern_ptr_gpu, w, h, r, num_pixels, + ) + + dctx.enqueue_copy(out_p, dev_out) + dctx.synchronize() + except e: + print("GPU Run Error (Convolve):", String(e)) + else: + var in_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin]( + UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=in_addr + ), + rt, + ) + var out_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + out_p, + rt, + ) + var kern_ptr_cpu = UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=kern_addr + ) + + for y in range(h): + var cur_y = y + def row_fn[width: Int](x: Int) {mut}: + var res = apply_convolve_vector[width]( + x, cur_y, w, h, r, in_t, kern_ptr_cpu + ) + for i in range(width): + var pixel = SIMD[DType.float32, 4]( + res[i * 4], + res[i * 4 + 1], + res[i * 4 + 2], + res[i * 4 + 3], + ) + out_t.store[width=4](Index(cur_y, x + i, 0), pixel) + + vectorize[SIMD_WIDTH](w, row_fn) diff --git a/mojo/iop/sigmoid/__init__.mojo b/mojo/iop/sigmoid/__init__.mojo new file mode 100644 index 0000000000..256cd8163f --- /dev/null +++ b/mojo/iop/sigmoid/__init__.mojo @@ -0,0 +1 @@ +# iop.sigmoid package diff --git a/mojo/iop/sigmoid/kernels.mojo b/mojo/iop/sigmoid/kernels.mojo new file mode 100644 index 0000000000..8bfeb61285 --- /dev/null +++ b/mojo/iop/sigmoid/kernels.mojo @@ -0,0 +1,204 @@ +from std.math import isnan, sqrt, pow, max, min + +@always_inline +def generalized_loglogistic_sigmoid_scalar( + value: Float32, + magnitude: Float32, + paper_exp: Float32, + film_fog: Float32, + film_power: Float32, + paper_power: Float32, +) -> Float32: + var clamped_value = max(value, Float32(0.0)) + var film_response = pow(film_fog + clamped_value, film_power) + var paper_response = magnitude * pow( + film_response / (paper_exp + film_response), paper_power + ) + if isnan(paper_response): + return magnitude + return paper_response + +@always_inline +def apply_sigmoid_rgb_ratio( + in_r: Float32, + in_g: Float32, + in_b: Float32, + in_a: Float32, + white_target: Float32, + black_target: Float32, + paper_exp: Float32, + film_fog: Float32, + film_power: Float32, + paper_power: Float32, +) -> SIMD[DType.float32, 4]: + var avg = max((in_r + in_g + in_b) / Float32(3.0), Float32(0.0)) + var min_v = min(min(in_r, in_g), in_b) + var sat = Float32(1.0) + if min_v < Float32(0.0): + sat = -avg / (min_v - avg) + + var p_r = avg + sat * (in_r - avg) + var p_g = avg + sat * (in_g - avg) + var p_b = avg + sat * (in_b - avg) + + var luma = (p_r + p_g + p_b) / Float32(3.0) + var mapped_luma = generalized_loglogistic_sigmoid_scalar( + luma, white_target, paper_exp, film_fog, film_power, paper_power + ) + + if luma > Float32(1e-9): + var scale = mapped_luma / luma + p_r *= scale + p_g *= scale + p_b *= scale + else: + p_r = mapped_luma + p_g = mapped_luma + p_b = mapped_luma + + var p_min = min(min(p_r, p_g), p_b) + var p_max = max(max(p_r, p_g), p_b) + var eps = Float32(1e-6) + var d_white = (white_target - mapped_luma) / (p_max - mapped_luma + eps) + var d_black = (black_target - mapped_luma) / (p_min - mapped_luma - eps) + var db_vs_chroma = min(d_white, d_black) + var cvm_border = (mapped_luma - p_min) / (mapped_luma + eps) + var p_chr_adj = Float32(1.0) / (cvm_border * db_vs_chroma + eps) + var h_chr = (Float32(2.0) * cvm_border / (Float32(1.0) - cvm_border * cvm_border + eps)) * p_chr_adj + var h_z = sqrt(h_chr * h_chr + Float32(1.0)) + var chroma_f = h_chr / (Float32(1.0) + h_z) * db_vs_chroma + + return SIMD[DType.float32, 4]( + mapped_luma + chroma_f * (p_r - mapped_luma), + mapped_luma + chroma_f * (p_g - mapped_luma), + mapped_luma + chroma_f * (p_b - mapped_luma), + in_a, + ) + +@always_inline +def apply_sigmoid_per_channel( + in_r: Float32, + in_g: Float32, + in_b: Float32, + in_a: Float32, + white_target: Float32, + paper_exp: Float32, + film_fog: Float32, + contrast_power: Float32, + skew_power: Float32, + hue_preservation: Float32, + pipe_to_base: SIMD[DType.float32, 16], + base_to_rendering: SIMD[DType.float32, 16], + rendering_to_pipe: SIMD[DType.float32, 16], +) -> SIMD[DType.float32, 4]: + var i_r = pipe_to_base[0] * in_r + pipe_to_base[1] * in_g + pipe_to_base[2] * in_b + var i_g = pipe_to_base[4] * in_r + pipe_to_base[5] * in_g + pipe_to_base[6] * in_b + var i_b = pipe_to_base[8] * in_r + pipe_to_base[9] * in_g + pipe_to_base[10] * in_b + + var avg = max((i_r + i_g + i_b) / Float32(3.0), Float32(0.0)) + var min_v = min(min(i_r, i_g), i_b) + var sat = Float32(1.0) + if min_v < Float32(0.0): + sat = -avg / (min_v - avg) + i_r = avg + sat * (i_r - avg) + i_g = avg + sat * (i_g - avg) + i_b = avg + sat * (i_b - avg) + + var r_r = base_to_rendering[0] * i_r + base_to_rendering[1] * i_g + base_to_rendering[2] * i_b + var r_g = base_to_rendering[4] * i_r + base_to_rendering[5] * i_g + base_to_rendering[6] * i_b + var r_b = base_to_rendering[8] * i_r + base_to_rendering[9] * i_g + base_to_rendering[10] * i_b + + var pc_r = generalized_loglogistic_sigmoid_scalar( + r_r, white_target, paper_exp, film_fog, contrast_power, skew_power + ) + var pc_g = generalized_loglogistic_sigmoid_scalar( + r_g, white_target, paper_exp, film_fog, contrast_power, skew_power + ) + var pc_b = generalized_loglogistic_sigmoid_scalar( + r_b, white_target, paper_exp, film_fog, contrast_power, skew_power + ) + + var p_min: Float32 + var p_mid: Float32 + var p_max: Float32 + var pc_min: Float32 + var pc_mid: Float32 + var pc_max: Float32 + + if r_r >= r_g: + if r_g >= r_b: + p_max = r_r; p_mid = r_g; p_min = r_b + pc_max = pc_r; pc_mid = pc_g; pc_min = pc_b + elif r_b >= r_r: + p_max = r_b; p_mid = r_r; p_min = r_g + pc_max = pc_b; pc_mid = pc_r; pc_min = pc_g + else: + p_max = r_r; p_mid = r_b; p_min = r_g + pc_max = pc_r; pc_mid = pc_b; pc_min = pc_g + else: + if r_r >= r_b: + p_max = r_g; p_mid = r_r; p_min = r_b + pc_max = pc_g; pc_mid = pc_r; pc_min = pc_b + elif r_b >= r_g: + p_max = r_b; p_mid = r_g; p_min = r_r + pc_max = pc_b; pc_mid = pc_g; pc_min = pc_r + else: + p_max = r_g; p_mid = r_b; p_min = r_r + pc_max = pc_g; pc_mid = pc_b; pc_min = pc_r + + var chroma = p_max - p_min + var midscale = Float32(0.0) + if chroma != Float32(0.0): + midscale = (p_mid - p_min) / chroma + + var f_hc = pc_min + (pc_max - pc_min) * midscale + var n_mid = (Float32(1.0) - hue_preservation) * pc_mid + hue_preservation * f_hc + + var blend = Float32(2.0) * p_min / (p_min + p_mid + Float32(1e-9)) + var target = blend * (pc_r + pc_g + pc_b) + (Float32(1.0) - blend) * ( + pc_min + n_mid + pc_max + ) + + var res_min: Float32 + var res_mid: Float32 + var res_max: Float32 + if n_mid <= pc_mid: + res_mid = ( + (Float32(1.0) - hue_preservation) * pc_mid + + hue_preservation + * (midscale * pc_max + (Float32(1.0) - midscale) * (target - pc_max)) + ) / (Float32(1.0) + hue_preservation * (Float32(1.0) - midscale)) + res_min = target - pc_max - res_mid + res_max = pc_max + else: + res_mid = ( + (Float32(1.0) - hue_preservation) * pc_mid + + hue_preservation + * (pc_min * (Float32(1.0) - midscale) + midscale * (target - pc_min)) + ) / (Float32(1.0) + hue_preservation * midscale) + res_min = pc_min + res_max = target - pc_min - res_mid + + var res_r: Float32 + var res_g: Float32 + var res_b: Float32 + if r_r >= r_g: + if r_g >= r_b: + res_r = res_max; res_g = res_mid; res_b = res_min + elif r_b >= r_r: + res_b = res_max; res_r = res_mid; res_g = res_min + else: + res_r = res_max; res_b = res_mid; res_g = res_min + else: + if r_r >= r_b: + res_g = res_max; res_r = res_mid; res_b = res_min + elif r_b >= r_g: + res_b = res_max; res_g = res_mid; res_r = res_min + else: + res_g = res_max; res_b = res_mid; res_r = res_min + + var out_r = rendering_to_pipe[0] * res_r + rendering_to_pipe[1] * res_g + rendering_to_pipe[2] * res_b + var out_g = rendering_to_pipe[4] * res_r + rendering_to_pipe[5] * res_g + rendering_to_pipe[6] * res_b + var out_b = rendering_to_pipe[8] * res_r + rendering_to_pipe[9] * res_g + rendering_to_pipe[10] * res_b + + return SIMD[DType.float32, 4](out_r, out_g, out_b, in_a) diff --git a/mojo/iop/sigmoid/lib.mojo b/mojo/iop/sigmoid/lib.mojo new file mode 100644 index 0000000000..3b9191e0f6 --- /dev/null +++ b/mojo/iop/sigmoid/lib.mojo @@ -0,0 +1,432 @@ +from std.gpu.host import DeviceContext +from layout import Layout, LayoutTensor, UNKNOWN_VALUE +from layout.runtime_layout import RuntimeLayout, IndexList +from layout.coord import Coord +from std.utils import Index +from std.algorithm.functional import elementwise +from std.memory import alloc +from std.memory.unsafe_pointer import UnsafePointer +from iop.sigmoid.kernels import ( + apply_sigmoid_rgb_ratio, + apply_sigmoid_per_channel, +) + +comptime CHANNELS = 4 +comptime IMAGE_LAYOUT = Layout.row_major(UNKNOWN_VALUE, UNKNOWN_VALUE, CHANNELS) +comptime DTYPE = DType.float32 + + +struct MojoCtx: + var use_gpu: Int + var dctx_addr: Int + + def __init__(out self, use_gpu: Int, dctx_addr: Int): + self.use_gpu = use_gpu + self.dctx_addr = dctx_addr + + +struct CParams: + var wt: Float32 + var bt: Float32 + var pe: Float32 + var ff: Float32 + var fp: Float32 + var pp: Float32 + var hp: Float32 + var ptb: SIMD[DType.float32, 16] + var btr: SIMD[DType.float32, 16] + var rtp: SIMD[DType.float32, 16] + + def __init__(out self, addr: Int): + var p = UnsafePointer[Float32, MutAnyOrigin](unsafe_from_address=addr) + self.wt = p[0] + self.bt = p[4] + self.pe = p[8] + self.ff = p[12] + self.fp = p[16] + self.pp = p[20] + self.hp = p[32] + var m_ptb = SIMD[DType.float32, 16]() + var m_btr = SIMD[DType.float32, 16]() + var m_rtp = SIMD[DType.float32, 16]() + for i in range(16): + m_ptb[i] = p[36 + i] + m_btr[i] = p[52 + i] + m_rtp[i] = p[68 + i] + self.ptb = m_ptb + self.btr = m_btr + self.rtp = m_rtp + + +def _launch_rgb_ratio_gpu( + dctx: DeviceContext, + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], + out_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + wt: Float32, + bt: Float32, + pe: Float32, + ff: Float32, + fp: Float32, + pp: Float32, + img_w: Int, + img_h: Int, + num_pixels: Int, +) raises: + @parameter + @always_inline + def gpu_kernel[simd_width: Int, alignment: Int](coord: Coord) capturing -> None: + var px_idx = Int(coord[0].value()) + var y = px_idx // img_w + var x = px_idx % img_w + var pix = in_t.load[width=4](Index(y, x, 0)) + var r = apply_sigmoid_rgb_ratio( + pix[0], + pix[1], + pix[2], + pix[3], + wt, + bt, + pe, + ff, + fp, + pp, + ) + out_t.store[width=4](Index(y, x, 0), r) + + elementwise[gpu_kernel, 1, target="gpu"](num_pixels, dctx) + dctx.synchronize() + + +def _launch_per_channel_gpu( + dctx: DeviceContext, + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin], + out_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + wt: Float32, + pe: Float32, + ff: Float32, + fp: Float32, + pp: Float32, + hp: Float32, + kptb: SIMD[DType.float32, 16], + kbtr: SIMD[DType.float32, 16], + krtp: SIMD[DType.float32, 16], + img_w: Int, + img_h: Int, + num_pixels: Int, +) raises: + @parameter + @always_inline + def gpu_kernel[simd_width: Int, alignment: Int](coord: Coord) capturing -> None: + var px_idx = Int(coord[0].value()) + var y = px_idx // img_w + var x = px_idx % img_w + var pix = in_t.load[width=4](Index(y, x, 0)) + var r = apply_sigmoid_per_channel( + pix[0], + pix[1], + pix[2], + pix[3], + wt, + pe, + ff, + fp, + pp, + hp, + kptb, + kbtr, + krtp, + ) + out_t.store[width=4](Index(y, x, 0), r) + + elementwise[gpu_kernel, 1, target="gpu"](num_pixels, dctx) + dctx.synchronize() + + +@export("sigmoid_mojo_init") +def sigmoid_mojo_init(ctx_out: UnsafePointer[Int, MutAnyOrigin], use_gpu: Int32) abi("C") -> None: + var gpu = use_gpu != 0 + var dctx_addr = 0 + if gpu: + try: + var d_ptr = alloc[DeviceContext](1) + d_ptr.unsafe_write(DeviceContext()) + dctx_addr = Int(d_ptr) + print("Mojo: GPU Context Initialized Successfully") + except e: + print("Mojo: GPU Init Error (falling back to CPU):", String(e)) + gpu = False + else: + print("Mojo: CPU Context Initialized") + var p = alloc[MojoCtx](1) + p[0].use_gpu = 1 if gpu else 0 + p[0].dctx_addr = dctx_addr + ctx_out[0] = Int(p) + + +@export("sigmoid_mojo_destroy") +def sigmoid_mojo_destroy(ctx_addr: Int) abi("C") -> None: + var p = UnsafePointer[MojoCtx, MutAnyOrigin](unsafe_from_address=ctx_addr) + if p[0].dctx_addr != 0: + var dctx_ptr = UnsafePointer[DeviceContext, MutAnyOrigin]( + unsafe_from_address=p[0].dctx_addr + ) + dctx_ptr.unsafe_deinit_pointee() + dctx_ptr.free() + p.free() + + +def _run_cpu_rgb_ratio( + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + out_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + wt: Float32, + bt: Float32, + pe: Float32, + ff: Float32, + fp: Float32, + pp: Float32, + width: Int, + height: Int, + num_pixels: Int, +): + for i in range(num_pixels): + var px = in_t.load[width=4](Index(i // width, i % width, 0)) + var r = apply_sigmoid_rgb_ratio( + px[0], px[1], px[2], px[3], wt, bt, pe, ff, fp, pp + ) + out_t.store[width=4](Index(i // width, i % width, 0), r) + + +def _run_cpu_per_channel( + in_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + out_t: LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin], + wt: Float32, + pe: Float32, + ff: Float32, + fp: Float32, + pp: Float32, + hp: Float32, + ptb: SIMD[DType.float32, 16], + btr: SIMD[DType.float32, 16], + rtp: SIMD[DType.float32, 16], + width: Int, + height: Int, + num_pixels: Int, +): + for i in range(num_pixels): + var px = in_t.load[width=4](Index(i // width, i % width, 0)) + var r = apply_sigmoid_per_channel( + px[0], + px[1], + px[2], + px[3], + wt, + pe, + ff, + fp, + pp, + hp, + ptb, + btr, + rtp, + ) + out_t.store[width=4](Index(i // width, i % width, 0), r) + + +@export("sigmoid_mojo_rgb_ratio") +def sigmoid_mojo_rgb_ratio( + ctx_addr: Int, + in_addr: Int, + out_addr: Int, + width: Int32, + height: Int32, + p_addr: Int, +) abi("C") -> None: + var ctx_p = UnsafePointer[MojoCtx, MutAnyOrigin]( + unsafe_from_address=ctx_addr + ) + var use_gpu = ctx_p[0].use_gpu != 0 + var dctx_addr = ctx_p[0].dctx_addr + var params = CParams(p_addr) + var h = Int(height) + var w = Int(width) + var num_pixels = h * w + + if use_gpu and dctx_addr != 0: + try: + var dctx = UnsafePointer[DeviceContext, MutAnyOrigin]( + unsafe_from_address=dctx_addr + )[0] + var dev_in = dctx.enqueue_create_buffer[DTYPE](num_pixels * 4) + var dev_out = dctx.enqueue_create_buffer[DTYPE](num_pixels * 4) + var in_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=in_addr + ) + dctx.enqueue_copy(dev_in, in_p) + + var rt_gpu = RuntimeLayout[IMAGE_LAYOUT].row_major( + IndexList[3](h, w, CHANNELS) + ) + var in_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin]( + UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=Int(dev_in.unsafe_ptr()) + ), + rt_gpu, + ) + var out_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=Int(dev_out.unsafe_ptr()) + ), + rt_gpu, + ) + + _launch_rgb_ratio_gpu( + dctx, + in_t, + out_t, + params.wt, + params.bt, + params.pe, + params.ff, + params.fp, + params.pp, + w, + h, + num_pixels, + ) + + var out_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=out_addr + ) + dctx.enqueue_copy(out_p, dev_out) + dctx.synchronize() + except e: + print("GPU Run Error (RGB Ratio):", String(e)) + else: + var rt = RuntimeLayout[IMAGE_LAYOUT].row_major( + IndexList[3](h, w, CHANNELS) + ) + var in_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=in_addr + ) + var out_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=out_addr + ) + var in_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin](in_p, rt) + var out_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin](out_p, rt) + + _run_cpu_rgb_ratio( + in_t, + out_t, + params.wt, + params.bt, + params.pe, + params.ff, + params.fp, + params.pp, + w, + h, + num_pixels, + ) + + +@export("sigmoid_mojo_per_channel") +def sigmoid_mojo_per_channel( + ctx_addr: Int, + in_addr: Int, + out_addr: Int, + width: Int32, + height: Int32, + p_addr: Int, +) abi("C") -> None: + var ctx_p = UnsafePointer[MojoCtx, MutAnyOrigin]( + unsafe_from_address=ctx_addr + ) + var use_gpu = ctx_p[0].use_gpu != 0 + var dctx_addr = ctx_p[0].dctx_addr + var params = CParams(p_addr) + var h = Int(height) + var w = Int(width) + var num_pixels = h * w + + if use_gpu and dctx_addr != 0: + try: + var dctx = UnsafePointer[DeviceContext, MutAnyOrigin]( + unsafe_from_address=dctx_addr + )[0] + var dev_in = dctx.enqueue_create_buffer[DTYPE](num_pixels * 4) + var dev_out = dctx.enqueue_create_buffer[DTYPE](num_pixels * 4) + var in_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=in_addr + ) + dctx.enqueue_copy(dev_in, in_p) + + var rt_gpu = RuntimeLayout[IMAGE_LAYOUT].row_major( + IndexList[3](h, w, CHANNELS) + ) + var in_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, ImmutAnyOrigin]( + UnsafePointer[Float32, ImmutAnyOrigin]( + unsafe_from_address=Int(dev_in.unsafe_ptr()) + ), + rt_gpu, + ) + var out_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin]( + UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=Int(dev_out.unsafe_ptr()) + ), + rt_gpu, + ) + + _launch_per_channel_gpu( + dctx, + in_t, + out_t, + params.wt, + params.pe, + params.ff, + params.fp, + params.pp, + params.hp, + params.ptb, + params.btr, + params.rtp, + w, + h, + num_pixels, + ) + + var out_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=out_addr + ) + dctx.enqueue_copy(out_p, dev_out) + dctx.synchronize() + except e: + print("GPU Error (Per Channel):", String(e)) + else: + var rt = RuntimeLayout[IMAGE_LAYOUT].row_major( + IndexList[3](h, w, CHANNELS) + ) + var in_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=in_addr + ) + var out_p = UnsafePointer[Float32, MutAnyOrigin]( + unsafe_from_address=out_addr + ) + var in_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin](in_p, rt) + var out_t = LayoutTensor[DTYPE, IMAGE_LAYOUT, MutAnyOrigin](out_p, rt) + + _run_cpu_per_channel( + in_t, + out_t, + params.wt, + params.pe, + params.ff, + params.fp, + params.pp, + params.hp, + params.ptb, + params.btr, + params.rtp, + w, + h, + num_pixels, + ) diff --git a/benchmark/pixi.lock b/mojo/pixi.lock similarity index 51% rename from benchmark/pixi.lock rename to mojo/pixi.lock index 38ce7ccdb5..f4dd8df35b 100644 --- a/benchmark/pixi.lock +++ b/mojo/pixi.lock @@ -1,82 +1,87 @@ -version: 6 +version: 7 +platforms: +- name: linux-64 + virtual-packages: + - __unix=0=0 + - __linux=4.18 + - __glibc=2.28 + - __archspec=0=x86_64 environments: default: channels: - url: https://conda.modular.com/max-nightly/ - url: https://conda.anaconda.org/conda-forge/ - options: - pypi-prerelease-mode: if-necessary-or-explicit packages: linux-64: - conda: https://conda.anaconda.org/conda-forge/linux-64/_openmp_mutex-4.5-20_gnu.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/_python_abi3_support-1.0-hd8ed1ab_2.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/bzip2-1.0.8-hda65f42_8.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/ca-certificates-2026.1.4-hbd8a1cb_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/click-8.3.1-pyh8f84b5b_1.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/cpython-3.14.3-py314hd8ed1ab_101.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/icu-78.2-h33c6efd_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/importlib-metadata-8.7.0-pyhe01879c_1.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_client-8.6.3-pyhd8ed1ab_1.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_core-5.9.1-pyhc90fa1f_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/bzip2-1.0.8-hda65f42_9.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/icu-78.3-h54a6638_2.conda - conda: https://conda.anaconda.org/conda-forge/linux-64/keyutils-1.6.3-hb9d3cd8_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/krb5-1.21.3-h659f571_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/ld_impl_linux-64-2.45.1-default_hbd61a6d_101.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libblas-3.11.0-5_h4a7cf45_openblas.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libcblas-3.11.0-5_h0358290_openblas.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/krb5-1.22.2-hbde042b_1.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/ld_impl_linux-64-2.46.1-default_hbd61a6d_102.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libblas-3.11.0-8_h4a7cf45_openblas.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libcblas-3.11.0-8_h0358290_openblas.conda - conda: https://conda.anaconda.org/conda-forge/linux-64/libedit-3.1.20250104-pl5321h7949ede_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libexpat-2.7.3-hecca717_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libexpat-2.8.1-hecca717_1.conda - conda: https://conda.anaconda.org/conda-forge/linux-64/libffi-3.5.2-h3435931_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libgcc-15.2.0-he0feb66_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libgcc-ng-15.2.0-h69a702a_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran-15.2.0-h69a702a_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran5-15.2.0-h68bc16d_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libgomp-15.2.0-he0feb66_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/liblapack-3.11.0-5_h47877c9_openblas.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/liblzma-5.8.2-hb03c661_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libgcc-15.2.0-he0feb66_20.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran-15.2.0-h69a702a_20.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran5-15.2.0-h68bc16d_20.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libgomp-15.2.0-he0feb66_20.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/liblapack-3.11.0-8_h47877c9_openblas.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/liblzma-5.8.3-hb03c661_0.conda - conda: https://conda.anaconda.org/conda-forge/linux-64/libmpdec-4.0.0-hb03c661_1.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libopenblas-0.3.30-pthreads_h94d23a6_4.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libsodium-1.0.20-h4ab18f5_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libsqlite-3.51.2-hf4e2dac_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libstdcxx-15.2.0-h934c35e_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libstdcxx-ng-15.2.0-hdf11a46_17.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libuuid-2.41.3-h5347b49_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/libzlib-1.3.1-hb9d3cd8_2.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/markdown-it-py-4.0.0-pyhd8ed1ab_0.conda - - conda: https://conda.modular.com/max-nightly/linux-64/max-26.2.0.dev2026022205-3.14release.conda - - conda: https://conda.modular.com/max-nightly/linux-64/max-core-26.2.0.dev2026022205-release.conda - - conda: https://conda.modular.com/max-nightly/noarch/mblack-26.2.0.dev2026022205-release.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libopenblas-0.3.33-pthreads_h94d23a6_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libsodium-1.0.22-h280c20c_1.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libsqlite-3.53.4-hf4e2dac_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libstdcxx-15.2.0-h934c35e_20.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libuuid-2.42.2-h5347b49_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/libzlib-1.3.2-h25fd6f3_2.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/msgspec-0.21.1-py314h5bd0f2a_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/ncurses-6.6-hdb14827_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/numpy-2.5.1-py314h2b28147_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/openssl-3.6.3-h35e630c_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/psutil-7.2.2-py314h0f05182_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/python-3.14.6-habeac84_101_cp314.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/pyzmq-27.1.0-py312hda471dd_3.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/readline-8.3-h853b02a_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/tk-8.6.13-noxft_hd70dff1_3.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/tornado-6.5.7-py314h5bd0f2a_0.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/zeromq-4.3.5-h09e67af_11.conda + - conda: https://conda.anaconda.org/conda-forge/linux-64/zstd-1.5.7-hb78ec9c_6.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/_python_abi3_support-1.0-hd8ed1ab_3.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/ca-certificates-2026.7.22-hbd8a1cb_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/click-8.4.2-pyhc90fa1f_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/cpython-3.14.6-py314hd8ed1ab_101.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/exceptiongroup-1.3.1-pyhd8ed1ab_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/importlib-metadata-9.0.0-pyhcf101f3_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_client-8.6.3-pyhd8ed1ab_1.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_core-5.9.1-pyhc90fa1f_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/markdown-it-py-4.2.0-pyhd8ed1ab_0.conda - conda: https://conda.anaconda.org/conda-forge/noarch/mdurl-0.1.2-pyhd8ed1ab_1.conda - - conda: https://conda.modular.com/max-nightly/linux-64/mojo-0.26.2.0.dev2026022205-release.conda - - conda: https://conda.modular.com/max-nightly/linux-64/mojo-compiler-0.26.2.0.dev2026022205-release.conda - - conda: https://conda.modular.com/max-nightly/noarch/mojo-python-0.26.2.0.dev2026022205-release.conda - conda: https://conda.anaconda.org/conda-forge/noarch/mypy_extensions-1.1.0-pyha770c72_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/ncurses-6.5-h2d0b736_3.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/numpy-2.4.2-py314h2b28147_1.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/openssl-3.6.1-h35e630c_1.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/packaging-26.0-pyhcf101f3_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/pathspec-1.0.4-pyhd8ed1ab_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/platformdirs-4.7.0-pyhcf101f3_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/pygments-2.19.2-pyhd8ed1ab_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/python-3.14.3-h32b2ec7_101_cp314.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/packaging-26.2-pyhc364b38_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/pathspec-1.1.1-pyhd8ed1ab_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/platformdirs-4.11.0-pyhcf101f3_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/pygments-2.20.0-pyhd8ed1ab_0.conda - conda: https://conda.anaconda.org/conda-forge/noarch/python-dateutil-2.9.0.post0-pyhe01879c_2.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/python-gil-3.14.3-h4df99d1_101.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/python-gil-3.14.6-h4df99d1_101.conda - conda: https://conda.anaconda.org/conda-forge/noarch/python_abi-3.14-8_cp314.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/pyyaml-6.0.3-py314h67df5f8_1.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/pyzmq-27.1.0-py312hfb55c3c_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/readline-8.3-h853b02a_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/rich-14.3.2-pyhcf101f3_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/rich-15.0.0-pyhcf101f3_0.conda - conda: https://conda.anaconda.org/conda-forge/noarch/six-1.17.0-pyhe01879c_1.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/tk-8.6.13-noxft_h366c992_103.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/tomli-2.4.0-pyhcf101f3_0.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/tornado-6.5.3-py314h5bd0f2a_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/traitlets-5.14.3-pyhd8ed1ab_1.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/typing-extensions-4.15.0-h396c80c_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/typing_extensions-4.15.0-pyhcf101f3_0.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/tzdata-2025c-hc9c84f9_1.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/yaml-0.2.5-h280c20c_3.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/zeromq-4.3.5-h387f397_9.conda - - conda: https://conda.anaconda.org/conda-forge/noarch/zipp-3.23.0-pyhcf101f3_1.conda - - conda: https://conda.anaconda.org/conda-forge/linux-64/zstd-1.5.7-hb78ec9c_6.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/taskgroup-0.2.2-pyhd8ed1ab_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/tomli-2.4.1-pyhcf101f3_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/traitlets-5.15.1-pyhcf101f3_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/typing-extensions-4.16.0-h69aa097_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/typing_extensions-4.16.0-pyhcf101f3_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/tzdata-2026c-h151e31d_0.conda + - conda: https://conda.anaconda.org/conda-forge/noarch/zipp-4.1.0-pyhcf101f3_0.conda + - conda: https://conda.modular.com/max-nightly/linux-64/max-26.5.0.dev2026072505-3.14release.conda + - conda: https://conda.modular.com/max-nightly/linux-64/max-core-26.5.0.dev2026072505-release.conda + - conda: https://conda.modular.com/max-nightly/linux-64/mojo-1.0.0b3.dev2026072505-release.conda + - conda: https://conda.modular.com/max-nightly/linux-64/mojo-compiler-1.0.0b3.dev2026072505-release.conda + - conda: https://conda.modular.com/max-nightly/noarch/mblack-26.5.0.dev2026072505-release.conda + - conda: https://conda.modular.com/max-nightly/noarch/mojo-python-1.0.0b3.dev2026072505-release.conda packages: - conda: https://conda.anaconda.org/conda-forge/linux-64/_openmp_mutex-4.5-20_gnu.conda build_number: 20 @@ -88,110 +93,39 @@ packages: constrains: - openmp_impl <0.0a0 license: BSD-3-Clause + license_family: BSD + run_exports: + strong: + - _openmp_mutex >=4.5 size: 28948 timestamp: 1770939786096 -- conda: https://conda.anaconda.org/conda-forge/noarch/_python_abi3_support-1.0-hd8ed1ab_2.conda - sha256: a3967b937b9abf0f2a99f3173fa4630293979bd1644709d89580e7c62a544661 - md5: aaa2a381ccc56eac91d63b6c1240312f - depends: - - cpython - - python-gil - license: MIT - license_family: MIT - size: 8191 - timestamp: 1744137672556 -- conda: https://conda.anaconda.org/conda-forge/linux-64/bzip2-1.0.8-hda65f42_8.conda - sha256: c30daba32ddebbb7ded490f0e371eae90f51e72db620554089103b4a6934b0d5 - md5: 51a19bba1b8ebfb60df25cde030b7ebc +- conda: https://conda.anaconda.org/conda-forge/linux-64/bzip2-1.0.8-hda65f42_9.conda + sha256: 0b75d45f0bba3e95dc693336fa51f40ea28c980131fec438afb7ce6118ed05f6 + md5: d2ffd7602c02f2b316fd921d39876885 depends: - __glibc >=2.17,<3.0.a0 - libgcc >=14 license: bzip2-1.0.6 license_family: BSD - size: 260341 - timestamp: 1757437258798 -- conda: https://conda.anaconda.org/conda-forge/noarch/ca-certificates-2026.1.4-hbd8a1cb_0.conda - sha256: b5974ec9b50e3c514a382335efa81ed02b05906849827a34061c496f4defa0b2 - md5: bddacf101bb4dd0e51811cb69c7790e2 + run_exports: + weak: + - bzip2 >=1.0.8,<2.0a0 + size: 260182 + timestamp: 1771350215188 +- conda: https://conda.anaconda.org/conda-forge/linux-64/icu-78.3-h54a6638_2.conda + sha256: d7c260b7e1cf22ce04d6ba8a86eabf4e6c50bc96a5c27fe2ecb32298af3e88eb + md5: 4ef4b977bb216a3001a3334696a80850 depends: - - __unix - license: ISC - size: 146519 - timestamp: 1767500828366 -- conda: https://conda.anaconda.org/conda-forge/noarch/click-8.3.1-pyh8f84b5b_1.conda - sha256: 38cfe1ee75b21a8361c8824f5544c3866f303af1762693a178266d7f198e8715 - md5: ea8a6c3256897cc31263de9f455e25d9 - depends: - - python >=3.10 - - __unix - - python - license: BSD-3-Clause - license_family: BSD - size: 97676 - timestamp: 1764518652276 -- conda: https://conda.anaconda.org/conda-forge/noarch/cpython-3.14.3-py314hd8ed1ab_101.conda - noarch: generic - sha256: 91b06300879df746214f7363d6c27c2489c80732e46a369eb2afc234bcafb44c - md5: 3bb89e4f795e5414addaa531d6b1500a - depends: - - python >=3.14,<3.15.0a0 - - python_abi * *_cp314 - license: Python-2.0 - size: 50078 - timestamp: 1770674447292 -- conda: https://conda.anaconda.org/conda-forge/linux-64/icu-78.2-h33c6efd_0.conda - sha256: 142a722072fa96cf16ff98eaaf641f54ab84744af81754c292cb81e0881c0329 - md5: 186a18e3ba246eccfc7cff00cd19a870 - depends: - - __glibc >=2.17,<3.0.a0 - - libgcc >=14 - libstdcxx >=14 + - libgcc >=14 + - __glibc >=2.17,<3.0.a0 license: MIT license_family: MIT - size: 12728445 - timestamp: 1767969922681 -- conda: https://conda.anaconda.org/conda-forge/noarch/importlib-metadata-8.7.0-pyhe01879c_1.conda - sha256: c18ab120a0613ada4391b15981d86ff777b5690ca461ea7e9e49531e8f374745 - md5: 63ccfdc3a3ce25b027b8767eb722fca8 - depends: - - python >=3.9 - - zipp >=3.20 - - python - license: Apache-2.0 - license_family: APACHE - size: 34641 - timestamp: 1747934053147 -- conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_client-8.6.3-pyhd8ed1ab_1.conda - sha256: 19d8bd5bb2fde910ec59e081eeb59529491995ce0d653a5209366611023a0b3a - md5: 4ebae00eae9705b0c3d6d1018a81d047 - depends: - - importlib-metadata >=4.8.3 - - jupyter_core >=4.12,!=5.0.* - - python >=3.9 - - python-dateutil >=2.8.2 - - pyzmq >=23.0 - - tornado >=6.2 - - traitlets >=5.3 - license: BSD-3-Clause - license_family: BSD - size: 106342 - timestamp: 1733441040958 -- conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_core-5.9.1-pyhc90fa1f_0.conda - sha256: 1d34b80e5bfcd5323f104dbf99a2aafc0e5d823019d626d0dce5d3d356a2a52a - md5: b38fe4e78ee75def7e599843ef4c1ab0 - depends: - - __unix - - python - - platformdirs >=2.5 - - python >=3.10 - - traitlets >=5.3 - - python - constrains: - - pywin32 >=300 - license: BSD-3-Clause - license_family: BSD - size: 65503 - timestamp: 1760643864586 + run_exports: + weak: + - icu >=78.3,<79.0a0 + size: 14455340 + timestamp: 1784916378180 - conda: https://conda.anaconda.org/conda-forge/linux-64/keyutils-1.6.3-hb9d3cd8_0.conda sha256: 0960d06048a7185d3542d850986d807c6e37ca2e644342dd0c72feefcf26c2a4 md5: b38117a3c920364aff79f870c984b4a3 @@ -199,65 +133,79 @@ packages: - __glibc >=2.17,<3.0.a0 - libgcc >=13 license: LGPL-2.1-or-later + run_exports: + weak: + - keyutils >=1.6.3,<2.0a0 size: 134088 timestamp: 1754905959823 -- conda: https://conda.anaconda.org/conda-forge/linux-64/krb5-1.21.3-h659f571_0.conda - sha256: 99df692f7a8a5c27cd14b5fb1374ee55e756631b9c3d659ed3ee60830249b238 - md5: 3f43953b7d3fb3aaa1d0d0723d91e368 - depends: - - keyutils >=1.6.1,<2.0a0 - - libedit >=3.1.20191231,<3.2.0a0 - - libedit >=3.1.20191231,<4.0a0 - - libgcc-ng >=12 - - libstdcxx-ng >=12 - - openssl >=3.3.1,<4.0a0 +- conda: https://conda.anaconda.org/conda-forge/linux-64/krb5-1.22.2-hbde042b_1.conda + sha256: 9b07046870772f28740e3f6149f09ff222843733087a33c5540b169c6289652d + md5: 54157a1c8c0bb70f62dd0b17fba7e7f2 + depends: + - __glibc >=2.17,<3.0.a0 + - keyutils >=1.6.3,<2.0a0 + - libedit >=3.1.20250104,<3.2.0a0 + - libedit >=3.1.20250104,<4.0a0 + - libgcc >=14 + - libstdcxx >=14 + - openssl >=3.5.7,<4.0a0 license: MIT license_family: MIT - size: 1370023 - timestamp: 1719463201255 -- conda: https://conda.anaconda.org/conda-forge/linux-64/ld_impl_linux-64-2.45.1-default_hbd61a6d_101.conda - sha256: 565941ac1f8b0d2f2e8f02827cbca648f4d18cd461afc31f15604cd291b5c5f3 - md5: 12bd9a3f089ee6c9266a37dab82afabd + run_exports: + weak: + - krb5 >=1.22.2,<1.23.0a0 + size: 1388990 + timestamp: 1781859420533 +- conda: https://conda.anaconda.org/conda-forge/linux-64/ld_impl_linux-64-2.46.1-default_hbd61a6d_102.conda + sha256: 27d83f1188cd19bcb7754a078b3fa7f4cfb8527f8eb2fde54dd01fc529d1adec + md5: 449500f2c089da11c40f5c21312e3e07 depends: - __glibc >=2.17,<3.0.a0 - zstd >=1.5.7,<1.6.0a0 constrains: - - binutils_impl_linux-64 2.45.1 + - binutils_impl_linux-64 2.46.1 license: GPL-3.0-only license_family: GPL - size: 725507 - timestamp: 1770267139900 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libblas-3.11.0-5_h4a7cf45_openblas.conda - build_number: 5 - sha256: 18c72545080b86739352482ba14ba2c4815e19e26a7417ca21a95b76ec8da24c - md5: c160954f7418d7b6e87eaf05a8913fa9 - depends: - - libopenblas >=0.3.30,<0.3.31.0a0 - - libopenblas >=0.3.30,<1.0a0 + run_exports: {} + size: 745303 + timestamp: 1784214507189 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libblas-3.11.0-8_h4a7cf45_openblas.conda + build_number: 8 + sha256: b2da6bfd72a1c9cb143ccf64bf5b28790cb4eb58bd1cb978f6537b2322f7d48b + md5: 00fc660ab1b2f5ca07e92b4900d10c79 + depends: + - libopenblas >=0.3.33,<0.3.34.0a0 + - libopenblas >=0.3.33,<1.0a0 constrains: - - mkl <2026 - - liblapack 3.11.0 5*_openblas - - libcblas 3.11.0 5*_openblas - - blas 2.305 openblas - - liblapacke 3.11.0 5*_openblas + - blas 2.308 openblas + - mkl <2027 + - libcblas 3.11.0 8*_openblas + - liblapack 3.11.0 8*_openblas + - liblapacke 3.11.0 8*_openblas license: BSD-3-Clause license_family: BSD - size: 18213 - timestamp: 1765818813880 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libcblas-3.11.0-5_h0358290_openblas.conda - build_number: 5 - sha256: 0cbdcc67901e02dc17f1d19e1f9170610bd828100dc207de4d5b6b8ad1ae7ad8 - md5: 6636a2b6f1a87572df2970d3ebc87cc0 - depends: - - libblas 3.11.0 5_h4a7cf45_openblas + run_exports: + weak: + - libblas >=3.11.0,<4.0a0 + size: 18804 + timestamp: 1779859100675 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libcblas-3.11.0-8_h0358290_openblas.conda + build_number: 8 + sha256: 1a2bc77bb26520255904a3d9b1f40e6bf0bf9d8d3405c7709dd162282820915a + md5: 33a413f1095f8325e5c30fde3b0d2445 + depends: + - libblas 3.11.0 8_h4a7cf45_openblas constrains: - - liblapacke 3.11.0 5*_openblas - - blas 2.305 openblas - - liblapack 3.11.0 5*_openblas + - blas 2.308 openblas + - liblapacke 3.11.0 8*_openblas + - liblapack 3.11.0 8*_openblas license: BSD-3-Clause license_family: BSD - size: 18194 - timestamp: 1765818837135 + run_exports: + weak: + - libcblas >=3.11.0,<4.0a0 + size: 18778 + timestamp: 1779859107964 - conda: https://conda.anaconda.org/conda-forge/linux-64/libedit-3.1.20250104-pl5321h7949ede_0.conda sha256: d789471216e7aba3c184cd054ed61ce3f6dac6f87a50ec69291b9297f8c18724 md5: c277e0a4d549b03ac1e9d6cbbe3d017b @@ -268,20 +216,24 @@ packages: - ncurses >=6.5,<7.0a0 license: BSD-2-Clause license_family: BSD + run_exports: + weak: + - libedit >=3.1.20250104,<3.2.0a0 size: 134676 timestamp: 1738479519902 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libexpat-2.7.3-hecca717_0.conda - sha256: 1e1b08f6211629cbc2efe7a5bca5953f8f6b3cae0eeb04ca4dacee1bd4e2db2f - md5: 8b09ae86839581147ef2e5c5e229d164 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libexpat-2.8.1-hecca717_1.conda + sha256: 16feffd9ddbbe5b718515d38ee376c685ba95491cd901244e24671d20b952a77 + md5: b24d3c612f71e7aa74158d92106318b2 depends: - __glibc >=2.17,<3.0.a0 - libgcc >=14 constrains: - - expat 2.7.3.* + - expat 2.8.1.* license: MIT license_family: MIT - size: 76643 - timestamp: 1763549731408 + run_exports: {} + size: 77856 + timestamp: 1781203599810 - conda: https://conda.anaconda.org/conda-forge/linux-64/libffi-3.5.2-h3435931_0.conda sha256: 31f19b6a88ce40ebc0d5a992c131f57d919f73c0b92cd1617a5bec83f6e961e6 md5: a360c33a5abe61c07959e449fa1453eb @@ -290,44 +242,40 @@ packages: - libgcc >=14 license: MIT license_family: MIT + run_exports: + weak: + - libffi >=3.5.2,<3.6.0a0 size: 58592 timestamp: 1769456073053 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libgcc-15.2.0-he0feb66_17.conda - sha256: 43860222cf3abf04ded0cf24541a105aa388e0e1d4d6ca46258e186d4e87ae3e - md5: 3c281169ea25b987311400d7a7e28445 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libgcc-15.2.0-he0feb66_20.conda + sha256: e3b7bb287d1685b15781a6d9e3408d6ddaff23f5a1d0d6c0140619dd3950fe72 + md5: 3533de187cf7283f96bfdb28ad73e2bc depends: - __glibc >=2.17,<3.0.a0 - _openmp_mutex >=4.5 constrains: - - libgcc-ng ==15.2.0=*_17 - - libgomp 15.2.0 he0feb66_17 - license: GPL-3.0-only WITH GCC-exception-3.1 - license_family: GPL - size: 1040478 - timestamp: 1770252533873 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libgcc-ng-15.2.0-h69a702a_17.conda - sha256: bdfe50501e4a2d904a5eae65a7ae26e2b7a29b473ab084ad55d96080b966502e - md5: 1478bfa85224a65ab096d69ffd2af1e5 - depends: - - libgcc 15.2.0 he0feb66_17 + - libgomp 15.2.0 he0feb66_20 + - libgcc-ng ==15.2.0=*_20 license: GPL-3.0-only WITH GCC-exception-3.1 license_family: GPL - size: 27541 - timestamp: 1770252546553 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran-15.2.0-h69a702a_17.conda - sha256: 1604c083dd65bc91e68b6cfe32c8610395088cb96af1acaf71f0dcaf83ac58f7 - md5: a6c682ac611cb1fa4d73478f9e6efb06 - depends: - - libgfortran5 15.2.0 h68bc16d_17 + run_exports: {} + size: 1041122 + timestamp: 1784923795784 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran-15.2.0-h69a702a_20.conda + sha256: 76d81032e264b74f519cc26962d5eb39547e0785fc9d2957bbc12e7a91214412 + md5: a450a08a63f940e9aa7b37692e71196a + depends: + - libgfortran5 15.2.0 h68bc16d_20 constrains: - - libgfortran-ng ==15.2.0=*_17 + - libgfortran-ng ==15.2.0=*_20 license: GPL-3.0-only WITH GCC-exception-3.1 license_family: GPL - size: 27515 - timestamp: 1770252591906 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran5-15.2.0-h68bc16d_17.conda - sha256: b1c77b85da9a3e204de986f59e262268805c6a35dffdf3953f1b98407db2aef3 - md5: 202fdf8cad9eea704c2b0d823d1732bf + run_exports: {} + size: 28103 + timestamp: 1784923831860 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libgfortran5-15.2.0-h68bc16d_20.conda + sha256: 95122f42566dc9a62b9615d2372b3f3c75fa7bd8bb1eda53aa7532ce60d545eb + md5: 4edbcbea1a8790a7d58e648523b69546 depends: - __glibc >=2.17,<3.0.a0 - libgcc >=15.2.0 @@ -335,42 +283,52 @@ packages: - libgfortran 15.2.0 license: GPL-3.0-only WITH GCC-exception-3.1 license_family: GPL - size: 2480824 - timestamp: 1770252563579 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libgomp-15.2.0-he0feb66_17.conda - sha256: b961b5dd9761907a7179678b58a69bb4fc16b940eb477f635aea3aec0a3f17a6 - md5: 51b78c6a757575c0d12f4401ffc67029 + run_exports: {} + size: 2483727 + timestamp: 1784923810904 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libgomp-15.2.0-he0feb66_20.conda + sha256: 30b8ff1bf43871a17c9de99e56171ef54cfbe690ffa86681628db5a00af97cf7 + md5: 49321086c41bb58fc4b6cd8cbb74679d depends: - __glibc >=2.17,<3.0.a0 license: GPL-3.0-only WITH GCC-exception-3.1 license_family: GPL - size: 603334 - timestamp: 1770252441199 -- conda: https://conda.anaconda.org/conda-forge/linux-64/liblapack-3.11.0-5_h47877c9_openblas.conda - build_number: 5 - sha256: c723b6599fcd4c6c75dee728359ef418307280fa3e2ee376e14e85e5bbdda053 - md5: b38076eb5c8e40d0106beda6f95d7609 - depends: - - libblas 3.11.0 5_h4a7cf45_openblas + run_exports: + strong: + - _openmp_mutex >=4.5 + size: 603998 + timestamp: 1784923730278 +- conda: https://conda.anaconda.org/conda-forge/linux-64/liblapack-3.11.0-8_h47877c9_openblas.conda + build_number: 8 + sha256: 168e327d737059553e15cc6ec36d76b9bbb3931c2a7721555fd68b4c9348b247 + md5: 809be8ba8712c77bc7d44c2d99390dc4 + depends: + - libblas 3.11.0 8_h4a7cf45_openblas constrains: - - blas 2.305 openblas - - liblapacke 3.11.0 5*_openblas - - libcblas 3.11.0 5*_openblas + - blas 2.308 openblas + - libcblas 3.11.0 8*_openblas + - liblapacke 3.11.0 8*_openblas license: BSD-3-Clause license_family: BSD - size: 18200 - timestamp: 1765818857876 -- conda: https://conda.anaconda.org/conda-forge/linux-64/liblzma-5.8.2-hb03c661_0.conda - sha256: 755c55ebab181d678c12e49cced893598f2bab22d582fbbf4d8b83c18be207eb - md5: c7c83eecbb72d88b940c249af56c8b17 + run_exports: + weak: + - liblapack >=3.11.0,<3.12.0a0 + size: 18790 + timestamp: 1779859115086 +- conda: https://conda.anaconda.org/conda-forge/linux-64/liblzma-5.8.3-hb03c661_0.conda + sha256: ec30e52a3c1bf7d0425380a189d209a52baa03f22fb66dd3eb587acaa765bd6d + md5: b88d90cad08e6bc8ad540cb310a761fb depends: - __glibc >=2.17,<3.0.a0 - libgcc >=14 constrains: - - xz 5.8.2.* + - xz 5.8.3.* license: 0BSD - size: 113207 - timestamp: 1768752626120 + run_exports: + weak: + - liblzma >=5.8.3,<6.0a0 + size: 113478 + timestamp: 1775825492909 - conda: https://conda.anaconda.org/conda-forge/linux-64/libmpdec-4.0.0-hb03c661_1.conda sha256: fe171ed5cf5959993d43ff72de7596e8ac2853e9021dec0344e583734f1e0843 md5: 2c21e66f50753a083cbe6b80f38268fa @@ -379,170 +337,395 @@ packages: - libgcc >=14 license: BSD-2-Clause license_family: BSD + run_exports: {} size: 92400 timestamp: 1769482286018 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libopenblas-0.3.30-pthreads_h94d23a6_4.conda - sha256: 199d79c237afb0d4780ccd2fbf829cea80743df60df4705202558675e07dd2c5 - md5: be43915efc66345cccb3c310b6ed0374 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libopenblas-0.3.33-pthreads_h94d23a6_0.conda + sha256: 3d9aa85648e5e18a6d66db98b8c4317cc426721ad7a220aa86330d1ccedc8903 + md5: 2d3278b721e40468295ca755c3b84070 depends: - __glibc >=2.17,<3.0.a0 - libgcc >=14 - libgfortran - libgfortran5 >=14.3.0 constrains: - - openblas >=0.3.30,<0.3.31.0a0 + - openblas >=0.3.33,<0.3.34.0a0 license: BSD-3-Clause license_family: BSD - size: 5927939 - timestamp: 1763114673331 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libsodium-1.0.20-h4ab18f5_0.conda - sha256: 0105bd108f19ea8e6a78d2d994a6d4a8db16d19a41212070d2d1d48a63c34161 - md5: a587892d3c13b6621a6091be690dbca2 + run_exports: + weak: + - libopenblas >=0.3.33,<1.0a0 + size: 5931919 + timestamp: 1776993658641 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libsodium-1.0.22-h280c20c_1.conda + sha256: b677bbf1c339d894757c3dcfbb2f88649e499e4991d70ae09a1466da9a6c92d6 + md5: 965e4d531b588b2e42f66fd8e48b056c depends: - - libgcc-ng >=12 + - __glibc >=2.17,<3.0.a0 + - libgcc >=14 license: ISC - size: 205978 - timestamp: 1716828628198 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libsqlite-3.51.2-hf4e2dac_0.conda - sha256: 04596fcee262a870e4b7c9807224680ff48d4d0cc0dac076a602503d3dc6d217 - md5: da5be73701eecd0e8454423fd6ffcf30 + run_exports: + weak: + - libsodium >=1.0.22,<1.0.23.0a0 + size: 269272 + timestamp: 1779163468406 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libsqlite-3.53.4-hf4e2dac_0.conda + sha256: 72023efc207fe681e26b65fc9d668062cf0b4f0eacf3431e6eb099b95c1f2efd + md5: df088a279cd5e6fd2790b4c196434da1 depends: - __glibc >=2.17,<3.0.a0 - - icu >=78.2,<79.0a0 + - icu >=78.3,<79.0a0 - libgcc >=14 - - libzlib >=1.3.1,<2.0a0 + - libzlib >=1.3.2,<2.0a0 license: blessing - size: 942808 - timestamp: 1768147973361 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libstdcxx-15.2.0-h934c35e_17.conda - sha256: 50c48cd3716a2e58e8e2e02edc78fef2d08fffe1e3b1ed40eb5f87e7e2d07889 - md5: 24c2fe35fa45cd71214beba6f337c071 + run_exports: + weak: + - libsqlite >=3.53.4,<4.0a0 + size: 964200 + timestamp: 1785016112246 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libstdcxx-15.2.0-h934c35e_20.conda + sha256: b5eadda8fb0df262f0d424c4a0c8c1c8d65d904ce622f07936c793ea1b8a39d9 + md5: fbd3d5506b11b5cfc916b29263b6b6f7 depends: - __glibc >=2.17,<3.0.a0 - - libgcc 15.2.0 he0feb66_17 + - libgcc 15.2.0 he0feb66_20 constrains: - - libstdcxx-ng ==15.2.0=*_17 + - libstdcxx-ng ==15.2.0=*_20 license: GPL-3.0-only WITH GCC-exception-3.1 license_family: GPL - size: 5852406 - timestamp: 1770252584235 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libstdcxx-ng-15.2.0-hdf11a46_17.conda - sha256: ca3fb322dab3373946b1064da686ec076f5b1b9caf0a2823dad00d0b0f704928 - md5: ea12f5a6bf12c88c06750d9803e1a570 + run_exports: {} + size: 5857690 + timestamp: 1784923825011 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libuuid-2.42.2-h5347b49_0.conda + sha256: 9b1bdce27a7e31f7d241aeecff67a1f3101d52a2b1e33ccc2cdf2613072bf81f + md5: 01bb81d12c957de066ea7362007df642 depends: - - libstdcxx 15.2.0 h934c35e_17 - license: GPL-3.0-only WITH GCC-exception-3.1 - license_family: GPL - size: 27573 - timestamp: 1770252638797 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libuuid-2.41.3-h5347b49_0.conda - sha256: 1a7539cfa7df00714e8943e18de0b06cceef6778e420a5ee3a2a145773758aee - md5: db409b7c1720428638e7c0d509d3e1b5 - depends: - - __glibc >=2.17,<3.0.a0 - libgcc >=14 + - __glibc >=2.17,<3.0.a0 license: BSD-3-Clause license_family: BSD - size: 40311 - timestamp: 1766271528534 -- conda: https://conda.anaconda.org/conda-forge/linux-64/libzlib-1.3.1-hb9d3cd8_2.conda - sha256: d4bfe88d7cb447768e31650f06257995601f89076080e76df55e3112d4e47dc4 - md5: edb0dca6bc32e4f4789199455a1dbeb8 + run_exports: + weak: + - libuuid >=2.42.2,<3.0a0 + size: 40017 + timestamp: 1781625522462 +- conda: https://conda.anaconda.org/conda-forge/linux-64/libzlib-1.3.2-h25fd6f3_2.conda + sha256: 55044c403570f0dc26e6364de4dc5368e5f3fc7ff103e867c487e2b5ab2bcda9 + md5: d87ff7921124eccd67248aa483c23fec depends: - __glibc >=2.17,<3.0.a0 - - libgcc >=13 constrains: - - zlib 1.3.1 *_2 + - zlib 1.3.2 *_2 license: Zlib license_family: Other - size: 60963 - timestamp: 1727963148474 -- conda: https://conda.anaconda.org/conda-forge/noarch/markdown-it-py-4.0.0-pyhd8ed1ab_0.conda - sha256: 7b1da4b5c40385791dbc3cc85ceea9fad5da680a27d5d3cb8bfaa185e304a89e - md5: 5b5203189eb668f042ac2b0826244964 + run_exports: + weak: + - libzlib >=1.3.2,<2.0a0 + size: 63629 + timestamp: 1774072609062 +- conda: https://conda.anaconda.org/conda-forge/linux-64/msgspec-0.21.1-py314h5bd0f2a_0.conda + sha256: 52565ceea81e801c59dcaeaf5a9c77fba2fade445e67e0864fda50d4b944e15b + md5: 4a8ea416a56e58f012e445f7af2bbcc8 depends: - - mdurl >=0.1,<1 - - python >=3.10 - license: MIT - license_family: MIT - size: 64736 - timestamp: 1754951288511 -- conda: https://conda.modular.com/max-nightly/linux-64/max-26.2.0.dev2026022205-3.14release.conda - sha256: abe9988bd7aa6881b07157512251af5b5e0b930e319030e7923effd6c7bbe8ce - md5: 724d343713de47a6210ba3c8127c0d3b + - __glibc >=2.17,<3.0.a0 + - libgcc >=14 + - python >=3.14,<3.15.0a0 + - python_abi 3.14.* *_cp314 + license: BSD-3-Clause + license_family: BSD + run_exports: {} + size: 220990 + timestamp: 1776337508167 +- conda: https://conda.anaconda.org/conda-forge/linux-64/ncurses-6.6-hdb14827_0.conda + sha256: fc89f74bbe362fb29fa3c037697a89bec140b346a2469a90f7936d1d7ea4d8a3 + md5: fc21868a1a5aacc937e7a18747acb8a5 depends: - - numpy >=1.18 - - typing-extensions >=4.12.2 - - pyyaml >=6.0.1 - - rich >=13.0.1 - - python-gil - - max-core ==26.2.0.dev2026022205 release + - __glibc >=2.17,<3.0.a0 + - libgcc >=14 + license: X11 AND BSD-3-Clause + run_exports: + weak: + - ncurses >=6.6,<7.0a0 + size: 918956 + timestamp: 1777422145199 +- conda: https://conda.anaconda.org/conda-forge/linux-64/numpy-2.5.1-py314h2b28147_0.conda + sha256: fb665b7e30f9472c58098983f614598cfcf34e7a26b6c419648803095c390aa7 + md5: 59044905d27ba41bfb280ed4692c15e2 + depends: + - python + - libgcc >=14 + - __glibc >=2.17,<3.0.a0 + - libstdcxx >=14 + - libblas >=3.9.0,<4.0a0 + - liblapack >=3.9.0,<4.0a0 - python_abi 3.14.* *_cp314 + - libcblas >=3.9.0,<4.0a0 constrains: - - click >=8.0.0 - - exceptiongroup >=0.2.2 - - gguf >=0.17.1 - - hf-transfer >=0.1.9 - - huggingface_hub >=0.28.0 - - llguidance >=0.7.30 - - jinja2 >=3.1.0 - - pillow >=11.0.0 - - psutil >=6.1.1 - - pydantic-settings >=2.7.1 - - pydantic - - pydantic-core - - requests >=2.32.3 - - sentencepiece >=0.2.0 - - taskgroup >=0.2.2 - - tqdm >=4.67.1 - - transformers >=4.57.0,<5.0.0 - - uvicorn >=0.34.0 - - uvloop >=0.21.0 - - aiofiles >=24.1.0 - - asgiref >=3.8.1 - - fastapi >=0.115.3 - - grpcio >=1.68.0 - - httpx >=0.28.1,<0.29 - - msgspec >=0.19.0 - - opentelemetry-api >=1.29.0 - - opentelemetry-exporter-otlp-proto-http >=1.27.0 - - opentelemetry-exporter-prometheus >=0.50b0 - - opentelemetry-sdk >=1.29.0,<1.36.0 - - prometheus_client >=0.21.0 - - protobuf >=6.33.5,<6.34.0 - - pyinstrument >=5.0.1 - - python-json-logger >=2.0.7 - - pyzmq >=26.3.0 - - regex >=2024.11.6 - - scipy >=1.13.0 - - sse-starlette >=2.1.2 - - starlette >=0.47.2 - license: LicenseRef-Modular-Proprietary - size: 6106368 - timestamp: 1771738763384 -- conda: https://conda.modular.com/max-nightly/linux-64/max-core-26.2.0.dev2026022205-release.conda - sha256: 898982c0b96cb40d83c08d84bf7a7295a58387f8c0be9571055f6fbbd5e9389b - md5: a25753c21792034b93c21cf09b47031e + - numpy-base <0a0 + license: BSD-3-Clause + license_family: BSD + run_exports: + weak: + - numpy >=1.25,<3 + size: 9089255 + timestamp: 1783206203765 +- conda: https://conda.anaconda.org/conda-forge/linux-64/openssl-3.6.3-h35e630c_0.conda + sha256: d48f5c22b9897c01e4dff3680f1f57ceb02711ab9c62f74339b080419dfad34b + md5: 79dd2074b5cd5c5c6b2930514a11e22d depends: - - mojo-compiler ==0.26.2.0.dev2026022205 release - license: LicenseRef-Modular-Proprietary - size: 131486199 - timestamp: 1771738763382 -- conda: https://conda.modular.com/max-nightly/noarch/mblack-26.2.0.dev2026022205-release.conda + - __glibc >=2.17,<3.0.a0 + - ca-certificates + - libgcc >=14 + license: Apache-2.0 + license_family: Apache + run_exports: + weak: + - openssl >=3.6.3,<4.0a0 + size: 3159683 + timestamp: 1781069855778 +- conda: https://conda.anaconda.org/conda-forge/linux-64/psutil-7.2.2-py314h0f05182_0.conda + sha256: f15574ed6c8c8ed8c15a0c5a00102b1efe8b867c0bd286b498cd98d95bd69ae5 + md5: 4f225a966cfee267a79c5cb6382bd121 + depends: + - python + - libgcc >=14 + - __glibc >=2.17,<3.0.a0 + - python_abi 3.14.* *_cp314 + license: BSD-3-Clause + license_family: BSD + run_exports: {} + size: 231303 + timestamp: 1769678156552 +- conda: https://conda.anaconda.org/conda-forge/linux-64/python-3.14.6-habeac84_101_cp314.conda + build_number: 101 + sha256: ee8f2006e1724b1f2e9e0ccc5a7cfdcab973460faa2f63ac1f6e44fdad4c0344 + md5: 78975a41cf3c525da654f17e35bfca9e + depends: + - __glibc >=2.17,<3.0.a0 + - bzip2 >=1.0.8,<2.0a0 + - ld_impl_linux-64 >=2.36.1 + - libexpat >=2.8.1,<3.0a0 + - libffi >=3.5.2,<3.6.0a0 + - libgcc >=14 + - liblzma >=5.8.3,<6.0a0 + - libmpdec >=4.0.0,<5.0a0 + - libsqlite >=3.53.3,<4.0a0 + - libuuid >=2.42.2,<3.0a0 + - libzlib >=1.3.2,<2.0a0 + - ncurses >=6.6,<7.0a0 + - openssl >=3.5.7,<4.0a0 + - python_abi 3.14.* *_cp314 + - readline >=8.3,<9.0a0 + - tk >=8.6.13,<8.7.0a0 + - tzdata + - zstd >=1.5.7,<1.6.0a0 + license: Python-2.0 + run_exports: + weak: + - python_abi 3.14.* *_cp314 + noarch: + - python + size: 36869055 + timestamp: 1784910110714 + python_site_packages_path: lib/python3.14/site-packages +- conda: https://conda.anaconda.org/conda-forge/linux-64/pyzmq-27.1.0-py312hda471dd_3.conda noarch: python - sha256: 1e044bdd68205536a6a8e2158240bb1e3fc2e6eee294fcd181a9957c0eb7ade6 - md5: aaa44d262d7410c49ab18353820f185f + sha256: 970b2a1d12983d8d1cc05d914ad88a0b6ef1fa14038c9649aa834dd6ebee65d7 + md5: acd216255e1370e9aeab5351b831f07c + depends: + - python + - libgcc >=14 + - libstdcxx >=14 + - __glibc >=2.17,<3.0.a0 + - _python_abi3_support 1.* + - cpython >=3.12 + - zeromq >=4.3.5,<4.4.0a0 + license: BSD-3-Clause + license_family: BSD + run_exports: {} + size: 210896 + timestamp: 1779483879367 +- conda: https://conda.anaconda.org/conda-forge/linux-64/readline-8.3-h853b02a_0.conda + sha256: 12ffde5a6f958e285aa22c191ca01bbd3d6e710aa852e00618fa6ddc59149002 + md5: d7d95fc8287ea7bf33e0e7116d2b95ec + depends: + - __glibc >=2.17,<3.0.a0 + - libgcc >=14 + - ncurses >=6.5,<7.0a0 + license: GPL-3.0-only + license_family: GPL + run_exports: + weak: + - readline >=8.3,<9.0a0 + size: 345073 + timestamp: 1765813471974 +- conda: https://conda.anaconda.org/conda-forge/linux-64/tk-8.6.13-noxft_hd70dff1_3.conda + build_number: 103 + sha256: 43624eab22f5f29df7d6ffe914cf442f28fd559b55b290906255492826e636e8 + md5: 48a1049e710857572fc2a832aa394d9f + depends: + - __glibc >=2.17,<3.0.a0 + - libgcc >=14 + - libzlib >=1.3.2,<2.0a0 + constrains: + - xorg-libx11 >=1.8.13,<2.0a0 + license: TCL + run_exports: + weak: + - tk >=8.6.13,<8.7.0a0 + size: 3550916 + timestamp: 1784229071544 +- conda: https://conda.anaconda.org/conda-forge/linux-64/tornado-6.5.7-py314h5bd0f2a_0.conda + sha256: bbb7056f7c5fd606df16ed73ee68687050de2c02fd69a3f69a1cb533a7ed2ae8 + md5: 4a8e5889712641aabdf6695e292857fe + depends: + - __glibc >=2.17,<3.0.a0 + - libgcc >=14 + - python >=3.14,<3.15.0a0 + - python_abi 3.14.* *_cp314 + license: Apache-2.0 + license_family: Apache + run_exports: {} + size: 918368 + timestamp: 1781006801436 +- conda: https://conda.anaconda.org/conda-forge/linux-64/zeromq-4.3.5-h09e67af_11.conda + sha256: dc9f28dedcb5f35a127fad2d847674d2833369dd616d294e423b8997df31d8a8 + md5: 96b08867e21d4694fa5c2c226e6581b0 + depends: + - libgcc >=14 + - __glibc >=2.17,<3.0.a0 + - libstdcxx >=14 + - krb5 >=1.22.2,<1.23.0a0 + - libsodium >=1.0.22,<1.0.23.0a0 + license: MPL-2.0 + license_family: MOZILLA + run_exports: + weak: + - zeromq >=4.3.5,<4.4.0a0 + size: 311184 + timestamp: 1779123989774 +- conda: https://conda.anaconda.org/conda-forge/linux-64/zstd-1.5.7-hb78ec9c_6.conda + sha256: 68f0206ca6e98fea941e5717cec780ed2873ffabc0e1ed34428c061e2c6268c7 + md5: 4a13eeac0b5c8e5b8ab496e6c4ddd829 + depends: + - __glibc >=2.17,<3.0.a0 + - libzlib >=1.3.1,<2.0a0 + license: BSD-3-Clause + license_family: BSD + run_exports: + weak: + - zstd >=1.5.7,<1.6.0a0 + size: 601375 + timestamp: 1764777111296 +- conda: https://conda.anaconda.org/conda-forge/noarch/_python_abi3_support-1.0-hd8ed1ab_3.conda + sha256: 2a7204314663eeda5dec482a956f0e2eaf289bd5b9953eaaaad0e81aa64638f2 + md5: 3845f3d75991bae0fb90884662f4327c + depends: + - cpython + - python-gil + license: MIT + license_family: MIT + run_exports: {} + size: 8144 + timestamp: 1784221492234 +- conda: https://conda.anaconda.org/conda-forge/noarch/ca-certificates-2026.7.22-hbd8a1cb_0.conda + sha256: 0a0544cf95f64394fe4959286f5c71f5444ad58feb0602e53becb27448d24da6 + md5: 0f51e2391ade309db462a55611263e9c + depends: + - __unix + license: ISC + run_exports: {} + size: 131780 + timestamp: 1784754889428 +- conda: https://conda.anaconda.org/conda-forge/noarch/click-8.4.2-pyhc90fa1f_0.conda + sha256: ccc4787f511964f9a1f2d2d2859c91c5d571fb60f7f09d4c4e092c9b7a94e671 + md5: 2c4bd6aeb90bb157456841c3270a0d92 depends: + - __unix + - python + - python >=3.10 + license: BSD-3-Clause + license_family: BSD + run_exports: {} + size: 107155 + timestamp: 1783085363526 +- conda: https://conda.anaconda.org/conda-forge/noarch/cpython-3.14.6-py314hd8ed1ab_101.conda + noarch: generic + sha256: 436618a5a090c9f7ade0c5f883e8602a130b3991bc88b06badd6f805d7dbff00 + md5: 424c465894c8af725105fe6ad74f6aec + depends: + - python >=3.14,<3.15.0a0 + - python_abi * *_cp314 + license: Python-2.0 + run_exports: {} + size: 49508 + timestamp: 1784909547134 +- conda: https://conda.anaconda.org/conda-forge/noarch/exceptiongroup-1.3.1-pyhd8ed1ab_0.conda + sha256: ee6cf346d017d954255bbcbdb424cddea4d14e4ed7e9813e429db1d795d01144 + md5: 8e662bd460bda79b1ea39194e3c4c9ab + depends: + - python >=3.10 + - typing_extensions >=4.6.0 + license: MIT and PSF-2.0 + run_exports: {} + size: 21333 + timestamp: 1763918099466 +- conda: https://conda.anaconda.org/conda-forge/noarch/importlib-metadata-9.0.0-pyhcf101f3_0.conda + sha256: 43e2a5497cad1598ff88a3e69f69bc88b7b8f141fa63c60eab5db296317318b8 + md5: ffc17e785d64e12fc311af9184221839 + depends: + - python >=3.10 + - zipp >=3.20 + - python + license: Apache-2.0 + license_family: APACHE + run_exports: {} + size: 34766 + timestamp: 1779714582554 +- conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_client-8.6.3-pyhd8ed1ab_1.conda + sha256: 19d8bd5bb2fde910ec59e081eeb59529491995ce0d653a5209366611023a0b3a + md5: 4ebae00eae9705b0c3d6d1018a81d047 + depends: + - importlib-metadata >=4.8.3 + - jupyter_core >=4.12,!=5.0.* + - python >=3.9 + - python-dateutil >=2.8.2 + - pyzmq >=23.0 + - tornado >=6.2 + - traitlets >=5.3 + license: BSD-3-Clause + license_family: BSD + run_exports: {} + size: 106342 + timestamp: 1733441040958 +- conda: https://conda.anaconda.org/conda-forge/noarch/jupyter_core-5.9.1-pyhc90fa1f_0.conda + sha256: 1d34b80e5bfcd5323f104dbf99a2aafc0e5d823019d626d0dce5d3d356a2a52a + md5: b38fe4e78ee75def7e599843ef4c1ab0 + depends: + - __unix + - python + - platformdirs >=2.5 + - python >=3.10 + - traitlets >=5.3 + - python + constrains: + - pywin32 >=300 + license: BSD-3-Clause + license_family: BSD + run_exports: {} + size: 65503 + timestamp: 1760643864586 +- conda: https://conda.anaconda.org/conda-forge/noarch/markdown-it-py-4.2.0-pyhd8ed1ab_0.conda + sha256: 0c4c35376fe920714390d46e4b8d31c876d65f18e1655899e0763ec25f2a902f + md5: 6d03368f2b2b0a5fb6839df53b2eb5e0 + depends: + - mdurl >=0.1,<1 - python >=3.10 - - click >=8.0.0 - - mypy_extensions >=0.4.3 - - packaging >=22.0 - - pathspec >=0.9.0 - - platformdirs >=2 - - tomli >=1.1.0 license: MIT - size: 135129 - timestamp: 1771738763383 + license_family: MIT + run_exports: {} + size: 69017 + timestamp: 1778169663339 - conda: https://conda.anaconda.org/conda-forge/noarch/mdurl-0.1.2-pyhd8ed1ab_1.conda sha256: 78c1bbe1723449c52b7a9df1af2ee5f005209f67e40b6e1d3c7619127c43b1c7 md5: 592132998493b3ff25fd7479396e8351 @@ -550,36 +733,9 @@ packages: - python >=3.9 license: MIT license_family: MIT + run_exports: {} size: 14465 timestamp: 1733255681319 -- conda: https://conda.modular.com/max-nightly/linux-64/mojo-0.26.2.0.dev2026022205-release.conda - sha256: 0d8ca62be9a970f4382c99549fb88c39c7dc28092ee812f03f1639990d4948b2 - md5: 1275bc6aedb6410368dd1bf6d98f6f90 - depends: - - python >=3.10 - - mojo-compiler ==0.26.2.0.dev2026022205 release - - mblack ==26.2.0.dev2026022205 release - - jupyter_client >=8.6.2,<8.7 - license: LicenseRef-Modular-Proprietary - size: 88491769 - timestamp: 1771738763383 -- conda: https://conda.modular.com/max-nightly/linux-64/mojo-compiler-0.26.2.0.dev2026022205-release.conda - sha256: 173dd0f2401b3342d7cf59c29745a84f2c1168976a6f102318631b81c370c6b4 - md5: c8db27767e0be9250b72e3414602ff37 - depends: - - mojo-python ==0.26.2.0.dev2026022205 release - license: LicenseRef-Modular-Proprietary - size: 87474357 - timestamp: 1771738763382 -- conda: https://conda.modular.com/max-nightly/noarch/mojo-python-0.26.2.0.dev2026022205-release.conda - noarch: python - sha256: 71199f2ae34cfd0cfcce57613448d3d54e000151908ca9564ccb8e1b6beb4ddd - md5: 7742228841ea05a513e667fbf57fdfb2 - depends: - - python - license: LicenseRef-Modular-Proprietary - size: 676288 - timestamp: 1771738763380 - conda: https://conda.anaconda.org/conda-forge/noarch/mypy_extensions-1.1.0-pyha770c72_0.conda sha256: 6ed158e4e5dd8f6a10ad9e525631e35cee8557718f83de7a4e3966b1f772c4b1 md5: e9c622e0d00fa24a6292279af3ab6d06 @@ -587,110 +743,51 @@ packages: - python >=3.9 license: MIT license_family: MIT + run_exports: {} size: 11766 timestamp: 1745776666688 -- conda: https://conda.anaconda.org/conda-forge/linux-64/ncurses-6.5-h2d0b736_3.conda - sha256: 3fde293232fa3fca98635e1167de6b7c7fda83caf24b9d6c91ec9eefb4f4d586 - md5: 47e340acb35de30501a76c7c799c41d7 - depends: - - __glibc >=2.17,<3.0.a0 - - libgcc >=13 - license: X11 AND BSD-3-Clause - size: 891641 - timestamp: 1738195959188 -- conda: https://conda.anaconda.org/conda-forge/linux-64/numpy-2.4.2-py314h2b28147_1.conda - sha256: 1d8377c8001c15ed12c2713b723213474b435706ab9d34ede69795d64af9e94d - md5: 4ea6b620fdf24a1a0bc4f1c7134dfafb - depends: - - python - - libstdcxx >=14 - - libgcc >=14 - - __glibc >=2.17,<3.0.a0 - - libcblas >=3.9.0,<4.0a0 - - python_abi 3.14.* *_cp314 - - libblas >=3.9.0,<4.0a0 - - liblapack >=3.9.0,<4.0a0 - constrains: - - numpy-base <0a0 - license: BSD-3-Clause - license_family: BSD - size: 8926994 - timestamp: 1770098474394 -- conda: https://conda.anaconda.org/conda-forge/linux-64/openssl-3.6.1-h35e630c_1.conda - sha256: 44c877f8af015332a5d12f5ff0fb20ca32f896526a7d0cdb30c769df1144fb5c - md5: f61eb8cd60ff9057122a3d338b99c00f - depends: - - __glibc >=2.17,<3.0.a0 - - ca-certificates - - libgcc >=14 - license: Apache-2.0 - license_family: Apache - size: 3164551 - timestamp: 1769555830639 -- conda: https://conda.anaconda.org/conda-forge/noarch/packaging-26.0-pyhcf101f3_0.conda - sha256: c1fc0f953048f743385d31c468b4a678b3ad20caffdeaa94bed85ba63049fd58 - md5: b76541e68fea4d511b1ac46a28dcd2c6 +- conda: https://conda.anaconda.org/conda-forge/noarch/packaging-26.2-pyhc364b38_0.conda + sha256: 3906abfb6511a3bb309e39b9b1b7bc38f50a723971de2395489fd1f379255890 + md5: 4c06a92e74452cfa53623a81592e8934 depends: - python >=3.8 - python license: Apache-2.0 license_family: APACHE - size: 72010 - timestamp: 1769093650580 -- conda: https://conda.anaconda.org/conda-forge/noarch/pathspec-1.0.4-pyhd8ed1ab_0.conda - sha256: 29ea20d0faf20374fcd61c25f6d32fb8e9a2c786a7f1473a0c3ead359470fbe1 - md5: 2908273ac396d2cd210a8127f5f1c0d6 + run_exports: {} + size: 91574 + timestamp: 1777103621679 +- conda: https://conda.anaconda.org/conda-forge/noarch/pathspec-1.1.1-pyhd8ed1ab_0.conda + sha256: 6eaee417d33f298db79bc7185ab1208604c0e6cf51dade34cd513c6f9db9c6f3 + md5: 11adc78451c998c0fd162584abfa3559 depends: - python >=3.10 license: MPL-2.0 license_family: MOZILLA - size: 53739 - timestamp: 1769677743677 -- conda: https://conda.anaconda.org/conda-forge/noarch/platformdirs-4.7.0-pyhcf101f3_0.conda - sha256: 5b14e300187919e783b24507dbeaeee891d97a3e9a9c80b5dcd3073753bac69c - md5: 2157d0900a4bc2e9a0ba3cccb8497e8c + run_exports: {} + size: 56559 + timestamp: 1777271601895 +- conda: https://conda.anaconda.org/conda-forge/noarch/platformdirs-4.11.0-pyhcf101f3_0.conda + sha256: a4b8c7d3b3703d10f93187986d59aec09b22033978945845899beb7f849a7be6 + md5: 1fadaa6dd1d03d062075f84157ac2cc7 depends: - python >=3.10 - python license: MIT - size: 24091 - timestamp: 1770990257318 -- conda: https://conda.anaconda.org/conda-forge/noarch/pygments-2.19.2-pyhd8ed1ab_0.conda - sha256: 5577623b9f6685ece2697c6eb7511b4c9ac5fb607c9babc2646c811b428fd46a - md5: 6b6ece66ebcae2d5f326c77ef2c5a066 + license_family: MIT + run_exports: {} + size: 26632 + timestamp: 1784661349391 +- conda: https://conda.anaconda.org/conda-forge/noarch/pygments-2.20.0-pyhd8ed1ab_0.conda + sha256: cf70b2f5ad9ae472b71235e5c8a736c9316df3705746de419b59d442e8348e86 + md5: 16c18772b340887160c79a6acc022db0 depends: - - python >=3.9 + - python >=3.10 license: BSD-2-Clause license_family: BSD - size: 889287 - timestamp: 1750615908735 -- conda: https://conda.anaconda.org/conda-forge/linux-64/python-3.14.3-h32b2ec7_101_cp314.conda - build_number: 101 - sha256: cb0628c5f1732f889f53a877484da98f5a0e0f47326622671396fb4f2b0cd6bd - md5: c014ad06e60441661737121d3eae8a60 - depends: - - __glibc >=2.17,<3.0.a0 - - bzip2 >=1.0.8,<2.0a0 - - ld_impl_linux-64 >=2.36.1 - - libexpat >=2.7.3,<3.0a0 - - libffi >=3.5.2,<3.6.0a0 - - libgcc >=14 - - liblzma >=5.8.2,<6.0a0 - - libmpdec >=4.0.0,<5.0a0 - - libsqlite >=3.51.2,<4.0a0 - - libuuid >=2.41.3,<3.0a0 - - libzlib >=1.3.1,<2.0a0 - - ncurses >=6.5,<7.0a0 - - openssl >=3.5.5,<4.0a0 - - python_abi 3.14.* *_cp314 - - readline >=8.3,<9.0a0 - - tk >=8.6.13,<8.7.0a0 - - tzdata - - zstd >=1.5.7,<1.6.0a0 - license: Python-2.0 - size: 36702440 - timestamp: 1770675584356 - python_site_packages_path: lib/python3.14/site-packages + run_exports: {} + size: 893031 + timestamp: 1774796815820 - conda: https://conda.anaconda.org/conda-forge/noarch/python-dateutil-2.9.0.post0-pyhe01879c_2.conda sha256: d6a17ece93bbd5139e02d2bd7dbfa80bee1a4261dced63f65f679121686bf664 md5: 5b8d21249ff20967101ffa321cab24e8 @@ -700,17 +797,19 @@ packages: - python license: Apache-2.0 license_family: APACHE + run_exports: {} size: 233310 timestamp: 1751104122689 -- conda: https://conda.anaconda.org/conda-forge/noarch/python-gil-3.14.3-h4df99d1_101.conda - sha256: 233aebd94c704ac112afefbb29cf4170b7bc606e22958906f2672081bc50638a - md5: 235765e4ea0d0301c75965985163b5a1 +- conda: https://conda.anaconda.org/conda-forge/noarch/python-gil-3.14.6-h4df99d1_101.conda + sha256: f96f61b33fe7d0f599ba5e23d9e5231fad9c8a37a1c141dfa8edbd0ba78de9e3 + md5: 7f742295acd62ee0688e7e6924b71e67 depends: - - cpython 3.14.3.* + - cpython 3.14.6.* - python_abi * *_cp314 license: Python-2.0 - size: 50062 - timestamp: 1770674497152 + run_exports: {} + size: 49484 + timestamp: 1784909578801 - conda: https://conda.anaconda.org/conda-forge/noarch/python_abi-3.14-8_cp314.conda build_number: 8 sha256: ad6d2e9ac39751cc0529dd1566a26751a0bf2542adb0c232533d32e176e21db5 @@ -719,51 +818,12 @@ packages: - python 3.14.* *_cp314 license: BSD-3-Clause license_family: BSD + run_exports: {} size: 6989 timestamp: 1752805904792 -- conda: https://conda.anaconda.org/conda-forge/linux-64/pyyaml-6.0.3-py314h67df5f8_1.conda - sha256: b318fb070c7a1f89980ef124b80a0b5ccf3928143708a85e0053cde0169c699d - md5: 2035f68f96be30dc60a5dfd7452c7941 - depends: - - __glibc >=2.17,<3.0.a0 - - libgcc >=14 - - python >=3.14,<3.15.0a0 - - python_abi 3.14.* *_cp314 - - yaml >=0.2.5,<0.3.0a0 - license: MIT - license_family: MIT - size: 202391 - timestamp: 1770223462836 -- conda: https://conda.anaconda.org/conda-forge/linux-64/pyzmq-27.1.0-py312hfb55c3c_0.conda - noarch: python - sha256: a00a41b66c12d9c60e66b391e9a4832b7e28743348cf4b48b410b91927cd7819 - md5: 3399d43f564c905250c1aea268ebb935 - depends: - - python - - __glibc >=2.17,<3.0.a0 - - libstdcxx >=14 - - libgcc >=14 - - _python_abi3_support 1.* - - cpython >=3.12 - - zeromq >=4.3.5,<4.4.0a0 - license: BSD-3-Clause - license_family: BSD - size: 212218 - timestamp: 1757387023399 -- conda: https://conda.anaconda.org/conda-forge/linux-64/readline-8.3-h853b02a_0.conda - sha256: 12ffde5a6f958e285aa22c191ca01bbd3d6e710aa852e00618fa6ddc59149002 - md5: d7d95fc8287ea7bf33e0e7116d2b95ec - depends: - - __glibc >=2.17,<3.0.a0 - - libgcc >=14 - - ncurses >=6.5,<7.0a0 - license: GPL-3.0-only - license_family: GPL - size: 345073 - timestamp: 1765813471974 -- conda: https://conda.anaconda.org/conda-forge/noarch/rich-14.3.2-pyhcf101f3_0.conda - sha256: ed17985cec5a0540002c6cabe67848f7cc17e5f4019c0e2a40534e9b7c0b38de - md5: 33950a076fd589a7655c6888cc3d2b34 +- conda: https://conda.anaconda.org/conda-forge/noarch/rich-15.0.0-pyhcf101f3_0.conda + sha256: 3d6ba2c0fcdac3196ba2f0615b4104e532525ffa1335b50a2878be5ff488814a + md5: 0242025a3c804966bf71aa04eee82f66 depends: - markdown-it-py >=2.2.0 - pygments >=2.13.0,<3.0.0 @@ -772,8 +832,9 @@ packages: - python license: MIT license_family: MIT - size: 208269 - timestamp: 1769971520792 + run_exports: {} + size: 208577 + timestamp: 1775991661559 - conda: https://conda.anaconda.org/conda-forge/noarch/six-1.17.0-pyhe01879c_1.conda sha256: 458227f759d5e3fcec5d9b7acce54e10c9e1f4f4b7ec978f3bfd54ce4ee9853d md5: 3339e3b65d58accf4ca4fb8748ab16b3 @@ -782,118 +843,154 @@ packages: - python license: MIT license_family: MIT + run_exports: {} size: 18455 timestamp: 1753199211006 -- conda: https://conda.anaconda.org/conda-forge/linux-64/tk-8.6.13-noxft_h366c992_103.conda - sha256: cafeec44494f842ffeca27e9c8b0c27ed714f93ac77ddadc6aaf726b5554ebac - md5: cffd3bdd58090148f4cfcd831f4b26ab +- conda: https://conda.anaconda.org/conda-forge/noarch/taskgroup-0.2.2-pyhd8ed1ab_0.conda + sha256: 6f8db6da8de445930de55b708e6a5d3ab5f076bc14a39578db0190b2a9b8e437 + md5: 9fa69537fb68a095fbac139210575bad depends: - - __glibc >=2.17,<3.0.a0 - - libgcc >=14 - - libzlib >=1.3.1,<2.0a0 - constrains: - - xorg-libx11 >=1.8.12,<2.0a0 - license: TCL - license_family: BSD - size: 3301196 - timestamp: 1769460227866 -- conda: https://conda.anaconda.org/conda-forge/noarch/tomli-2.4.0-pyhcf101f3_0.conda - sha256: 62940c563de45790ba0f076b9f2085a842a65662268b02dd136a8e9b1eaf47a8 - md5: 72e780e9aa2d0a3295f59b1874e3768b + - exceptiongroup + - python >=3.9 + - typing_extensions >=4.12.2,<5 + license: MIT + license_family: MIT + run_exports: {} + size: 17330 + timestamp: 1736003478648 +- conda: https://conda.anaconda.org/conda-forge/noarch/tomli-2.4.1-pyhcf101f3_0.conda + sha256: 91cafdb64268e43e0e10d30bd1bef5af392e69f00edd34dfaf909f69ab2da6bd + md5: b5325cf06a000c5b14970462ff5e4d58 depends: - python >=3.10 - python license: MIT license_family: MIT - size: 21453 - timestamp: 1768146676791 -- conda: https://conda.anaconda.org/conda-forge/linux-64/tornado-6.5.3-py314h5bd0f2a_0.conda - sha256: b8f9f9ae508d79c9c697eb01b6a8d2ed4bc1899370f44aa6497c8abbd15988ea - md5: e35f08043f54d26a1be93fdbf90d30c3 - depends: - - __glibc >=2.17,<3.0.a0 - - libgcc >=14 - - python >=3.14,<3.15.0a0 - - python_abi 3.14.* *_cp314 - license: Apache-2.0 - license_family: Apache - size: 905436 - timestamp: 1765458949518 -- conda: https://conda.anaconda.org/conda-forge/noarch/traitlets-5.14.3-pyhd8ed1ab_1.conda - sha256: f39a5620c6e8e9e98357507262a7869de2ae8cc07da8b7f84e517c9fd6c2b959 - md5: 019a7385be9af33791c989871317e1ed + run_exports: {} + size: 21561 + timestamp: 1774492402955 +- conda: https://conda.anaconda.org/conda-forge/noarch/traitlets-5.15.1-pyhcf101f3_0.conda + sha256: b89a823edf524956b94a2a4db974866e4501f05c68976eff458c5dcf07f88431 + md5: 37e3be7b6e2977d37b8fa5da229f5dc0 depends: - - python >=3.9 + - python >=3.10 + - python license: BSD-3-Clause license_family: BSD - size: 110051 - timestamp: 1733367480074 -- conda: https://conda.anaconda.org/conda-forge/noarch/typing-extensions-4.15.0-h396c80c_0.conda - sha256: 7c2df5721c742c2a47b2c8f960e718c930031663ac1174da67c1ed5999f7938c - md5: edd329d7d3a4ab45dcf905899a7a6115 - depends: - - typing_extensions ==4.15.0 pyhcf101f3_0 + run_exports: {} + size: 115158 + timestamp: 1780507822178 +- conda: https://conda.anaconda.org/conda-forge/noarch/typing-extensions-4.16.0-h69aa097_0.conda + sha256: b141933ece3518f6d7b75dfb59451e2f26b405a44c18e2518a83e9a02e09315c + md5: c680b5747e8c4c8f23dca0bb7042a8fc + depends: + - typing_extensions ==4.16.0 pyhcf101f3_0 license: PSF-2.0 license_family: PSF - size: 91383 - timestamp: 1756220668932 -- conda: https://conda.anaconda.org/conda-forge/noarch/typing_extensions-4.15.0-pyhcf101f3_0.conda - sha256: 032271135bca55aeb156cee361c81350c6f3fb203f57d024d7e5a1fc9ef18731 - md5: 0caa1af407ecff61170c9437a808404d + run_exports: {} + size: 94080 + timestamp: 1783002732887 +- conda: https://conda.anaconda.org/conda-forge/noarch/typing_extensions-4.16.0-pyhcf101f3_0.conda + sha256: 2d888f90af0686044882c74193ec80a90ec1943145d94a7b1b048958acda1848 + md5: c70ad746c22219b9700931707482992c depends: - python >=3.10 - python license: PSF-2.0 license_family: PSF - size: 51692 - timestamp: 1756220668932 -- conda: https://conda.anaconda.org/conda-forge/noarch/tzdata-2025c-hc9c84f9_1.conda - sha256: 1d30098909076af33a35017eed6f2953af1c769e273a0626a04722ac4acaba3c - md5: ad659d0a2b3e47e38d829aa8cad2d610 + run_exports: {} + size: 52631 + timestamp: 1783002732887 +- conda: https://conda.anaconda.org/conda-forge/noarch/tzdata-2026c-h151e31d_0.conda + sha256: b928c30ddcb0e3f544c6eade8352737e6e610e263276b90232db6a578ef899d8 + md5: fcb489df604d100968b737f2cb6076c6 license: LicenseRef-Public-Domain - size: 119135 - timestamp: 1767016325805 -- conda: https://conda.anaconda.org/conda-forge/linux-64/yaml-0.2.5-h280c20c_3.conda - sha256: 6d9ea2f731e284e9316d95fa61869fe7bbba33df7929f82693c121022810f4ad - md5: a77f85f77be52ff59391544bfe73390a + run_exports: {} + size: 118849 + timestamp: 1784250406640 +- conda: https://conda.anaconda.org/conda-forge/noarch/zipp-4.1.0-pyhcf101f3_0.conda + sha256: 210bd31c22bb88f5e2a167df24c95bb5f152b2ada7502f9b8c49d1f5366db423 + md5: ba3dcdc8584155c97c648ae9c044b7a3 depends: - - libgcc >=14 - - __glibc >=2.17,<3.0.a0 + - python >=3.10 + - python license: MIT license_family: MIT - size: 85189 - timestamp: 1753484064210 -- conda: https://conda.anaconda.org/conda-forge/linux-64/zeromq-4.3.5-h387f397_9.conda - sha256: 47cfe31255b91b4a6fa0e9dbaf26baa60ac97e033402dbc8b90ba5fee5ffe184 - md5: 8035e5b54c08429354d5d64027041cad + run_exports: {} + size: 24190 + timestamp: 1779159948016 +- conda: https://conda.modular.com/max-nightly/linux-64/max-26.5.0.dev2026072505-3.14release.conda + sha256: 735cc65e1938c0c0e5182d96110a144239782b5d954ebddbad4dfcb0a246637b + md5: ffe92be5964a56b949a8c782a3f0a98d depends: - - libstdcxx >=14 - - libgcc >=14 - - __glibc >=2.17,<3.0.a0 - - libgcc >=14 - - libsodium >=1.0.20,<1.0.21.0a0 - - krb5 >=1.21.3,<1.22.0a0 - license: MPL-2.0 - license_family: MOZILLA - size: 310648 - timestamp: 1757370847287 -- conda: https://conda.anaconda.org/conda-forge/noarch/zipp-3.23.0-pyhcf101f3_1.conda - sha256: b4533f7d9efc976511a73ef7d4a2473406d7f4c750884be8e8620b0ce70f4dae - md5: 30cd29cb87d819caead4d55184c1d115 + - click >=8.0.0 + - exceptiongroup >=0.2.2 + - msgspec >=0.19.0 + - numpy >=1.18 + - psutil >=7.0.0 + - rich >=13.0.1 + - taskgroup >=0.2.2 + - typing-extensions >=4.12.2 + - python 3.14.* + - python-gil + - max-core ==26.5.0.dev2026072505 + license: LicenseRef-Modular-Proprietary + run_exports: {} + size: 18225480 + timestamp: 1784959682216 +- conda: https://conda.modular.com/max-nightly/linux-64/max-core-26.5.0.dev2026072505-release.conda + sha256: a2caecf6f29447375ce1bb42979d998c13acf5f48f5b16c7dc6166825c005cf0 + md5: 20809390abf8b0e993718841eb16813b + depends: + - mojo-compiler ==1.0.0b3.dev2026072505 + license: LicenseRef-Modular-Proprietary + run_exports: {} + size: 103408905 + timestamp: 1784959675208 +- conda: https://conda.modular.com/max-nightly/linux-64/mojo-1.0.0b3.dev2026072505-release.conda + sha256: 9fa11e681371d1994ffbe66d0f7e1502e7b22381c5ca8d5f844b77a335c8e19b + md5: 6a0c4ee8b9b53aba98644702a4c2c682 depends: - python >=3.10 - - python - license: MIT - license_family: MIT - size: 24194 - timestamp: 1764460141901 -- conda: https://conda.anaconda.org/conda-forge/linux-64/zstd-1.5.7-hb78ec9c_6.conda - sha256: 68f0206ca6e98fea941e5717cec780ed2873ffabc0e1ed34428c061e2c6268c7 - md5: 4a13eeac0b5c8e5b8ab496e6c4ddd829 + - mojo-compiler ==1.0.0b3.dev2026072505 + - mblack ==26.5.0.dev2026072505 + - jupyter_client >=8.6.2,<8.7 + license: LicenseRef-Modular-Proprietary + run_exports: {} + size: 114479370 + timestamp: 1784959545074 +- conda: https://conda.modular.com/max-nightly/linux-64/mojo-compiler-1.0.0b3.dev2026072505-release.conda + sha256: 307aae574f073cf914c9efe3cfadeeb2f3992c3a426154731e934b81f7b2f725 + md5: b9d5de41a9c940f47090818f8d9d0c18 + depends: + - mojo-python ==1.0.0b3.dev2026072505 + license: LicenseRef-Modular-Proprietary + run_exports: {} + size: 80560693 + timestamp: 1784959541118 +- conda: https://conda.modular.com/max-nightly/noarch/mblack-26.5.0.dev2026072505-release.conda + noarch: python + sha256: a66e15918fe18928e4a35c08e16f0517c25d5bc10f836a2240c7cbf9788f9501 + md5: 19378925651cd95d631a7f580cb2846d depends: - - __glibc >=2.17,<3.0.a0 - - libzlib >=1.3.1,<2.0a0 - license: BSD-3-Clause - license_family: BSD - size: 601375 - timestamp: 1764777111296 + - python >=3.10 + - click >=8.0.0 + - mypy_extensions >=0.4.3 + - packaging >=22.0 + - pathspec >=0.9.0 + - platformdirs >=2 + - tomli >=1.1.0 + license: LicenseRef-Modular-Proprietary + run_exports: {} + size: 136731 + timestamp: 1784959402726 +- conda: https://conda.modular.com/max-nightly/noarch/mojo-python-1.0.0b3.dev2026072505-release.conda + noarch: python + sha256: 9d90f9d14440c864e96c331fb54e5a2be5fc724bdcdec0adf9ddd52d1dabdcb9 + md5: f5e33c79f95b3cb3bd7e7d4bfe227eb2 + depends: + - python >=3.10 + license: LicenseRef-Modular-Proprietary + run_exports: {} + size: 24107 + timestamp: 1784959402688 diff --git a/benchmark/pixi.toml b/mojo/pixi.toml similarity index 67% rename from benchmark/pixi.toml rename to mojo/pixi.toml index 46776c37af..66bc00e163 100644 --- a/benchmark/pixi.toml +++ b/mojo/pixi.toml @@ -1,12 +1,12 @@ [workspace] authors = ["maxchisto "] channels = ["https://conda.modular.com/max-nightly", "conda-forge"] -name = "benchmark" +name = "mojo" platforms = ["linux-64"] version = "0.1.0" [tasks] [dependencies] -mojo = ">=0.26.2.0.dev2026022205,<0.27" -max = ">=26.2.0.dev2026022205,<27" +mojo = ">=1.0.0b3.dev2026072505,<2" +max = ">=26.5.0.dev2026072505" diff --git a/mojo/sigmoid_build_instructions.md b/mojo/sigmoid_build_instructions.md new file mode 100644 index 0000000000..dc107c6639 --- /dev/null +++ b/mojo/sigmoid_build_instructions.md @@ -0,0 +1,39 @@ +# Build Instructions + +## 1. Compile the Mojo Module (`libsigmoid_mojo.so`) + +The Mojo code provides the core processing logic and kernels. + +```bash +cd ~/code/darktable/mojo +make build +``` + +The resulting `libsigmoid_mojo.so` contains the following exported symbols used by the C bridge: +- `sigmoid_mojo_init` +- `sigmoid_mojo_destroy` +- `sigmoid_mojo_rgb_ratio` +- `sigmoid_mojo_per_channel` + +## 2. Compile the C-Bridge Plugin (`libsigmoid.so`) + +The C part (`src/iop/sigmoid.c`) handles the darktable user interface and parameter management. It loads the Mojo shared library at runtime. + +```bash +./build_sigmoid_iop.sh +``` + + +## 3. Installation + +Both shared libraries must be placed in the darktable plugins directory so that darktable can find the plugin and the plugin can find the Mojo library. + +```bash +# Copy the Mojo library +sudo cp libsigmoid_mojo.so /usr/lib/darktable/plugins/ + +# Copy the C bridge plugin +sudo cp libsigmoid.so /usr/lib/darktable/plugins/ +``` + +> The C code in `src/iop/sigmoid.c` uses `dlopen("libsigmoid_mojo.so", RTLD_LAZY | RTLD_LOCAL)` in `init_global` to load the Mojo module. diff --git a/src/external/OpenCL b/src/external/OpenCL deleted file mode 160000 index 8a97ebc88d..0000000000 --- a/src/external/OpenCL +++ /dev/null @@ -1 +0,0 @@ -Subproject commit 8a97ebc88daa3495d6f57ec10bb515224400186f diff --git a/src/external/OpenCL/.gitignore b/src/external/OpenCL/.gitignore new file mode 100644 index 0000000000..e69de29bb2 diff --git a/src/external/lua-scripts b/src/external/lua-scripts deleted file mode 160000 index db505f1e0a..0000000000 --- a/src/external/lua-scripts +++ /dev/null @@ -1 +0,0 @@ -Subproject commit db505f1e0a089b210f5afa653e53ebd32e5a82b0 diff --git a/src/external/lua-scripts/.gitignore b/src/external/lua-scripts/.gitignore new file mode 100644 index 0000000000..e69de29bb2 diff --git a/src/iop/sigmoid.c b/src/iop/sigmoid.c index cbd96d7725..b9301431b7 100644 --- a/src/iop/sigmoid.c +++ b/src/iop/sigmoid.c @@ -27,9 +27,11 @@ #include "gui/gtk.h" #include "gui/presets.h" #include "iop/iop_api.h" +#include "common/opencl.h" #include #include +#include DT_MODULE_INTROSPECTION(3, dt_iop_sigmoid_params_t) @@ -180,10 +182,38 @@ typedef struct dt_iop_sigmoid_gui_data_t dt_gui_collapsible_section_t display_luminance_section, primaries_section; } dt_iop_sigmoid_gui_data_t; +typedef struct { + float white_target[4]; + float black_target[4]; + float paper_exposure[4]; + float film_fog[4]; + float film_power[4]; + float paper_power[4]; + float contrast_power[4]; + float skew_power[4]; + float hue_preservation[4]; + float pipe_to_base[16]; + float base_to_rendering[16]; + float rendering_to_pipe[16]; +} SigmoidMojoParams; + +typedef void (*mojo_init_fn)(uintptr_t *ctx, int use_gpu); +typedef void (*mojo_destroy_fn)(uintptr_t ctx); +typedef void (*mojo_rgb_ratio_fn)(uintptr_t ctx, float *in, float *out, int32_t width, int32_t height, void *params); +typedef void (*mojo_per_channel_fn)(uintptr_t ctx, float *in, float *out, int32_t width, int32_t height, void *params); + typedef struct dt_iop_sigmoid_global_data_t { int kernel_sigmoid_loglogistic_per_channel; int kernel_sigmoid_loglogistic_rgb_ratio; + + void *mojo_lib; + uintptr_t mojo_ctx_cpu; + uintptr_t mojo_ctx_gpu; + mojo_init_fn mojo_init; + mojo_destroy_fn mojo_destroy; + mojo_rgb_ratio_fn mojo_rgb_ratio; + mojo_per_channel_fn mojo_per_channel; } dt_iop_sigmoid_global_data_t; @@ -760,6 +790,40 @@ void process_loglogistic_per_channel(dt_develop_t *dev, } } +static SigmoidMojoParams _build_mojo_params(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece) +{ + const dt_iop_sigmoid_data_t *d = piece->data; + SigmoidMojoParams p = {0}; + + for(int i = 0; i < 4; i++) { + p.white_target[i] = d->white_target; + p.black_target[i] = d->black_target; + p.paper_exposure[i] = d->paper_exposure; + p.film_fog[i] = d->film_fog; + p.film_power[i] = d->film_power; + p.paper_power[i] = d->paper_power; + p.contrast_power[i] = d->film_power; + p.skew_power[i] = d->paper_power; + p.hue_preservation[i] = d->hue_preservation; + } + + const dt_iop_order_iccprofile_info_t *pipe_work_profile = dt_ioppr_get_pipe_work_profile_info(piece->pipe); + const dt_iop_order_iccprofile_info_t *base_profile = _get_base_profile(self->dev, pipe_work_profile, d->base_primaries); + dt_colormatrix_t pipe_to_base_transposed, base_to_rendering_transposed, rendering_to_pipe_transposed; + dt_colormatrix_t pipe_to_base, base_to_rendering, rendering_to_pipe; + + _calculate_adjusted_primaries(d, pipe_work_profile, base_profile, pipe_to_base_transposed, base_to_rendering_transposed, rendering_to_pipe_transposed); + transpose_3xSSE(pipe_to_base_transposed, pipe_to_base); + transpose_3xSSE(base_to_rendering_transposed, base_to_rendering); + transpose_3xSSE(rendering_to_pipe_transposed, rendering_to_pipe); + + memcpy(p.pipe_to_base, pipe_to_base, sizeof(pipe_to_base)); + memcpy(p.base_to_rendering, base_to_rendering, sizeof(base_to_rendering)); + memcpy(p.rendering_to_pipe, rendering_to_pipe, sizeof(rendering_to_pipe)); + + return p; +} + /** process, all real work is done here. */ void process(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, @@ -768,8 +832,22 @@ void process(dt_iop_module_t *self, const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out) { - // this is called for preview and full pipe separately, each with its own pixelpipe piece. const dt_iop_sigmoid_data_t *module_data = piece->data; + const dt_iop_sigmoid_global_data_t *gd = self->global_data; + + if(gd && gd->mojo_lib && gd->mojo_per_channel && gd->mojo_rgb_ratio) + { + SigmoidMojoParams p = _build_mojo_params(self, piece); + const int w = roi_in->width, h = roi_in->height; + int use_gpu = dt_opencl_is_enabled() ? 1 : 0; + uintptr_t ctx = use_gpu ? gd->mojo_ctx_gpu : gd->mojo_ctx_cpu; + + if(module_data->color_processing == DT_SIGMOID_METHOD_PER_CHANNEL) + gd->mojo_per_channel(ctx, (float *)ivoid, (float *)ovoid, w, h, &p); + else + gd->mojo_rgb_ratio(ctx, (float *)ivoid, (float *)ovoid, w, h, &p); + return; + } if(module_data->color_processing == DT_SIGMOID_METHOD_PER_CHANNEL) { @@ -781,87 +859,42 @@ void process(dt_iop_module_t *self, } } -#ifdef HAVE_OPENCL -int process_cl(dt_iop_module_t *self, - dt_dev_pixelpipe_iop_t *piece, - cl_mem dev_in, - cl_mem dev_out, - const dt_iop_roi_t *const roi_in, - const dt_iop_roi_t *const roi_out) -{ - const dt_iop_sigmoid_data_t *const d = piece->data; - const dt_iop_sigmoid_global_data_t *const gd = self->global_data; - cl_int err = CL_MEM_OBJECT_ALLOCATION_FAILURE; - const int devid = piece->pipe->devid; - const int width = roi_in->width; - const int height = roi_in->height; - const float white_target = d->white_target; - const float paper_exp = d->paper_exposure; - const float film_fog = d->film_fog; - const float contrast_power = d->film_power; - const float skew_power = d->paper_power; +void init_global(dt_iop_module_so_t *self) +{ + dt_iop_sigmoid_global_data_t *gd = calloc(1, sizeof(dt_iop_sigmoid_global_data_t)); + self->data = gd; - const dt_iop_order_iccprofile_info_t *pipe_work_profile = dt_ioppr_get_pipe_work_profile_info(piece->pipe); - const dt_iop_order_iccprofile_info_t *base_profile = _get_base_profile(self->dev, pipe_work_profile, d->base_primaries); - dt_colormatrix_t pipe_to_base_transposed, base_to_rendering_transposed, - rendering_to_pipe_transposed, pipe_to_base, base_to_rendering, rendering_to_pipe; - _calculate_adjusted_primaries(d, pipe_work_profile, base_profile, pipe_to_base_transposed, base_to_rendering_transposed, rendering_to_pipe_transposed); - transpose_3xSSE(pipe_to_base_transposed, pipe_to_base); - transpose_3xSSE(base_to_rendering_transposed, base_to_rendering); - transpose_3xSSE(rendering_to_pipe_transposed, rendering_to_pipe); - const cl_mem dev_pipe_to_base - = dt_opencl_copy_host_to_device_constant(devid, sizeof(pipe_to_base), pipe_to_base); - const cl_mem dev_base_to_rendering - = dt_opencl_copy_host_to_device_constant(devid, sizeof(base_to_rendering), base_to_rendering); - const cl_mem dev_rendering_to_pipe - = dt_opencl_copy_host_to_device_constant(devid, sizeof(rendering_to_pipe), rendering_to_pipe); - if(dev_pipe_to_base == NULL || dev_base_to_rendering == NULL || dev_rendering_to_pipe == NULL) - goto cleanup; - - if(d->color_processing == DT_SIGMOID_METHOD_PER_CHANNEL) + gd->mojo_lib = dlopen("libsigmoid_mojo.so", RTLD_LAZY | RTLD_LOCAL); + if(gd->mojo_lib) { - const float hue_preservation = d->hue_preservation; - err = dt_opencl_enqueue_kernel_2d_args( - devid, gd->kernel_sigmoid_loglogistic_per_channel, width, height, CLARG(dev_in), CLARG(dev_out), - CLARG(width), CLARG(height), CLARG(white_target), CLARG(paper_exp), CLARG(film_fog), CLARG(contrast_power), - CLARG(skew_power), CLARG(hue_preservation), CLARG(dev_pipe_to_base), CLARG(dev_base_to_rendering), CLARG(dev_rendering_to_pipe)); - } - else - { - const float black_target = d->black_target; + gd->mojo_init = (mojo_init_fn)dlsym(gd->mojo_lib, "sigmoid_mojo_init"); + gd->mojo_destroy = (mojo_destroy_fn)dlsym(gd->mojo_lib, "sigmoid_mojo_destroy"); + gd->mojo_rgb_ratio = (mojo_rgb_ratio_fn)dlsym(gd->mojo_lib, "sigmoid_mojo_rgb_ratio"); + gd->mojo_per_channel = (mojo_per_channel_fn)dlsym(gd->mojo_lib, "sigmoid_mojo_per_channel"); - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_sigmoid_loglogistic_rgb_ratio, width, height, - CLARG(dev_in), CLARG(dev_out), CLARG(width), CLARG(height), - CLARG(white_target), CLARG(black_target), CLARG(paper_exp), - CLARG(film_fog), CLARG(contrast_power), CLARG(skew_power)); + if (gd->mojo_init && gd->mojo_destroy) + { + gd->mojo_init(&gd->mojo_ctx_cpu, 0); // CPU context + gd->mojo_init(&gd->mojo_ctx_gpu, 1); // GPU context + } } - -cleanup: - dt_opencl_release_mem_object(dev_pipe_to_base); - dt_opencl_release_mem_object(dev_base_to_rendering); - dt_opencl_release_mem_object(dev_rendering_to_pipe); - return err; -} -#endif // HAVE_OPENCL - -void init_global(dt_iop_module_so_t *self) -{ - const int program = 36; // sigmoid.cl, from programs.conf - dt_iop_sigmoid_global_data_t *gd = malloc(sizeof(dt_iop_sigmoid_global_data_t)); - - self->data = gd; - gd->kernel_sigmoid_loglogistic_per_channel = dt_opencl_create_kernel(program, "sigmoid_loglogistic_per_channel"); - gd->kernel_sigmoid_loglogistic_rgb_ratio = dt_opencl_create_kernel(program, "sigmoid_loglogistic_rgb_ratio"); } void cleanup_global(dt_iop_module_so_t *self) { - const dt_iop_sigmoid_global_data_t *gd = self->data; - dt_opencl_free_kernel(gd->kernel_sigmoid_loglogistic_per_channel); - dt_opencl_free_kernel(gd->kernel_sigmoid_loglogistic_rgb_ratio); - free(self->data); + dt_iop_sigmoid_global_data_t *gd = self->data; + if(gd) + { + if(gd->mojo_destroy) + { + if(gd->mojo_ctx_cpu) gd->mojo_destroy(gd->mojo_ctx_cpu); + if(gd->mojo_ctx_gpu) gd->mojo_destroy(gd->mojo_ctx_gpu); + } + if(gd->mojo_lib) dlclose(gd->mojo_lib); + free(gd); + } self->data = NULL; }