diff --git a/src/maxdiffusion/configs/base_flux2klein.yml b/src/maxdiffusion/configs/base_flux2klein.yml index f2813c8fd..16a8a5cfa 100644 --- a/src/maxdiffusion/configs/base_flux2klein.yml +++ b/src/maxdiffusion/configs/base_flux2klein.yml @@ -40,6 +40,8 @@ max_sequence_length: 512 time_shift: True base_shift: 0.5 max_shift: 1.15 +image_paths: [] +use_base2_exp: True unet_checkpoint: '' @@ -74,7 +76,20 @@ mask_padding_tokens: True # in cross attention q. attention_sharding_uniform: True -flash_block_sizes: {} +flash_block_sizes: { + "block_q": 4608, + "block_kv": 1024, + "block_kv_compute": 1024, +} +ulysses_shards: 2 +ulysses_attention_chunks: 1 +text_encoder_attention: 'flash' +text_encoder_flash_block_sizes: { + "block_q": 512, + "block_kv": 512, + "block_kv_compute": 512, +} +text_encoder_max_layer: 27 # GroupNorm groups norm_num_groups: 32 @@ -154,12 +169,12 @@ data_sharding: [['data', 'fsdp', 'context', 'tensor']] # value to auto-shard based on available slices and devices. # By default, product of the DCN axes should equal number of slices # and product of the ICI axes should equal number of devices per slice. -dcn_data_parallelism: 1 # recommended DCN axis to be auto-sharded -dcn_fsdp_parallelism: -1 +dcn_data_parallelism: 1 +dcn_fsdp_parallelism: 1 dcn_context_parallelism: 1 dcn_tensor_parallelism: 1 ici_data_parallelism: 1 -ici_fsdp_parallelism: -1 +ici_fsdp_parallelism: 1 ici_context_parallelism: 1 ici_tensor_parallelism: 1 @@ -203,7 +218,7 @@ num_train_epochs: 1 seed: 0 output_dir: 'output/' output_name: "flux2klein_generated_image.png" -per_device_batch_size: 1 +per_device_batch_size: 1.0 warmup_steps_fraction: 0.1 learning_rate_schedule_steps: -1 # By default the length of the schedule is set to the number of steps. @@ -231,6 +246,7 @@ do_classifier_free_guidance: True guidance_scale: 4.0 guidance_rescale: 0.0 num_inference_steps: 4 +num_reps: 1 save_final_checkpoint: False # SDXL Lightning parameters diff --git a/src/maxdiffusion/configs/base_flux2klein_9B.yml b/src/maxdiffusion/configs/base_flux2klein_9B.yml index a6c670a69..c6dc2689b 100644 --- a/src/maxdiffusion/configs/base_flux2klein_9B.yml +++ b/src/maxdiffusion/configs/base_flux2klein_9B.yml @@ -40,6 +40,8 @@ max_sequence_length: 512 time_shift: True base_shift: 0.5 max_shift: 1.15 +image_paths: [] +use_base2_exp: True unet_checkpoint: '' @@ -74,7 +76,20 @@ mask_padding_tokens: True # in cross attention q. attention_sharding_uniform: True -flash_block_sizes: {} +flash_block_sizes: { + "block_q": 4608, + "block_kv": 1024, + "block_kv_compute": 1024, +} +ulysses_shards: 2 +ulysses_attention_chunks: 1 +text_encoder_attention: 'flash' +text_encoder_flash_block_sizes: { + "block_q": 512, + "block_kv": 512, + "block_kv_compute": 512, +} +text_encoder_max_layer: 27 # GroupNorm groups norm_num_groups: 32 @@ -154,12 +169,12 @@ data_sharding: [['data', 'fsdp', 'context', 'tensor']] # value to auto-shard based on available slices and devices. # By default, product of the DCN axes should equal number of slices # and product of the ICI axes should equal number of devices per slice. -dcn_data_parallelism: 1 # recommended DCN axis to be auto-sharded -dcn_fsdp_parallelism: -1 +dcn_data_parallelism: 1 +dcn_fsdp_parallelism: 1 dcn_context_parallelism: 1 dcn_tensor_parallelism: 1 ici_data_parallelism: 1 -ici_fsdp_parallelism: -1 # recommended ICI axis to be auto-sharded +ici_fsdp_parallelism: 1 ici_context_parallelism: 1 ici_tensor_parallelism: 1 @@ -203,7 +218,7 @@ num_train_epochs: 1 seed: 0 output_dir: 'output/' output_name: "flux2klein_generated_image.png" -per_device_batch_size: 1 +per_device_batch_size: 1.0 warmup_steps_fraction: 0.1 learning_rate_schedule_steps: -1 # By default the length of the schedule is set to the number of steps. @@ -231,6 +246,7 @@ do_classifier_free_guidance: True guidance_scale: 4.0 guidance_rescale: 0.0 num_inference_steps: 4 +num_reps: 1 save_final_checkpoint: False # SDXL Lightning parameters diff --git a/src/maxdiffusion/generate_flux2klein.py b/src/maxdiffusion/generate_flux2klein.py index 7956c850d..ad84d3bbd 100644 --- a/src/maxdiffusion/generate_flux2klein.py +++ b/src/maxdiffusion/generate_flux2klein.py @@ -20,6 +20,7 @@ import sys from typing import List +from PIL import Image, UnidentifiedImageError from absl import app import jax import jax.numpy as jnp @@ -35,11 +36,11 @@ from maxdiffusion.max_utils import create_device_mesh from maxdiffusion.train_utils import transformer_engine_context -from maxdiffusion.models.flux.transformers.transformer_flux_flax import Flux2KleinTransformer2DModel from maxdiffusion.models.vae_flax import FlaxAutoencoderKL from maxdiffusion.models.qwen3_flax import FlaxQwen3Config, FlaxQwen3Model from maxdiffusion.models.qwen3_utils import load_and_convert_qwen3_weights from maxdiffusion.schedulers.scheduling_flow_match_flax import FlaxFlowMatchScheduler +from maxdiffusion.pipelines.flux.flux2klein_pipeline import FlaxFlux2KleinPipeline def partition_prompts(prompt_str: str, batch_size: int) -> List[str]: @@ -79,8 +80,22 @@ def encode_prompt(prompt: str, snapshot_dir: str = None, repo_id: str = "black-f text_encoder_path = os.path.join(snapshot_dir, "text_encoder") tokenizer_path = os.path.join(snapshot_dir, "tokenizer") - if not os.path.exists(tokenizer_path): - tokenizer_path = text_encoder_path + + if not os.path.exists(os.path.join(text_encoder_path, "config.json")) or not os.path.exists(tokenizer_path): + try: + fb_dir = snapshot_download(repo_id=repo_id, local_files_only=True) + if not os.path.exists(os.path.join(text_encoder_path, "config.json")): + text_encoder_path = os.path.join(fb_dir, "text_encoder") + if not os.path.exists(tokenizer_path): + tokenizer_path = ( + os.path.join(fb_dir, "tokenizer") + if os.path.exists(os.path.join(fb_dir, "tokenizer")) + else os.path.join(fb_dir, "text_encoder") + ) + except Exception: + if not os.path.exists(tokenizer_path): + tokenizer_path = text_encoder_path + tokenizer = AutoTokenizer.from_pretrained(tokenizer_path) text_encoder = AutoModelForCausalLM.from_pretrained(text_encoder_path, torch_dtype=torch.float32) text_encoder.eval() @@ -132,25 +147,56 @@ def main(argv): # Import modules after jax.distributed.initialize() has run via pyconfig.initialize() from maxdiffusion.models.flux.util import ( - load_and_convert_flux_klein_weights, load_and_convert_vae_weights, - cast_dict_to_bfloat16_inplace, ) - from maxdiffusion.pipelines.flux.flux2klein_pipeline import FlaxFlux2KleinPipeline config = pyconfig.config os.makedirs(config.output_dir, exist_ok=True) + num_devices_to_use = getattr(config, "num_devices", None) + if num_devices_to_use is not None and num_devices_to_use > 0: + active_devices = jax.devices()[:num_devices_to_use] + else: + active_devices = jax.devices() + active_device_count = len(active_devices) + + if hasattr(config, "per_device_batch_size") and config.per_device_batch_size > 0: + calculated_batch_size = int(config.per_device_batch_size * active_device_count) + assert calculated_batch_size >= 1, ( + f"Calculated global batch_size is {calculated_batch_size}, which is invalid (must be >= 1). " + f"per_device_batch_size={config.per_device_batch_size} multiplied by active_device_count={active_device_count} " + f"evaluated to {config.per_device_batch_size * active_device_count}, which truncates to 0. " + f"Please increase per_device_batch_size or specify an explicit batch_size in your configuration." + ) + if calculated_batch_size != config.batch_size: + max_logging.log( + f"ℹ️ Updating batch_size from {config.batch_size} to {calculated_batch_size} " + f"based on per_device_batch_size={config.per_device_batch_size} and active_device_count={active_device_count}." + ) + pyconfig._config.keys["batch_size"] = calculated_batch_size + # 2. Setup device mesh - if config.batch_size == 1 and config.ici_tensor_parallelism == 1 and jax.device_count() > 1: + custom_parallelism_set = any( + any(arg.startswith(f"{k}=") for arg in sys.argv) + for k in [ + "ici_data_parallelism", + "ici_fsdp_parallelism", + "ici_context_parallelism", + "ici_tensor_parallelism", + ] + ) + + if not custom_parallelism_set and active_device_count > 1: max_logging.log( - f"ℹ️ Auto-configuring Tensor Parallelism: ici_tensor_parallelism={jax.device_count()}, ici_fsdp_parallelism=1 for batch_size=1 on {jax.device_count()} TPU devices." + f"ℹ️ Defaulting to Tensor Parallelism: ici_tensor_parallelism={active_device_count} on {active_device_count} TPU devices." ) - pyconfig._config.keys["ici_tensor_parallelism"] = jax.device_count() + pyconfig._config.keys["ici_tensor_parallelism"] = active_device_count + pyconfig._config.keys["ici_data_parallelism"] = 1 pyconfig._config.keys["ici_fsdp_parallelism"] = 1 + pyconfig._config.keys["ici_context_parallelism"] = 1 max_logging.log("Setting up JAX device mesh...") - devices_array = create_device_mesh(config) + devices_array = create_device_mesh(config, devices=active_devices) mesh = Mesh(devices_array, config.mesh_axes) # Check compatibility of batch dimension sharding @@ -174,8 +220,7 @@ def main(argv): # 3. Resolve weights repository snapshots repo_id = getattr(config, "pretrained_model_name_or_path", None) if not repo_id: - depth_val = getattr(config, "depth", None) - repo_id = "black-forest-labs/FLUX.2-klein-9B" if depth_val == 24 else "black-forest-labs/FLUX.2-klein-4B" + raise ValueError("pretrained_model_name_or_path must be specified in configuration YAML or CLI.") max_logging.log(f"Target model detected: {repo_id}") if os.path.exists(repo_id): @@ -184,8 +229,13 @@ def main(argv): else: from huggingface_hub import snapshot_download - max_logging.log(f"Resolving snapshot directory for model '{repo_id}' from HF Hub...") - snapshot_dir = snapshot_download(repo_id=repo_id) + rev = getattr(config, "revision", None) + if not rev or rev == "refs/pr/95": + rev = "main" + try: + snapshot_dir = snapshot_download(repo_id=repo_id, revision=rev, local_files_only=True) + except Exception: + snapshot_dir = snapshot_download(repo_id=repo_id, revision=rev) max_logging.log(f"Host {jax.process_index()} using HF snapshot directory: {snapshot_dir}") safetensors_path = os.path.join(snapshot_dir, "transformer") @@ -194,9 +244,23 @@ def main(argv): # 4. Load Qwen3 Config & Setup model layout from transformers import AutoConfig - - max_logging.log(f"Loading Qwen3 config from text_encoder path: {text_encoder_path}...") - pt_config = AutoConfig.from_pretrained(text_encoder_path, local_files_only=True) + from maxdiffusion.max_utils import get_flash_block_sizes + from flax import nnx + from maxdiffusion.models.flux.transformers.transformer_flux_flax import NNXFlux2KleinTransformer2DModel + from maxdiffusion.models.flux.util import load_and_convert_flux_klein_nnx_weights + + pt_config = AutoConfig.from_pretrained(text_encoder_path) + + te_bs = get_flash_block_sizes( + type( + "Config", + (), + { + "flash_block_sizes": getattr(config, "text_encoder_flash_block_sizes", {}) or {}, + "attention": getattr(config, "text_encoder_attention", "flash"), + }, + )() + ) qwen3_config = FlaxQwen3Config( vocab_size=pt_config.vocab_size, @@ -209,32 +273,35 @@ def main(argv): rms_norm_eps=pt_config.rms_norm_eps, rope_theta=pt_config.rope_theta, dtype=jnp.bfloat16 if config.weights_dtype == "bfloat16" else jnp.float32, + attention_kernel=getattr(config, "text_encoder_attention", "flash"), + flash_block_sizes=te_bs, + mesh=mesh, + ulysses_shards=getattr(config, "ulysses_shards", -1), + ulysses_attention_chunks=getattr(config, "ulysses_attention_chunks", 1), + max_layer_to_run=getattr(config, "text_encoder_max_layer", 27), + is_causal=getattr(config, "text_encoder_is_causal", True), ) qwen3_model = FlaxQwen3Model(qwen3_config) - # Load Transformer HF config.json directly for model architecture parameters - import json - - transformer_config_json = os.path.join(safetensors_path, "config.json") + # Load Transformer config for layer counts if present transformer_pt_cfg = {} + transformer_config_json = os.path.join(safetensors_path, "config.json") if os.path.exists(transformer_config_json): - with open(transformer_config_json, "r") as f: - transformer_pt_cfg = json.load(f) - - num_double_layers = getattr(config, "num_double_layers", -1) - if num_double_layers is None or num_double_layers <= 0: - num_double_layers = transformer_pt_cfg.get("num_layers", 5) + try: + import json - depth = getattr(config, "depth", -1) - if depth is None or depth <= 0: - depth = transformer_pt_cfg.get("num_single_layers", 20) + with open(transformer_config_json, "r") as f: + transformer_pt_cfg = json.load(f) + except Exception: + pass - num_attention_heads = getattr(config, "num_attention_heads", -1) - if num_attention_heads is None or num_attention_heads <= 0: - num_attention_heads = transformer_pt_cfg.get("num_attention_heads", 24) + num_double_layers = getattr(config, "num_double_layers", None) or transformer_pt_cfg.get("num_layers", 5) + depth = getattr(config, "depth", None) or transformer_pt_cfg.get("num_single_layers", 20) + num_attention_heads = getattr(config, "num_attention_heads", None) or transformer_pt_cfg.get("num_attention_heads", 24) - # 5. Instantiate JAX Flux2KleinTransformer2DModel - transformer = Flux2KleinTransformer2DModel( + # 5. Instantiate JAX NNXFlux2KleinTransformer2DModel + transformer = NNXFlux2KleinTransformer2DModel( + rngs=nnx.Rngs(0), in_channels=128, num_layers=num_double_layers, num_single_layers=depth, @@ -242,20 +309,20 @@ def main(argv): num_attention_heads=num_attention_heads, joint_attention_dim=3 * pt_config.hidden_size, pooled_projection_dim=768, + guidance_embeds=True, + axes_dim=(32, 32, 32, 32), + theta=2000.0, mlp_ratio=3.0, - qkv_bias=False, - joint_attention_bias=False, - x_embedder_bias=False, - proj_out_bias=False, - use_global_modulation=True, - use_swiglu=True, - axes_dims_rope=(32, 32, 32, 32), - theta=2000, + attention_kernel=config.attention, + flash_min_seq_length=512, + flash_block_sizes=get_flash_block_sizes(config), mesh=mesh, dtype=jnp.bfloat16 if config.weights_dtype == "bfloat16" else jnp.float32, weights_dtype=jnp.bfloat16 if config.weights_dtype == "bfloat16" else jnp.float32, - attention_kernel=config.attention, - scale_shift_order=getattr(config, "scale_shift_order", "shift_scale"), + scale_shift_order=getattr(config, "scale_shift_order", "scale_shift"), + ulysses_shards=getattr(config, "ulysses_shards", -1), + ulysses_attention_chunks=getattr(config, "ulysses_attention_chunks", 1), + use_base2_exp=getattr(config, "use_base2_exp", True), ) # 6. Instantiate JAX VAE @@ -277,36 +344,15 @@ def main(argv): # 7. Evaluate shapes & extract mesh shardings max_logging.log("Evaluating model shapes and shardings...") - h_packed = config.height // 16 - w_packed = config.width // 16 - seq_len_img = h_packed * w_packed seq_len_txt = config.max_sequence_length - - img_dummy = jnp.zeros((config.batch_size, seq_len_img, 128)) - img_ids_dummy = jnp.zeros((config.batch_size, seq_len_img, 4)) - txt_dummy = jnp.zeros((config.batch_size, seq_len_txt, 3 * pt_config.hidden_size)) - txt_ids_dummy = jnp.zeros((config.batch_size, seq_len_txt, 4)) - vec_dummy = jnp.zeros((config.batch_size, 768)) - t_vec_dummy = jnp.zeros((config.batch_size,)) - guidance_vec_dummy = jnp.zeros((config.batch_size,)) dummy_img = jnp.zeros((config.batch_size, 3, 512, 512)) dummy_ids = jnp.zeros((config.batch_size, seq_len_txt), dtype=jnp.int32) dummy_mask = jnp.zeros((config.batch_size, seq_len_txt), dtype=jnp.int32) key = jax.random.PRNGKey(0) - key, vae_key, qwen_key = jax.random.split(key, 3) - - def transformer_init_fn(): - return transformer.init( - key, - hidden_states=img_dummy, - img_ids=img_ids_dummy, - encoder_hidden_states=txt_dummy, - txt_ids=txt_ids_dummy, - pooled_projections=vec_dummy, - timestep=t_vec_dummy, - guidance=guidance_vec_dummy, - ) + vae_key, qwen_key = jax.random.split(key, 2) + + abstract_state = nnx.state(transformer, nnx.Param) def vae_init_fn(): return vae.init(vae_key, dummy_img) @@ -315,11 +361,10 @@ def qwen3_init_fn(): return qwen3_model.init(qwen_key, dummy_ids, dummy_mask) with mesh, nn_partitioning.axis_rules(config.logical_axis_rules): - abstract_transformer_vars = jax.eval_shape(transformer_init_fn) + logical_transformer_specs = nnx.get_partition_spec(abstract_state) abstract_vae_vars = jax.eval_shape(vae_init_fn) abstract_qwen3_vars = jax.eval_shape(qwen3_init_fn) - logical_transformer_specs = nn.get_partition_spec(abstract_transformer_vars) logical_vae_specs = nn.get_partition_spec(abstract_vae_vars) logical_qwen3_specs = nn.get_partition_spec(abstract_qwen3_vars) @@ -327,9 +372,9 @@ def qwen3_init_fn(): vae_mesh_shardings = nn.logical_to_mesh_sharding(logical_vae_specs, mesh, config.logical_axis_rules) qwen3_mesh_shardings = nn.logical_to_mesh_sharding(logical_qwen3_specs, mesh, config.logical_axis_rules) - transformer_shardings = flax.core.freeze(transformer_mesh_shardings["params"]) vae_shardings = flax.core.freeze(vae_mesh_shardings["params"]) qwen3_shardings = flax.core.freeze(qwen3_mesh_shardings["params"]) + transformer_shardings = transformer_mesh_shardings # 8. Load weights on Host CPU max_logging.log("Loading parameters on Host CPU...") @@ -342,11 +387,7 @@ def qwen3_init_fn(): def unbox_fn(x): return x.unbox() if isinstance(x, flax_spmd.LogicallyPartitioned) else x - params = jax.tree_util.tree_map( - unbox_fn, abstract_transformer_vars["params"], is_leaf=lambda k: isinstance(k, flax_spmd.LogicallyPartitioned) - ) - params = flax.core.unfreeze(params) - + t_sub0 = time.time() vae_params = jax.tree_util.tree_map( unbox_fn, abstract_vae_vars["params"], is_leaf=lambda k: isinstance(k, flax_spmd.LogicallyPartitioned) ) @@ -357,49 +398,43 @@ def unbox_fn(x): ) qwen3_params = flax.core.unfreeze(qwen3_params) - params = load_and_convert_flux_klein_weights(safetensors_path, params, num_double_layers, depth) - vae_params, vae_bn_mean, vae_bn_std = load_and_convert_vae_weights(vae_safetensors_path, vae_params) - qwen3_params = load_and_convert_qwen3_weights(text_encoder_path, qwen3_params, qwen3_config) + max_logging.log(f" -> [SUB-TIMING 1/3] PyTree unboxing template setup: {time.time() - t_sub0:.2f}s") + t_sub1 = time.time() - if config.weights_dtype == "bfloat16": - max_logging.log("Casting JAX parameters to bfloat16 in-place...") - cast_dict_to_bfloat16_inplace(params, exclude_keywords=("norm",)) - cast_dict_to_bfloat16_inplace(vae_params, exclude_keywords=("norm",)) - cast_dict_to_bfloat16_inplace(qwen3_params, exclude_keywords=("norm",)) - vae_bn_mean = vae_bn_mean.astype(jnp.bfloat16) - vae_bn_std = vae_bn_std.astype(jnp.bfloat16) + weight_dtype = jnp.bfloat16 if config.weights_dtype == "bfloat16" else jnp.float32 + + params = load_and_convert_flux_klein_nnx_weights( + safetensors_path, abstract_state, num_double_layers, depth, dtype=weight_dtype + ) + vae_params, vae_bn_mean, vae_bn_std = load_and_convert_vae_weights( + vae_safetensors_path, vae_params, dtype=weight_dtype + ) + qwen3_params = load_and_convert_qwen3_weights(text_encoder_path, qwen3_params, qwen3_config) + max_logging.log( + f" -> [SUB-TIMING 2/3] Safetensors loading & key mapping (in target dtype): {time.time() - t_sub1:.4f}s" + ) - params = flax.core.freeze(params) vae_params = flax.core.freeze(vae_params) qwen3_params = flax.core.freeze(qwen3_params) max_logging.log("\n" + "=" * 80) max_logging.log("🚀 Pinning all parameters to TPU HBM permanently...") max_logging.log("=" * 80 + "\n") + t_sub3 = time.time() max_logging.log("Putting params on TPU HBM...") with mesh, nn_partitioning.axis_rules(config.logical_axis_rules): - try: - params = jax.tree_util.tree_map(max_utils.device_put_replicated, params, transformer_shardings) - except Exception as err: - max_logging.log("\n❌ jax.device_put(params, transformer_shardings) FAILED!") - flat_p = flax.traverse_util.flatten_dict(params) - flat_s = flax.traverse_util.flatten_dict(transformer_shardings) - k_p = set(flat_p.keys()) - k_s = set(flat_s.keys()) - max_logging.log(f"Keys in sharding spec but missing in params: {k_s - k_p}") - max_logging.log(f"Keys in params but missing in sharding spec: {k_p - k_s}") - sys.stdout.flush() - raise err + params = jax.tree_util.tree_map(max_utils.device_put_replicated, params, transformer_shardings) max_logging.log("Putting vae_params on TPU HBM...") vae_params = jax.tree_util.tree_map(max_utils.device_put_replicated, vae_params, vae_shardings) max_logging.log("Putting qwen3_params on TPU HBM...") qwen3_params = jax.tree_util.tree_map(max_utils.device_put_replicated, qwen3_params, qwen3_shardings) + max_logging.log(f" -> [SUB-TIMING 3/3] TPU HBM device_put placement: {time.time() - t_sub3:.4f}s") max_logging.log("All parameters placed on TPU HBM successfully!") gc.collect() jax.effects_barrier() load_time = time.time() - t_load_start - max_logging.log(f" -> [TIMING] Total Model Loading & Device Placement: {load_time:.2f} seconds ⏱️\n") + max_logging.log(f" -> [TIMING] Total Model Loading & Device Placement: {load_time:.4f} seconds ⏱️\n") # 9. Setup FlowMatch Scheduler scheduler = FlaxFlowMatchScheduler( @@ -426,16 +461,47 @@ def unbox_fn(x): mesh=mesh, ) - active_prompts = partition_prompts(config.prompt, config.batch_size) + prompt_str = getattr(config, "prompt", None) + if not prompt_str: + raise ValueError("Prompt must be specified in the configuration YAML or passed via CLI prompt='...'") + active_prompts = partition_prompts(prompt_str, config.batch_size) + + # Parse reference image paths for multi-image editing if provided + images = None + image_paths = getattr(config, "image_paths", None) + if image_paths is not None: + if isinstance(image_paths, str) and image_paths.strip(): + import ast + + try: + image_paths = ast.literal_eval(image_paths) + except Exception: + image_paths = [p.strip() for p in image_paths.split(",") if p.strip()] + if isinstance(image_paths, (list, tuple)) and len(image_paths) > 0: + max_logging.log(f" -> Loading {len(image_paths)} reference image(s) for multi-image editing...") + images = [] + for p in image_paths: + try: + if not os.path.exists(p): + raise FileNotFoundError(f"Reference image file not found: {p}") + with Image.open(p) as img_raw: + img = img_raw.convert("RGB").resize((config.width, config.height), Image.Resampling.BICUBIC) + images.append(img) + except (UnidentifiedImageError, OSError, FileNotFoundError) as e: + max_logging.log(f"❌ Error loading reference image '{p}': {e}") + raise ValueError(f"Failed to load reference image '{p}': {e}") from e + except Exception as e: + max_logging.log(f"❌ Unexpected error loading reference image '{p}': {e}") + raise ValueError(f"Failed to load reference image '{p}': {e}") from e if getattr(config, "interactive", False): - print("\n" + "=" * 80) - print(" BATCHED INTERACTIVE GENERATION MODE ENABLED 🎮") - print("The model has been fully loaded and compiled on the TPU.") - print(f"Batch size: {config.batch_size} parallel images.") - print("Enter prompts separated by '||' (e.g. A cute cat || A red car)") - print("Type 'exit' to quit.") - print("=" * 80) + max_logging.log("\n" + "=" * 80) + max_logging.log(" BATCHED INTERACTIVE GENERATION MODE ENABLED 🎮") + max_logging.log("The model has been fully loaded and compiled on the TPU.") + max_logging.log(f"Batch size: {config.batch_size} parallel images.") + max_logging.log("Enter prompts separated by '||' (e.g. A cute cat || A red car)") + max_logging.log("Type 'exit' to quit.") + max_logging.log("=" * 80) image_idx = 1 while True: @@ -465,6 +531,7 @@ def unbox_fn(x): width=config.width, num_inference_steps=config.num_inference_steps, batch_size=config.batch_size, + images=images, use_latents=False, output_dir=config.output_dir, output_name=output_file, @@ -481,37 +548,24 @@ def unbox_fn(x): max_logging.log(f" -> Custom latents shape: {latents_to_use.shape} | sum: {latents_to_use.sum():.6f}") max_logging.log("\n" + "=" * 80) - max_logging.log("🚀 Running initial dry run (Warmup Pass) to compile XLA graphs...") + max_logging.log("🚀 Pre-compiling XLA graphs concurrently (AOT Compilation)...") max_logging.log("=" * 80) - _, warmup_trace = pipeline( - prompt=active_prompts, + aot_time = pipeline.compile_aot_async( params=params, vae_params=vae_params, qwen3_params=qwen3_params, vae_bn_mean=vae_bn_mean, vae_bn_std=vae_bn_std, - transformer_shardings=transformer_shardings, - vae_shardings=vae_shardings, - qwen3_shardings=qwen3_shardings, + batch_size=config.batch_size, height=config.height, width=config.width, - num_inference_steps=config.num_inference_steps, - batch_size=config.batch_size, - use_latents=use_latents_flag, - latents=latents_to_use, - output_dir=config.output_dir, - output_name="flux2klein_warmup.png", - ) - warmup_time = ( - warmup_trace.get("prompt_encoding", 0.0) - + warmup_trace.get("denoise_loop", 0.0) - + warmup_trace.get("vae_decode", 0.0) + images=images, ) max_logging.log("\n" + "=" * 80) - max_logging.log("⏱️ Running timed pass at full TPU speed...") + max_logging.log("🚀 Running initial dry run (Warmup Pass) to verify compiled graph execution...") max_logging.log("=" * 80) - _, main_trace = pipeline( + _, warmup_trace = pipeline( prompt=active_prompts, params=params, vae_params=vae_params, @@ -525,27 +579,141 @@ def unbox_fn(x): width=config.width, num_inference_steps=config.num_inference_steps, batch_size=config.batch_size, + images=images, use_latents=use_latents_flag, latents=latents_to_use, output_dir=config.output_dir, - output_name=config.output_name, + output_name="flux2klein_warmup.png", + warmup=True, ) - main_time = ( - main_trace.get("prompt_encoding", 0.0) + main_trace.get("denoise_loop", 0.0) + main_trace.get("vae_decode", 0.0) + warmup_time = ( + warmup_trace.get("vae_encode", 0.0) + + warmup_trace.get("prompt_encoding", 0.0) + + warmup_trace.get("denoise_loop", 0.0) + + warmup_trace.get("vae_decode", 0.0) ) + num_reps = int(getattr(config, "num_reps", 1)) + max_logging.log("\n" + "=" * 80) + max_logging.log(f"⏱️ Running timed pass at full TPU speed (num_reps={num_reps})...") + max_logging.log("=" * 80) + + main_traces = [] + main_times = [] + + for rep in range(num_reps): + rep_str = f" [Rep {rep+1}/{num_reps}]" if num_reps > 1 else "" + if rep > 0: + max_logging.log(f"⏱️ Running timed pass{rep_str}...") + + if max_utils.profiler_enabled(config) and rep == 0: + max_logging.log(f"🚀 XProf / JAX Profiler active! Capturing trace into: {config.tensorboard_dir}") + with max_utils.Profiler(config, session_name="flux2klein_inference"): + _, trace_i = pipeline( + prompt=active_prompts, + params=params, + vae_params=vae_params, + qwen3_params=qwen3_params, + vae_bn_mean=vae_bn_mean, + vae_bn_std=vae_bn_std, + transformer_shardings=transformer_shardings, + vae_shardings=vae_shardings, + qwen3_shardings=qwen3_shardings, + height=config.height, + width=config.width, + num_inference_steps=config.num_inference_steps, + batch_size=config.batch_size, + images=images, + use_latents=use_latents_flag, + latents=latents_to_use, + output_dir=config.output_dir, + output_name=f"rep_{rep+1}_{config.output_name}" if num_reps > 1 else config.output_name, + ) + else: + _, trace_i = pipeline( + prompt=active_prompts, + params=params, + vae_params=vae_params, + qwen3_params=qwen3_params, + vae_bn_mean=vae_bn_mean, + vae_bn_std=vae_bn_std, + transformer_shardings=transformer_shardings, + vae_shardings=vae_shardings, + qwen3_shardings=qwen3_shardings, + height=config.height, + width=config.width, + num_inference_steps=config.num_inference_steps, + batch_size=config.batch_size, + images=images, + use_latents=use_latents_flag, + latents=latents_to_use, + output_dir=config.output_dir, + output_name=f"rep_{rep+1}_{config.output_name}" if num_reps > 1 else config.output_name, + ) + + tot_time_i = trace_i.get( + "e2e_pipeline_total", + trace_i.get("vae_encode", 0.0) + + trace_i.get("prompt_encoding", 0.0) + + trace_i.get("denoise_loop", 0.0) + + trace_i.get("vae_decode", 0.0), + ) + main_traces.append(trace_i) + main_times.append(tot_time_i) + if num_reps > 1: + vae_enc_str = f" | VAE_Enc={trace_i.get('vae_encode', 0.0):.4f}s" if trace_i.get("vae_encode", 0.0) > 0 else "" + max_logging.log( + f" -> Rep {rep+1}/{num_reps} Completed: Total={tot_time_i:.4f}s{vae_enc_str} | Qwen3={trace_i.get('qwen3_encoding', 0.0):.4f}s | Denoise={trace_i.get('denoise_loop', 0.0):.4f}s | VAE_Dec={trace_i.get('vae_decode', 0.0):.4f}s" + ) + + avg_main_time = sum(main_times) / num_reps + avg_vae_encode = sum(tr.get("vae_encode", 0.0) for tr in main_traces) / num_reps + avg_vae_to_qwen3 = sum(tr.get("vae_encode_to_qwen3", 0.0) for tr in main_traces) / num_reps + avg_start_to_qwen3 = sum(tr.get("start_to_qwen3", 0.0) for tr in main_traces) / num_reps + avg_prompt_enc = sum(tr.get("qwen3_encoding", tr.get("prompt_encoding", 0.0)) for tr in main_traces) / num_reps + avg_qwen3_to_denoise = sum(tr.get("qwen3_to_denoise", 0.0) for tr in main_traces) / num_reps + avg_denoise = sum(tr.get("denoise_loop", 0.0) for tr in main_traces) / num_reps + avg_denoise_to_vae = sum(tr.get("denoise_to_vae", 0.0) for tr in main_traces) / num_reps + avg_vae_decode = sum(tr.get("vae_decode", 0.0) for tr in main_traces) / num_reps + avg_image_saving = sum(tr.get("image_saving", 0.0) for tr in main_traces) / num_reps + + total_cold_start = load_time + aot_time + warmup_time + max_logging.log("\n" + "=" * 80) - max_logging.log("📊 FLUX.2-KLEIN LATENCY & TIMING BREAKDOWN (PURE MODEL INFERENCE)") + max_logging.log("📊 FLUX.2-KLEIN COMPLETE LATENCY & TIMING BREAKDOWN") max_logging.log("=" * 80) - max_logging.log(f"1) Total Model Loading & Placement Time: {load_time:.2f} seconds ⏱️") - max_logging.log(f"2) Cold-Start / Warmup Pass (XLA Compilation): {warmup_time:.2f} seconds ⏱️") - max_logging.log(f" - Qwen3 Encoding: {warmup_trace.get('prompt_encoding', 0.0):.2f}s") - max_logging.log(f" - Flux Denoising: {warmup_trace.get('denoise_loop', 0.0):.2f}s") - max_logging.log(f" - VAE Decoding: {warmup_trace.get('vae_decode', 0.0):.2f}s") - max_logging.log(f"3) Main Warmed-Up Pass (Pure Model Inference): {main_time:.2f} seconds ⏱️") - max_logging.log(f" - Qwen3 Encoding: {main_trace.get('prompt_encoding', 0.0):.2f}s") - max_logging.log(f" - Flux Denoising: {main_trace.get('denoise_loop', 0.0):.2f}s") - max_logging.log(f" - VAE Decoding: {main_trace.get('vae_decode', 0.0):.2f}s") + max_logging.log(f"1) Model Loading & Placement Time: {load_time:.4f} seconds ⏱️") + max_logging.log(f"2) Concurrent AOT XLA Compilation Time: {aot_time:.4f} seconds ⚡") + max_logging.log(f"3) Warmup Pass Execution Time: {warmup_time:.4f} seconds ⏱️") + if warmup_trace.get("vae_encode", 0.0) > 0: + max_logging.log(f" - VAE Encoding: {warmup_trace.get('vae_encode', 0.0):.4f}s") + max_logging.log(f" - Qwen3 Encoding: {warmup_trace.get('prompt_encoding', 0.0):.4f}s") + max_logging.log(f" - Flux Denoising: {warmup_trace.get('denoise_loop', 0.0):.4f}s") + max_logging.log(f" - VAE Decoding: {warmup_trace.get('vae_decode', 0.0):.4f}s") + max_logging.log(f"👉 TOTAL COLD-START TIME (Loading + AOT + Warmup): {total_cold_start:.4f} seconds 🎯") + rep_label = f" (Average across {num_reps} reps)" if num_reps > 1 else "" + max_logging.log(f"4) Main Warmed-Up Pass (Pure Inference Latency){rep_label}: {avg_main_time:.4f} seconds ⏱️") + step_num = 1 + if avg_vae_encode > 0: + max_logging.log(f" - {step_num}. VAE Image Encoding: {avg_vae_encode*1000:.2f} ms ({avg_vae_encode:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. VAE -> Qwen3: {avg_vae_to_qwen3*1000:.2f} ms ({avg_vae_to_qwen3:.4f}s)") + step_num += 1 + else: + max_logging.log(f" - {step_num}. Start -> Qwen3: {avg_start_to_qwen3*1000:.2f} ms ({avg_start_to_qwen3:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. Qwen3 Encoding: {avg_prompt_enc*1000:.2f} ms ({avg_prompt_enc:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. Qwen3 -> Denoising: {avg_qwen3_to_denoise*1000:.2f} ms ({avg_qwen3_to_denoise:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. Flux Denoising Loop: {avg_denoise*1000:.2f} ms ({avg_denoise:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. Denoising -> VAE: {avg_denoise_to_vae*1000:.2f} ms ({avg_denoise_to_vae:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. VAE Decoding: {avg_vae_decode*1000:.2f} ms ({avg_vae_decode:.4f}s)") + step_num += 1 + max_logging.log(f" - {step_num}. Image Saving: {avg_image_saving*1000:.2f} ms ({avg_image_saving:.4f}s)") + max_logging.log(f" - 👉 TOTAL E2E PIPELINE: {avg_main_time*1000:.2f} ms ({avg_main_time:.4f}s)") max_logging.log("=" * 80) max_logging.log("\n=======================================================") diff --git a/src/maxdiffusion/max_utils.py b/src/maxdiffusion/max_utils.py index 37027c27d..ac44ce0a0 100644 --- a/src/maxdiffusion/max_utils.py +++ b/src/maxdiffusion/max_utils.py @@ -379,11 +379,17 @@ def walk_and_upload_blobs(config, output_dir): def device_put_replicated(x, sharding): - """ - Although the name indicates replication, this function can be used + """Although the name indicates replication, this function can be used + to also shard an array based on sharding. """ - return jax.make_array_from_callback(x.shape, sharding, lambda index: x[index]) + arr = getattr(x, "value", x) + shd = getattr(sharding, "value", sharding) + res = jax.make_array_from_callback(arr.shape, shd, lambda index: arr[index]) + if hasattr(x, "set_value"): + x.set_value(res) + return x + return res def fill_unspecified_mesh_axes(parallelism_vals, target_product, parallelism_type): @@ -764,18 +770,19 @@ def get_flash_block_sizes(config): f"block_kv_dq: {user_block_sizes.get('block_kv_dq')}," f"use_fused_bwd_kernel: {user_block_sizes.get('use_fused_bwd_kernel')}" ) + use_fused_bwd = True if attention_is_tokamax else bool(user_block_sizes.get("use_fused_bwd_kernel", False)) flash_block_sizes = splash_attention_kernel.BlockSizes( - block_q=user_block_sizes.get("block_q_dkv", user_block_sizes["block_kv"]) + block_q=user_block_sizes.get("block_q_dkv", user_block_sizes.get("block_kv")) if attention_is_tokamax - else user_block_sizes["block_q"], - block_kv_compute=user_block_sizes["block_kv_compute"], - block_kv=user_block_sizes["block_kv"], - block_q_dkv=user_block_sizes["block_q_dkv"], - block_kv_dkv=user_block_sizes["block_kv_dkv"], - block_kv_dkv_compute=user_block_sizes["block_kv_dkv_compute"], - block_q_dq=None if attention_is_tokamax else value_or_none(user_block_sizes, "block_q_dq"), - block_kv_dq=None if attention_is_tokamax else value_or_none(user_block_sizes, "block_kv_dq"), - use_fused_bwd_kernel=True if attention_is_tokamax else value_or_none(user_block_sizes, "use_fused_bwd_kernel"), + else user_block_sizes.get("block_q"), + block_kv_compute=user_block_sizes.get("block_kv_compute"), + block_kv=user_block_sizes.get("block_kv"), + block_q_dkv=user_block_sizes.get("block_q_dkv", user_block_sizes.get("block_q")), + block_kv_dkv=user_block_sizes.get("block_kv_dkv", user_block_sizes.get("block_kv")), + block_kv_dkv_compute=user_block_sizes.get("block_kv_dkv_compute", user_block_sizes.get("block_kv_compute")), + block_q_dq=None if use_fused_bwd else user_block_sizes.get("block_q_dq", user_block_sizes.get("block_q")), + block_kv_dq=None if use_fused_bwd else user_block_sizes.get("block_kv_dq", user_block_sizes.get("block_kv")), + use_fused_bwd_kernel=use_fused_bwd, ) return flash_block_sizes @@ -898,15 +905,23 @@ def initialize_jax_for_gpu(): def maybe_initialize_jax_distributed_system(raw_keys): - if raw_keys["skip_jax_distributed_system"]: + if raw_keys.get("skip_jax_distributed_system", False): max_logging.log("Skipping jax distributed system due to skip_jax_distributed_system=True flag.") return + from jax._src.xla_bridge import backends_are_initialized + + if backends_are_initialized(): + max_logging.log("XLA backends already initialized; skipping jax.distributed.initialize().") + return if is_gpu_backend(raw_keys): max_logging.log("Attempting to initialize the jax distributed system for GPU backend...") initialize_jax_for_gpu() max_logging.log("Jax distributed system initialized on GPU!") else: - jax.distributed.initialize() + try: + jax.distributed.initialize() + except Exception as e: + max_logging.log(f"Warning: jax.distributed.initialize() skipped or failed: {e}") def safe_getattr(obj: Any, name: str, default: Any) -> Any: diff --git a/src/maxdiffusion/models/attention_flax.py b/src/maxdiffusion/models/attention_flax.py index 8b84ae057..5a2754b6d 100644 --- a/src/maxdiffusion/models/attention_flax.py +++ b/src/maxdiffusion/models/attention_flax.py @@ -127,8 +127,7 @@ def _reshape_batch_dim_to_heads(tensor, heads): tensor = tensor.reshape(batch_size // head_size, head_size, seq_len, dim) tensor = jnp.transpose(tensor, (0, 2, 1, 3)) reshaped_tensor = tensor.reshape(batch_size // head_size, seq_len, dim * head_size) - axis_names = nn.logical_to_mesh_axes((BATCH, LENGTH, HEAD)) - return jax.lax.with_sharding_constraint(reshaped_tensor, axis_names) + return nn.with_logical_constraint(reshaped_tensor, (BATCH, LENGTH, HEAD)) def _reshape_heads_to_batch_dim(tensor, heads): @@ -141,8 +140,7 @@ def _reshape_heads_to_batch_dim(tensor, heads): else: batch_size, head_size, seq_len, head_dim = tensor.shape reshaped_tensor = tensor.reshape(batch_size * head_size, seq_len, head_dim) - axis_names = nn.logical_to_mesh_axes((BATCH, LENGTH, HEAD)) - return jax.lax.with_sharding_constraint(reshaped_tensor, axis_names) + return nn.with_logical_constraint(reshaped_tensor, (BATCH, LENGTH, HEAD)) def _reshape_heads_to_head_dim(tensor): @@ -151,8 +149,7 @@ def _reshape_heads_to_head_dim(tensor): b, h, s, d = tensor.shape tensor = jnp.transpose(tensor, axes=[0, 2, 1, 3]) reshaped_tensor = jnp.reshape(tensor, (b, -1, h * d)) - axis_names = nn.logical_to_mesh_axes((BATCH, LENGTH, HEAD)) - return jax.lax.with_sharding_constraint(reshaped_tensor, axis_names) + return nn.with_logical_constraint(reshaped_tensor, (BATCH, LENGTH, HEAD)) def _unflatten_heads(tensor, heads): @@ -579,6 +576,7 @@ def _tpu_flash_attention( attention_mask: jax.Array = None, use_base2_exp: bool = False, use_experimental_scheduler: bool = False, + is_causal: bool = False, ) -> jax.Array: """TPU Flash Attention""" @@ -656,8 +654,12 @@ def wrap_flash_attention(query, key, value, attention_mask): key, _, key_seq_len = _pad_data_for_flash(key, heads, block_kv) value, _, _ = _pad_data_for_flash(value, heads, block_kv) - mask = splash_attention_mask.FullMask(_shape=(query.shape[2], key.shape[2])) - multi_head_mask = splash_attention_mask.MultiHeadMask(masks=(mask,) * query.shape[1]) + if is_causal: + mask = splash_attention_mask.CausalMask((query.shape[2], key.shape[2])) + multi_head_mask = splash_attention_mask.MultiHeadMask(masks=(mask,) * query.shape[1]) + else: + mask = splash_attention_mask.FullMask(_shape=(query.shape[2], key.shape[2])) + multi_head_mask = splash_attention_mask.MultiHeadMask(masks=(mask,) * query.shape[1]) segment_ids_cls = ( tokamax_splash_base.SegmentIds if attention_kernel == "tokamax_ring" else splash_attention_kernel.SegmentIds @@ -674,9 +676,14 @@ def wrap_flash_attention(query, key, value, attention_mask): # make_splash_mha is wrapped around shardmap and seq and head is already # sharded based on in_specs, therefore setting head_shards=1 and q_seq_shards=1. if attention_kernel == "tokamax_flash": - mask = tokamax_splash_attention_mask.FullMask( - _shape=(query.shape[2], key.shape[2]), - ) + if is_causal: + mask = tokamax_splash_attention_mask.CausalMask( + (query.shape[2], key.shape[2]), + ) + else: + mask = tokamax_splash_attention_mask.FullMask( + _shape=(query.shape[2], key.shape[2]), + ) splash_kernel = tokamax_splash_attention_kernel.make_splash_mha( mask=mask, q_seq_shards=1, # the sizes of the axis is sharding over seq_len @@ -1571,10 +1578,9 @@ def _cudnn_flash_attention(query: Array, key: Array, value: Array, heads: int, m key = _reshape_data_for_cudnn_flash(key, heads) value = _reshape_data_for_cudnn_flash(value, heads) - axis_names = nn.logical_to_mesh_axes((BATCH, LENGTH, HEAD, D_KV)) - query = jax.lax.with_sharding_constraint(query, axis_names) - key = jax.lax.with_sharding_constraint(key, axis_names) - value = jax.lax.with_sharding_constraint(value, axis_names) + query = nn.with_logical_constraint(query, (BATCH, LENGTH, HEAD, D_KV)) + key = nn.with_logical_constraint(key, (BATCH, LENGTH, HEAD, D_KV)) + value = nn.with_logical_constraint(value, (BATCH, LENGTH, HEAD, D_KV)) out = dpa_layer(query, key, value, mask=None) return _reshape_data_from_cudnn_flash(out) @@ -1677,6 +1683,7 @@ def ulysses_ring_custom_fixed_m_kernel(q, k, v, context): use_base2_exp=context.get("use_base2_exp", True), use_experimental_scheduler=context.get("use_experimental_scheduler", False), use_fixed_m=True, + ulysses_attention_chunks=context.get("ulysses_attention_chunks", 1), ) @@ -1787,6 +1794,7 @@ def flash_kernel(q, k, v, context): attention_mask=context["attention_mask"], use_base2_exp=context["use_base2_exp"], use_experimental_scheduler=context["use_experimental_scheduler"], + is_causal=context.get("is_causal", False), ) @@ -1808,6 +1816,7 @@ def tokamax_flash_kernel(q, k, v, context): attention_mask=context["attention_mask"], use_base2_exp=context["use_base2_exp"], use_experimental_scheduler=context["use_experimental_scheduler"], + is_causal=context.get("is_causal", False), ) @@ -1829,6 +1838,7 @@ def tokamax_ring_kernel(q, k, v, context): attention_mask=context["attention_mask"], use_base2_exp=context["use_base2_exp"], use_experimental_scheduler=context["use_experimental_scheduler"], + is_causal=context.get("is_causal", False), ) @@ -1882,6 +1892,7 @@ def _apply_attention( use_experimental_scheduler: bool = False, ulysses_shards: int = -1, ulysses_attention_chunks: int = 1, + is_causal: bool = False, ): """Routes to different attention kernels using a module-level registry.""" @@ -1947,6 +1958,7 @@ def _apply_attention( "float32_qk_product": float32_qk_product, "use_memory_efficient_attention": use_memory_efficient_attention, "dpa_layer": dpa_layer, + "is_causal": is_causal, } # Module-level Registry lookup @@ -2280,6 +2292,7 @@ class AttentionOp(nn.Module): use_experimental_scheduler: bool = False ulysses_shards: int = -1 ulysses_attention_chunks: int = 1 + is_causal: bool = False def setup(self): self.dpa_layer = None @@ -2329,6 +2342,7 @@ def apply_attention(self, query: Array, key: Array, value: Array, attention_mask use_experimental_scheduler=self.use_experimental_scheduler, ulysses_shards=self.ulysses_shards, ulysses_attention_chunks=self.ulysses_attention_chunks, + is_causal=self.is_causal, ) @@ -2619,9 +2633,9 @@ def __call__( rngs: nnx.Rngs = None, cached_kv: Optional[Dict[str, Tuple[jax.Array, jax.Array]]] = None, ) -> jax.Array: - axis_names = nn.logical_to_mesh_axes((BATCH, LENGTH, HEAD)) - hidden_states = jax.lax.with_sharding_constraint(hidden_states, axis_names) - encoder_hidden_states = jax.lax.with_sharding_constraint(encoder_hidden_states, axis_names) + hidden_states = nn.with_logical_constraint(hidden_states, (BATCH, LENGTH, HEAD)) + if encoder_hidden_states is not None: + encoder_hidden_states = nn.with_logical_constraint(encoder_hidden_states, (BATCH, LENGTH, HEAD)) dtype = hidden_states.dtype is_self_attention = encoder_hidden_states is None if encoder_hidden_states is None: @@ -2854,6 +2868,8 @@ class FlaxFluxAttention(nn.Module): qkv_bias: bool = False use_base2_exp: bool = False use_experimental_scheduler: bool = False + ulysses_shards: int = -1 + ulysses_attention_chunks: int = 1 def setup(self): if self.attention_kernel in {"flash", "cudnn_flash_te"} and self.mesh is None: @@ -2875,6 +2891,8 @@ def setup(self): float32_qk_product=False, use_base2_exp=self.use_base2_exp, use_experimental_scheduler=self.use_experimental_scheduler, + ulysses_shards=self.ulysses_shards, + ulysses_attention_chunks=self.ulysses_attention_chunks, ) kernel_axes = ("embed", "heads") diff --git a/src/maxdiffusion/models/embeddings_flax.py b/src/maxdiffusion/models/embeddings_flax.py index 526ca7071..17ac2b07a 100644 --- a/src/maxdiffusion/models/embeddings_flax.py +++ b/src/maxdiffusion/models/embeddings_flax.py @@ -97,7 +97,7 @@ def __init__( in_features=in_channels, out_features=time_embed_dim, use_bias=sample_proj_bias, - dtype=jnp.float32, + dtype=dtype, param_dtype=weights_dtype, precision=precision, kernel_init=nnx.with_partitioning( @@ -126,7 +126,7 @@ def __init__( in_features=time_embed_dim, out_features=time_embed_dim_out, use_bias=sample_proj_bias, - dtype=jnp.float32, + dtype=dtype, param_dtype=weights_dtype, precision=precision, kernel_init=nnx.with_partitioning( @@ -355,7 +355,7 @@ def __init__( in_features=in_features, out_features=hidden_size, use_bias=True, - dtype=jnp.float32, + dtype=dtype, param_dtype=weights_dtype, precision=precision, kernel_init=nnx.with_partitioning( @@ -371,7 +371,7 @@ def __init__( in_features=hidden_size, out_features=out_features, use_bias=True, - dtype=jnp.float32, + dtype=dtype, param_dtype=weights_dtype, precision=precision, kernel_init=nnx.with_partitioning( @@ -615,7 +615,7 @@ def __init__( weights_dtype=weights_dtype, ) - if pooled_projection_dim > 0: + if pooled_projection_dim is not None and pooled_projection_dim > 0: self.pooled_embedder = NNXPixArtAlphaTextProjection( rngs=rngs, in_features=pooled_projection_dim, @@ -633,7 +633,7 @@ def __call__( pooled_projection: Optional[jax.Array] = None, ) -> jax.Array: timesteps_proj = self.time_proj(timestep) - dtype = pooled_projection.dtype if pooled_projection is not None else jnp.float32 + dtype = pooled_projection.dtype if pooled_projection is not None else self.dtype timestep_emb = self.timestep_embedder(timesteps_proj.astype(dtype)) if self.guidance_embeds and guidance is not None: @@ -643,7 +643,7 @@ def __call__( else: time_guidance_emb = timestep_emb - if pooled_projection is not None and self.pooled_projection_dim > 0: + if pooled_projection is not None and self.pooled_projection_dim is not None and self.pooled_projection_dim > 0: pooled_projections = self.pooled_embedder(pooled_projection) conditioning = time_guidance_emb + pooled_projections else: diff --git a/src/maxdiffusion/models/flux/transformers/transformer_flux_flax.py b/src/maxdiffusion/models/flux/transformers/transformer_flux_flax.py index af8e3763a..3183f462c 100644 --- a/src/maxdiffusion/models/flux/transformers/transformer_flux_flax.py +++ b/src/maxdiffusion/models/flux/transformers/transformer_flux_flax.py @@ -14,7 +14,7 @@ limitations under the License. """ -from typing import Dict, Optional, Tuple +from typing import Dict, Optional, Tuple, Union import jax import math import jax.numpy as jnp @@ -27,11 +27,9 @@ AdaLayerNormZeroSingle, AdaLayerNormContinuous, AdaLayerNormZero, - NNXAdaLayerNormZeroSingle, NNXAdaLayerNormContinuous, - NNXAdaLayerNormZero, ) -from ...attention_flax import FlaxFluxAttention as FluxAttention, FlaxFluxAttention, apply_rope +from ...attention_flax import FlaxFluxAttention as FluxAttention, FlaxFluxAttention, apply_rope, NNXAttentionOp from flax import nnx from ...embeddings_flax import ( FluxPosEmbed, @@ -1278,64 +1276,123 @@ def __call__( return Transformer2DModelOutput(sample=output) -# ============================================================================= -# FLAX NNX MODEL IMPLEMENTATIONS FOR FLUX.2-KLEIN -# ============================================================================= +class NNXFlaxSwiGluFeedForward(nnx.Module): + """Flax NNX SwiGLU FeedForward module.""" + + def __init__( + self, + rngs: nnx.Rngs, + dim: int, + dim_out: int, + mult: float = 3.0, + dtype: jnp.dtype = jnp.float32, + weights_dtype: jnp.dtype = jnp.float32, + ): + inner_dim = int(dim * mult) + self.linear_in = nnx.Linear( + in_features=dim, + out_features=inner_dim * 2, + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "mlp")), + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) + self.linear_out = nnx.Linear( + in_features=inner_dim, + out_features=dim_out, + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("mlp", "embed")), + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + x = self.linear_in(x) + x1, x2 = jnp.split(x, 2, axis=-1) + hidden = nnx.silu(x1) * x2 + return self.linear_out(hidden) -class NNXFluxDoubleAttention(nnx.Module): +class NNXFluxAttention(nnx.Module): + """Flax NNX Double-Stream Joint Attention for FLUX.2-Klein.""" def __init__( self, rngs: nnx.Rngs, query_dim: int, - heads: int, - dim_head: int, - qkv_bias: bool = False, + heads: int = 8, + dim_head: int = 64, + attention_kernel: str = "dot_product", + flash_min_seq_length: int = 512, + flash_block_sizes: Optional[Dict[str, int]] = None, + mesh: Optional[jax.sharding.Mesh] = None, dtype: jnp.dtype = jnp.float32, weights_dtype: jnp.dtype = jnp.float32, + qkv_bias: bool = False, + ulysses_shards: int = -1, + ulysses_attention_chunks: int = 1, + use_base2_exp: bool = False, ): - self.query_dim = query_dim self.heads = heads self.dim_head = dim_head - inner_dim = heads * dim_head + inner_dim = dim_head * heads + scale = dim_head**-0.5 + + self.attention_op = NNXAttentionOp( + mesh=mesh, + attention_kernel=attention_kernel, + scale=scale, + heads=heads, + dim_head=dim_head, + flash_min_seq_length=flash_min_seq_length, + flash_block_sizes=flash_block_sizes, + dtype=dtype, + float32_qk_product=False, + split_head_dim=False, + ulysses_shards=ulysses_shards, + ulysses_attention_chunks=ulysses_attention_chunks, + use_base2_exp=use_base2_exp, + ) + + kernel_axes = ("embed", "heads") + proj_attn_kernel_axes = ("heads", "embed") - self.qkv = nnx.Linear( + self.i_qkv = nnx.Linear( in_features=query_dim, out_features=inner_dim * 3, use_bias=qkv_bias, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "heads")), + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), kernel_axes), bias_init=nnx.with_partitioning(nnx.initializers.zeros, ("heads",)), dtype=dtype, param_dtype=weights_dtype, rngs=rngs, ) - self.encoder_qkv = nnx.Linear( + self.e_qkv = nnx.Linear( in_features=query_dim, out_features=inner_dim * 3, use_bias=qkv_bias, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "heads")), + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), kernel_axes), bias_init=nnx.with_partitioning(nnx.initializers.zeros, ("heads",)), dtype=dtype, param_dtype=weights_dtype, rngs=rngs, ) - self.proj_attn = nnx.Linear( + self.i_proj = nnx.Linear( in_features=inner_dim, out_features=query_dim, - use_bias=True, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("heads", "embed")), - bias_init=nnx.with_partitioning(nnx.initializers.zeros, ("embed",)), + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), proj_attn_kernel_axes), dtype=dtype, param_dtype=weights_dtype, rngs=rngs, ) - self.encoder_proj_attn = nnx.Linear( + self.e_proj = nnx.Linear( in_features=inner_dim, out_features=query_dim, - use_bias=True, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("heads", "embed")), - bias_init=nnx.with_partitioning(nnx.initializers.zeros, ("embed",)), + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), proj_attn_kernel_axes), dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1356,59 +1413,84 @@ def __init__( param_dtype=weights_dtype, rngs=rngs, ) + self.encoder_query_norm = nnx.RMSNorm( + num_features=dim_head, + epsilon=1e-6, + scale_init=nnx.with_partitioning(nnx.initializers.ones, ("heads",)), + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) + self.encoder_key_norm = nnx.RMSNorm( + num_features=dim_head, + epsilon=1e-6, + scale_init=nnx.with_partitioning(nnx.initializers.ones, ("heads",)), + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) def __call__( self, hidden_states: jax.Array, - encoder_hidden_states: jax.Array, - image_rotary_emb: Tuple[jax.Array, jax.Array], - ) -> Tuple[jax.Array, jax.Array]: - batch_size, img_len, _ = hidden_states.shape - txt_len = encoder_hidden_states.shape[1] - - qkv_img = self.qkv(hidden_states) - qkv_txt = self.encoder_qkv(encoder_hidden_states) - - q_img, k_img, v_img = jnp.split(qkv_img, 3, axis=-1) - q_txt, k_txt, v_txt = jnp.split(qkv_txt, 3, axis=-1) - - q_img = rearrange(q_img, "b l (h d) -> b l h d", h=self.heads) - k_img = rearrange(k_img, "b l (h d) -> b l h d", h=self.heads) - v_img = rearrange(v_img, "b l (h d) -> b l h d", h=self.heads) + encoder_hidden_states: Optional[jax.Array] = None, + image_rotary_emb: Optional[Tuple[jax.Array, jax.Array]] = None, + ) -> Tuple[jax.Array, Optional[jax.Array]]: + B, L = hidden_states.shape[:2] + H, D = self.heads, self.dim_head - q_txt = rearrange(q_txt, "b l (h d) -> b l h d", h=self.heads) - k_txt = rearrange(k_txt, "b l (h d) -> b l h d", h=self.heads) - v_txt = rearrange(v_txt, "b l (h d) -> b l h d", h=self.heads) + qkv_proj = self.i_qkv(hidden_states).reshape(B, L, 3, H, D) + query_proj, key_proj, value_proj = jnp.split(qkv_proj, 3, axis=2) + query_proj = self.query_norm(query_proj.squeeze(2)) + key_proj = self.key_norm(key_proj.squeeze(2)) + value_proj = value_proj.squeeze(2) - q_img = self.query_norm(q_img) - k_img = self.key_norm(k_img) - q_txt = self.query_norm(q_txt) - k_txt = self.key_norm(k_txt) + if encoder_hidden_states is not None: + B_enc, L_txt = encoder_hidden_states.shape[:2] + encoder_qkv_proj = self.e_qkv(encoder_hidden_states).reshape(B_enc, L_txt, 3, H, D) + enc_query_proj, enc_key_proj, enc_value_proj = jnp.split(encoder_qkv_proj, 3, axis=2) + enc_query_proj = self.encoder_query_norm(enc_query_proj.squeeze(2)) + enc_key_proj = self.encoder_key_norm(enc_key_proj.squeeze(2)) + enc_value_proj = enc_value_proj.squeeze(2) - q = jnp.concatenate([q_txt, q_img], axis=1) - k = jnp.concatenate([k_txt, k_img], axis=1) - v = jnp.concatenate([v_txt, v_img], axis=1) + query_proj = jnp.concatenate((enc_query_proj, query_proj), axis=1) + key_proj = jnp.concatenate((enc_key_proj, key_proj), axis=1) + value_proj = jnp.concatenate((enc_value_proj, value_proj), axis=1) if image_rotary_emb is not None: - q, k = apply_rope(q, k, image_rotary_emb) + if not isinstance(image_rotary_emb, (tuple, list)): + image_rotary_emb_reordered = rearrange(image_rotary_emb, "n d (i j) -> n d i j", i=2, j=2) + else: + image_rotary_emb_reordered = image_rotary_emb + query_proj = query_proj.swapaxes(1, 2) + key_proj = key_proj.swapaxes(1, 2) + query_proj, key_proj = apply_rope(query_proj, key_proj, image_rotary_emb_reordered) + query_proj = query_proj.swapaxes(1, 2) + key_proj = key_proj.swapaxes(1, 2) - scale = self.dim_head**-0.5 - attn_weights = jnp.einsum("b q h d, b k h d -> b h q k", q, k, precision=None) * scale - attn_weights = jax.nn.softmax(attn_weights, axis=-1) - out = jnp.einsum("b h q k, b k h d -> b q h d", attn_weights, v, precision=None) + query_proj = query_proj.reshape(B, -1, H * D) + key_proj = key_proj.reshape(B, -1, H * D) + value_proj = value_proj.reshape(B, -1, H * D) - out = rearrange(out, "b l h d -> b l (h d)") + if encoder_hidden_states is not None: + query_proj = nn.with_logical_constraint(query_proj, ("activation_batch", "activation_length", "activation_heads")) + key_proj = nn.with_logical_constraint(key_proj, ("activation_batch", "activation_length", "activation_heads")) + value_proj = nn.with_logical_constraint(value_proj, ("activation_batch", "activation_length", "activation_heads")) - out_txt = out[:, :txt_len, :] - out_img = out[:, txt_len:, :] + attn_output = self.attention_op.apply_attention(query_proj, key_proj, value_proj) + context_attn_output = None - out_img = self.proj_attn(out_img) - out_txt = self.encoder_proj_attn(out_txt) + if encoder_hidden_states is not None: + context_attn_output = attn_output[:, : encoder_hidden_states.shape[1]] + attn_output = attn_output[:, encoder_hidden_states.shape[1] :] + attn_output = self.i_proj(attn_output) + context_attn_output = self.e_proj(context_attn_output) - return out_img, out_txt + return attn_output, context_attn_output class NNXFluxSingleAttention(nnx.Module): + """Flax NNX Single-Stream Attention for FLUX.2-Klein.""" def __init__( self, @@ -1416,33 +1498,36 @@ def __init__( dim: int, num_attention_heads: int, attention_head_dim: int, + attention_kernel: str = "dot_product", + flash_min_seq_length: int = 512, + flash_block_sizes: Optional[Dict[str, int]] = None, + mesh: Optional[jax.sharding.Mesh] = None, dtype: jnp.dtype = jnp.float32, weights_dtype: jnp.dtype = jnp.float32, + ulysses_shards: int = -1, + ulysses_attention_chunks: int = 1, + use_base2_exp: bool = False, ): - self.dim = dim - self.heads = num_attention_heads - self.dim_head = attention_head_dim - inner_dim = num_attention_heads * attention_head_dim + self.num_attention_heads = num_attention_heads + self.attention_head_dim = attention_head_dim + scale = attention_head_dim**-0.5 - self.to_qkv_mlp_proj = nnx.Linear( - in_features=dim, - out_features=inner_dim * 3 + int(dim * 4.0), - use_bias=False, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "mlp")), - dtype=dtype, - param_dtype=weights_dtype, - rngs=rngs, - ) - self.to_out = nnx.Linear( - in_features=inner_dim + int(dim * 4.0), - out_features=dim, - use_bias=False, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("mlp", "embed")), + self.attention_op = NNXAttentionOp( + mesh=mesh, + attention_kernel=attention_kernel, + scale=scale, + heads=num_attention_heads, + dim_head=attention_head_dim, + flash_min_seq_length=flash_min_seq_length, + flash_block_sizes=flash_block_sizes, dtype=dtype, - param_dtype=weights_dtype, - rngs=rngs, + float32_qk_product=False, + split_head_dim=False, + ulysses_shards=ulysses_shards, + ulysses_attention_chunks=ulysses_attention_chunks, + use_base2_exp=use_base2_exp, ) - self.norm_q = nnx.RMSNorm( + self.query_norm = nnx.RMSNorm( num_features=attention_head_dim, epsilon=1e-6, scale_init=nnx.with_partitioning(nnx.initializers.ones, ("heads",)), @@ -1450,7 +1535,7 @@ def __init__( param_dtype=weights_dtype, rngs=rngs, ) - self.norm_k = nnx.RMSNorm( + self.key_norm = nnx.RMSNorm( num_features=attention_head_dim, epsilon=1e-6, scale_init=nnx.with_partitioning(nnx.initializers.ones, ("heads",)), @@ -1459,42 +1544,9 @@ def __init__( rngs=rngs, ) - def __call__( - self, - hidden_states: jax.Array, - image_rotary_emb: Tuple[jax.Array, jax.Array], - ) -> jax.Array: - batch_size, seq_len, _ = hidden_states.shape - inner_dim = self.heads * self.dim_head - - qkv_mlp = self.to_qkv_mlp_proj(hidden_states) - qkv, mlp = jnp.split(qkv_mlp, [inner_dim * 3], axis=-1) - - q, k, v = jnp.split(qkv, 3, axis=-1) - q = rearrange(q, "b l (h d) -> b l h d", h=self.heads) - k = rearrange(k, "b l (h d) -> b l h d", h=self.heads) - v = rearrange(v, "b l (h d) -> b l h d", h=self.heads) - - q = self.norm_q(q) - k = self.norm_k(k) - - if image_rotary_emb is not None: - q, k = apply_rope(q, k, image_rotary_emb) - - scale = self.dim_head**-0.5 - attn_weights = jnp.einsum("b q h d, b k h d -> b h q k", q, k, precision=None) * scale - attn_weights = jax.nn.softmax(attn_weights, axis=-1) - attn_out = jnp.einsum("b h q k, b k h d -> b q h d", attn_weights, v, precision=None) - attn_out = rearrange(attn_out, "b l h d -> b l (h d)") - - mlp_act = jax.nn.gelu(mlp, approximate=True) - attn_mlp = jnp.concatenate([attn_out, mlp_act], axis=-1) - - out = self.to_out(attn_mlp) - return out - class NNXFluxDoubleTransformerBlock(nnx.Module): + """Flax NNX Double-Stream Transformer Block for FLUX.2-Klein.""" def __init__( self, @@ -1502,66 +1554,91 @@ def __init__( dim: int, num_attention_heads: int, attention_head_dim: int, - mlp_ratio: float = 4.0, + mlp_ratio: float = 3.0, + attention_kernel: str = "dot_product", + flash_min_seq_length: int = 512, + flash_block_sizes: Optional[Dict[str, int]] = None, + mesh: Optional[jax.sharding.Mesh] = None, dtype: jnp.dtype = jnp.float32, weights_dtype: jnp.dtype = jnp.float32, + qkv_bias: bool = False, + ulysses_shards: int = -1, + ulysses_attention_chunks: int = 1, + use_base2_exp: bool = False, ): self.dim = dim self.num_heads = num_attention_heads self.head_dim = attention_head_dim - mlp_hidden_dim = int(dim * mlp_ratio) - - self.img_norm1 = NNXAdaLayerNormZero(dim, dtype=dtype, weights_dtype=weights_dtype) - self.txt_norm1 = NNXAdaLayerNormZero(dim, dtype=dtype, weights_dtype=weights_dtype) - self.attn = NNXFluxDoubleAttention( - rngs=rngs, - query_dim=dim, - heads=num_attention_heads, - dim_head=attention_head_dim, + self.norm1 = nnx.LayerNorm( + num_features=dim, + use_bias=False, + use_scale=False, + epsilon=1e-6, dtype=dtype, - weights_dtype=weights_dtype, + param_dtype=weights_dtype, + rngs=rngs, ) - - self.img_mlp = nnx.Linear( - in_features=dim, - out_features=mlp_hidden_dim, - use_bias=True, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "mlp")), - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + self.norm1_context = nnx.LayerNorm( + num_features=dim, + use_bias=False, + use_scale=False, + epsilon=1e-6, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, ) - self.img_mlp_out = nnx.Linear( - in_features=mlp_hidden_dim, - out_features=dim, - use_bias=True, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("mlp", "embed")), - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + self.norm2 = nnx.LayerNorm( + num_features=dim, + use_bias=False, + use_scale=False, + epsilon=1e-6, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, ) - self.txt_mlp = nnx.Linear( - in_features=dim, - out_features=mlp_hidden_dim, - use_bias=True, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "mlp")), - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + self.norm2_context = nnx.LayerNorm( + num_features=dim, + use_bias=False, + use_scale=False, + epsilon=1e-6, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, ) - self.txt_mlp_out = nnx.Linear( - in_features=mlp_hidden_dim, - out_features=dim, - use_bias=True, - kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("mlp", "embed")), - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + + self.attn = NNXFluxAttention( + rngs=rngs, + query_dim=dim, + heads=num_attention_heads, + dim_head=attention_head_dim, + attention_kernel=attention_kernel, + flash_min_seq_length=flash_min_seq_length, + flash_block_sizes=flash_block_sizes, + mesh=mesh, dtype=dtype, - param_dtype=weights_dtype, + weights_dtype=weights_dtype, + qkv_bias=qkv_bias, + ulysses_shards=ulysses_shards, + ulysses_attention_chunks=ulysses_attention_chunks, + use_base2_exp=use_base2_exp, + ) + + self.ff = NNXFlaxSwiGluFeedForward( rngs=rngs, + dim=dim, + dim_out=dim, + mult=mlp_ratio, + dtype=dtype, + weights_dtype=weights_dtype, + ) + self.ff_context = NNXFlaxSwiGluFeedForward( + rngs=rngs, + dim=dim, + dim_out=dim, + mult=mlp_ratio, + dtype=dtype, + weights_dtype=weights_dtype, ) def __call__( @@ -1573,33 +1650,49 @@ def __call__( temb_mod_img: Optional[jax.Array] = None, temb_mod_txt: Optional[jax.Array] = None, ) -> Tuple[jax.Array, jax.Array]: - norm_h, c_gate_msa, c_shift_mlp, c_scale_mlp, c_gate_mlp = self.img_norm1(hidden_states, emb=temb_mod_img) - norm_enc, c_gate_msa_txt, c_shift_mlp_txt, c_scale_mlp_txt, c_gate_mlp_txt = self.txt_norm1( - encoder_hidden_states, emb=temb_mod_txt - ) + shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = jnp.split(temb_mod_img, 6, axis=-1) + c_shift_msa, c_scale_msa, c_gate_msa, c_shift_mlp, c_scale_mlp, c_gate_mlp = jnp.split(temb_mod_txt, 6, axis=-1) + + shift_msa = jnp.expand_dims(shift_msa, axis=1) + scale_msa = jnp.expand_dims(scale_msa, axis=1) + gate_msa = jnp.expand_dims(gate_msa, axis=1) + shift_mlp = jnp.expand_dims(shift_mlp, axis=1) + scale_mlp = jnp.expand_dims(scale_mlp, axis=1) + gate_mlp = jnp.expand_dims(gate_mlp, axis=1) + + c_shift_msa = jnp.expand_dims(c_shift_msa, axis=1) + c_scale_msa = jnp.expand_dims(c_scale_msa, axis=1) + c_gate_msa = jnp.expand_dims(c_gate_msa, axis=1) + c_shift_mlp = jnp.expand_dims(c_shift_mlp, axis=1) + c_scale_mlp = jnp.expand_dims(c_scale_mlp, axis=1) + c_gate_mlp = jnp.expand_dims(c_gate_mlp, axis=1) + + norm1_h = self.norm1(hidden_states) * (1.0 + scale_msa) + shift_msa + norm1_enc = self.norm1_context(encoder_hidden_states) * (1.0 + c_scale_msa) + c_shift_msa attn_img, attn_txt = self.attn( - hidden_states=norm_h, - encoder_hidden_states=norm_enc, + hidden_states=norm1_h, + encoder_hidden_states=norm1_enc, image_rotary_emb=image_rotary_emb, ) - hidden_states = hidden_states + c_gate_msa * attn_img - encoder_hidden_states = encoder_hidden_states + c_gate_msa_txt * attn_txt + hidden_states = hidden_states + gate_msa * attn_img + encoder_hidden_states = encoder_hidden_states + c_gate_msa * attn_txt - norm_h_mlp = norm_h * (1.0 + c_scale_mlp) + c_shift_mlp - norm_enc_mlp = norm_enc * (1.0 + c_scale_mlp_txt) + c_shift_mlp_txt + norm2_h = self.norm2(hidden_states) * (1.0 + scale_mlp) + shift_mlp + norm2_enc = self.norm2_context(encoder_hidden_states) * (1.0 + c_scale_mlp) + c_shift_mlp - img_ff = self.img_mlp_out(jax.nn.gelu(self.img_mlp(norm_h_mlp), approximate=True)) - txt_ff = self.txt_mlp_out(jax.nn.gelu(self.txt_mlp(norm_enc_mlp), approximate=True)) + mlp_output = self.ff(norm2_h) + encoder_mlp_output = self.ff_context(norm2_enc) - hidden_states = hidden_states + c_gate_mlp * img_ff - encoder_hidden_states = encoder_hidden_states + c_gate_mlp_txt * txt_ff + hidden_states = hidden_states + gate_mlp * mlp_output + encoder_hidden_states = encoder_hidden_states + c_gate_mlp * encoder_mlp_output return encoder_hidden_states, hidden_states class NNXFluxSingleTransformerBlock(nnx.Module): + """Flax NNX Single-Stream Transformer Block for FLUX.2-Klein.""" def __init__( self, @@ -1607,18 +1700,67 @@ def __init__( dim: int, num_attention_heads: int, attention_head_dim: int, + mlp_ratio: float = 3.0, + attention_kernel: str = "dot_product", + flash_min_seq_length: int = 512, + flash_block_sizes: Optional[Dict[str, int]] = None, + mesh: Optional[jax.sharding.Mesh] = None, dtype: jnp.dtype = jnp.float32, weights_dtype: jnp.dtype = jnp.float32, + ulysses_shards: int = -1, + ulysses_attention_chunks: int = 1, + use_base2_exp: bool = False, ): self.dim = dim - self.norm = NNXAdaLayerNormZeroSingle(dim, dtype=dtype, weights_dtype=weights_dtype) + self.num_attention_heads = num_attention_heads + self.attention_head_dim = attention_head_dim + mlp_hidden_dim = int(dim * mlp_ratio) + + self.norm = nnx.LayerNorm( + num_features=dim, + use_bias=False, + use_scale=False, + epsilon=1e-6, + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) + + out_dim = dim * 3 + 2 * mlp_hidden_dim + self.linear1 = nnx.Linear( + in_features=dim, + out_features=out_dim, + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "mlp")), + bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) + self.linear2 = nnx.Linear( + in_features=dim + mlp_hidden_dim, + out_features=dim, + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("mlp", "embed")), + bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + dtype=dtype, + param_dtype=weights_dtype, + rngs=rngs, + ) self.attn = NNXFluxSingleAttention( rngs=rngs, dim=dim, num_attention_heads=num_attention_heads, attention_head_dim=attention_head_dim, + attention_kernel=attention_kernel, + flash_min_seq_length=flash_min_seq_length, + flash_block_sizes=flash_block_sizes, + mesh=mesh, dtype=dtype, weights_dtype=weights_dtype, + ulysses_shards=ulysses_shards, + ulysses_attention_chunks=ulysses_attention_chunks, + use_base2_exp=use_base2_exp, ) def __call__( @@ -1628,22 +1770,59 @@ def __call__( image_rotary_emb: Tuple[jax.Array, jax.Array], temb_mod: Optional[jax.Array] = None, ) -> jax.Array: - norm_hidden_states, gate_msa = self.norm(hidden_states, emb=temb_mod) - attn_output = self.attn( - hidden_states=norm_hidden_states, - image_rotary_emb=image_rotary_emb, - ) - hidden_states = hidden_states + gate_msa * attn_output + residual = hidden_states + shift_msa, scale_msa, gate = jnp.split(temb_mod, 3, axis=-1) + shift_msa = jnp.expand_dims(shift_msa, axis=1) + scale_msa = jnp.expand_dims(scale_msa, axis=1) + gate = jnp.expand_dims(gate, axis=1) + + norm_hidden_states = self.norm(hidden_states) + norm_hidden_states = (1 + scale_msa) * norm_hidden_states + shift_msa + + qkv, mlp = jnp.split(self.linear1(norm_hidden_states), [3 * self.dim], axis=-1) + qkv = nn.with_logical_constraint(qkv, ("activation_batch", "activation_length", "activation_embed")) + mlp = nn.with_logical_constraint(mlp, ("activation_batch", "activation_length", "activation_embed")) + + B, L = hidden_states.shape[:2] + H, D = self.num_attention_heads, qkv.shape[-1] // (self.num_attention_heads * 3) + qkv_proj = qkv.reshape(B, L, 3, H, D).transpose(2, 0, 3, 1, 4) + q, k, v = qkv_proj + + q = self.attn.query_norm(q) + k = self.attn.key_norm(k) + + if image_rotary_emb is not None: + if isinstance(image_rotary_emb, (tuple, list)): + image_rotary_emb_reordered = image_rotary_emb + else: + image_rotary_emb_reordered = rearrange(image_rotary_emb, "n d (i j) -> n d i j", i=2, j=2) + q, k = apply_rope(q, k, image_rotary_emb_reordered) + + q = q.transpose(0, 2, 1, 3).reshape(q.shape[0], q.shape[2], -1) + k = k.transpose(0, 2, 1, 3).reshape(k.shape[0], k.shape[2], -1) + v = v.transpose(0, 2, 1, 3).reshape(v.shape[0], v.shape[2], -1) + + attn_output = self.attn.attention_op.apply_attention(q, k, v) + + mlp1, mlp2 = jnp.split(mlp, 2, axis=-1) + mlp_activated = nnx.silu(mlp1) * mlp2 + + attn_mlp = jnp.concatenate([attn_output, mlp_activated], axis=2) + attn_mlp = nn.with_logical_constraint(attn_mlp, ("activation_batch", "activation_length", "activation_embed")) + hidden_states = self.linear2(attn_mlp) + hidden_states = gate * hidden_states + hidden_states = residual + hidden_states return hidden_states -class NNXFluxTransformer2DModel(nnx.Module): +class NNXFlux2KleinTransformer2DModel(nnx.Module): + """Flax NNX Top-Level FLUX.2-Klein Transformer 2D Model.""" def __init__( self, rngs: nnx.Rngs, patch_size: int = 1, - in_channels: int = 64, + in_channels: int = 128, num_layers: int = 5, num_single_layers: int = 20, attention_head_dim: int = 128, @@ -1651,10 +1830,19 @@ def __init__( joint_attention_dim: int = 4096, pooled_projection_dim: int = 768, guidance_embeds: bool = True, - axes_dim: Tuple[int, ...] = (16, 56, 56), - theta: float = 10000.0, + axes_dim: Tuple[int, ...] = (32, 32, 32, 32), + theta: float = 2000.0, + mlp_ratio: float = 3.0, + attention_kernel: str = "dot_product", + flash_min_seq_length: int = 512, + flash_block_sizes: Optional[Dict[str, int]] = None, + mesh: Optional[jax.sharding.Mesh] = None, dtype: jnp.dtype = jnp.float32, weights_dtype: jnp.dtype = jnp.float32, + scale_shift_order: str = "scale_shift", + ulysses_shards: int = -1, + ulysses_attention_chunks: int = 1, + use_base2_exp: bool = False, ): self.in_channels = in_channels self.out_channels = in_channels @@ -1663,6 +1851,7 @@ def __init__( self.num_single_layers = num_single_layers self.attention_head_dim = attention_head_dim self.num_attention_heads = num_attention_heads + self.joint_attention_dim = joint_attention_dim self.inner_dim = num_attention_heads * attention_head_dim self.dtype = dtype @@ -1679,7 +1868,7 @@ def __init__( self.double_stream_modulation_img = nnx.Linear( in_features=self.inner_dim, out_features=6 * self.inner_dim, - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + use_bias=False, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1687,7 +1876,7 @@ def __init__( self.double_stream_modulation_txt = nnx.Linear( in_features=self.inner_dim, out_features=6 * self.inner_dim, - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + use_bias=False, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1695,7 +1884,7 @@ def __init__( self.single_stream_modulation = nnx.Linear( in_features=self.inner_dim, out_features=3 * self.inner_dim, - bias_init=nnx.with_partitioning(nnx.initializers.zeros, (None,)), + use_bias=False, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1704,6 +1893,7 @@ def __init__( self.x_embedder = nnx.Linear( in_features=in_channels, out_features=self.inner_dim, + use_bias=False, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1711,6 +1901,7 @@ def __init__( self.context_embedder = nnx.Linear( in_features=joint_attention_dim, out_features=self.inner_dim, + use_bias=False, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1723,8 +1914,16 @@ def __init__( dim=self.inner_dim, num_attention_heads=num_attention_heads, attention_head_dim=attention_head_dim, + mlp_ratio=mlp_ratio, + attention_kernel=attention_kernel, + flash_min_seq_length=flash_min_seq_length, + flash_block_sizes=flash_block_sizes, + mesh=mesh, dtype=dtype, weights_dtype=weights_dtype, + ulysses_shards=ulysses_shards, + ulysses_attention_chunks=ulysses_attention_chunks, + use_base2_exp=use_base2_exp, ) for _ in range(num_layers) ] @@ -1737,8 +1936,16 @@ def __init__( dim=self.inner_dim, num_attention_heads=num_attention_heads, attention_head_dim=attention_head_dim, + mlp_ratio=mlp_ratio, + attention_kernel=attention_kernel, + flash_min_seq_length=flash_min_seq_length, + flash_block_sizes=flash_block_sizes, + mesh=mesh, dtype=dtype, weights_dtype=weights_dtype, + ulysses_shards=ulysses_shards, + ulysses_attention_chunks=ulysses_attention_chunks, + use_base2_exp=use_base2_exp, ) for _ in range(num_single_layers) ] @@ -1748,13 +1955,14 @@ def __init__( rngs=rngs, embedding_dim=self.inner_dim, eps=1e-6, + scale_shift_order=scale_shift_order, dtype=dtype, weights_dtype=weights_dtype, ) self.proj_out = nnx.Linear( in_features=self.inner_dim, out_features=in_channels, - use_bias=True, + use_bias=False, dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -1764,12 +1972,13 @@ def __call__( self, hidden_states: jax.Array, encoder_hidden_states: jax.Array, - pooled_projections: jax.Array, - timestep: jax.Array, - img_ids: jax.Array, - txt_ids: jax.Array, + pooled_projections: Optional[jax.Array] = None, + timestep: Optional[jax.Array] = None, + img_ids: Optional[jax.Array] = None, + txt_ids: Optional[jax.Array] = None, guidance: Optional[jax.Array] = None, - ) -> jax.Array: + return_dict: bool = True, + ) -> Union[jax.Array, Transformer2DModelOutput]: hidden_states = self.x_embedder(hidden_states) timestep = timestep * 1000.0 if guidance is not None: @@ -1777,7 +1986,7 @@ def __call__( temb = self.time_text_embed(timestep, guidance, pooled_projections) temb = temb.astype(hidden_states.dtype) - temb_silu = jax.nn.silu(temb) + temb_silu = nnx.silu(temb) double_stream_mod_img = self.double_stream_modulation_img(temb_silu) double_stream_mod_txt = self.double_stream_modulation_txt(temb_silu) single_stream_mod = self.single_stream_modulation(temb_silu) @@ -1821,4 +2030,7 @@ def __call__( hidden_states = hidden_states[:, num_txt_tokens:, ...] hidden_states = self.norm_out(hidden_states, temb) output = self.proj_out(hidden_states) - return output + + if not return_dict: + return (output,) + return Transformer2DModelOutput(sample=output) diff --git a/src/maxdiffusion/models/flux/util.py b/src/maxdiffusion/models/flux/util.py index 952519776..5eb25572c 100644 --- a/src/maxdiffusion/models/flux/util.py +++ b/src/maxdiffusion/models/flux/util.py @@ -17,6 +17,7 @@ # copied from https://github.com/ml-gde/jflux/blob/main/jflux/util.py import os from dataclasses import dataclass +from typing import Any, Optional import jax from jax.typing import DTypeLike @@ -300,17 +301,17 @@ def unpack_latents(latents, batch_size, num_channels_latents, height, width): Unpacks packed sequence of shape (batch_size, (height//16)*(width//16), channels*4) back to the unpacked spatial grid shape (batch_size, channels, height//8, width//8). """ - import numpy as np + import jax.numpy as jnp h_latent = height // 8 w_latent = width // 8 # 1. Reshape to split spatial grid and packed channel blocks - latents = np.reshape(latents, (batch_size, h_latent // 2, w_latent // 2, num_channels_latents, 2, 2)) + latents = jnp.reshape(latents, (batch_size, h_latent // 2, w_latent // 2, num_channels_latents, 2, 2)) # 2. Permute dimensions back to unpacked order - latents = np.transpose(latents, (0, 3, 1, 4, 2, 5)) + latents = jnp.transpose(latents, (0, 3, 1, 4, 2, 5)) # 3. Flatten back to 4D unpacked latent shape - latents = np.reshape(latents, (batch_size, num_channels_latents, h_latent, w_latent)) + latents = jnp.reshape(latents, (batch_size, num_channels_latents, h_latent, w_latent)) return latents @@ -398,11 +399,12 @@ def cast_dict_to_bfloat16_inplace(d, device=None, exclude_keywords=None, parent_ is_excluded = exclude_keywords and any(kw.lower() in current_key.lower() for kw in exclude_keywords) target_dtype = jnp.float32 if is_excluded else jnp.bfloat16 - d[k] = v.astype(target_dtype) - if hasattr(d[k], "block_until_ready"): - d[k].block_until_ready() - del v - gc.collect() + if v.dtype != target_dtype: + d[k] = v.astype(target_dtype) + if hasattr(d[k], "block_until_ready"): + d[k].block_until_ready() + del v + gc.collect() # ----------------------------------------------------------------------------- @@ -410,7 +412,9 @@ def cast_dict_to_bfloat16_inplace(d, device=None, exclude_keywords=None, parent_ # ----------------------------------------------------------------------------- -def load_and_convert_flux_klein_weights(safetensors_path, params, num_double_layers, num_single_layers): +def load_and_convert_flux_klein_weights( + safetensors_path, params, num_double_layers, num_single_layers, dtype=None, pt_state_dict=None +): """ Loads weights from safetensors via zero-copy safetensors.numpy and converts them to JAX parameter dictionary. Supports dynamic layer counts (double and single stream blocks) and sharded safetensors directories. @@ -422,28 +426,30 @@ def load_and_convert_flux_klein_weights(safetensors_path, params, num_double_lay import os import gc - pt_state_dict = {} - if os.path.isdir(safetensors_path): - shards = glob.glob(os.path.join(safetensors_path, "*.safetensors")) - max_logging.log(f"Loading sharded weights from directory: {safetensors_path} (Found {len(shards)} shards)...") - for shard in sorted(shards): - max_logging.log(f"Loading shard: {shard}...") - pt_state_dict.update(load_file(shard)) - else: - max_logging.log(f"Loading weights from: {safetensors_path}") - pt_state_dict = load_file(safetensors_path) + if pt_state_dict is None: + pt_state_dict = {} + if os.path.isdir(safetensors_path): + shards = glob.glob(os.path.join(safetensors_path, "*.safetensors")) + max_logging.log(f"Loading sharded weights from directory: {safetensors_path} (Found {len(shards)} shards)...") + for shard in sorted(shards): + max_logging.log(f"Loading shard: {shard}...") + pt_state_dict.update(load_file(shard)) + else: + max_logging.log(f"Loading weights from: {safetensors_path}") + pt_state_dict = load_file(safetensors_path) max_logging.log("Mapping weights to JAX parameters...") expected_pytree = jax.tree_util.tree_map(lambda leaf: leaf, params) first_leaf = jax.tree_util.tree_leaves(params)[0] - target_dtype = first_leaf.dtype + target_dtype = dtype if dtype is not None else first_leaf.dtype - def convert_and_transpose_tensor(tensor, transpose=False): + def convert_and_transpose_tensor(tensor, transpose=False, is_norm=False): if transpose and len(tensor.shape) == 2: tensor = tensor.T - return jnp.array(tensor, dtype=target_dtype) + leaf_dtype = jnp.float32 if is_norm else target_dtype + return jnp.array(tensor, dtype=leaf_dtype) # Global layers params["context_embedder"]["kernel"] = convert_and_transpose_tensor( @@ -562,77 +568,305 @@ def convert_and_transpose_tensor(tensor, transpose=False): return params -def load_and_convert_vae_weights(safetensors_path, jax_params): +def load_and_convert_flux_klein_nnx_weights( + safetensors_path: str, + nnx_state: Any, + num_double_layers: int, + num_single_layers: int, + dtype=None, + pt_state_dict: Optional[dict] = None, +): + """Loads FLUX.2-Klein weights directly into an NNX State PyTree in target dtype.""" + import glob + import gc + from safetensors.numpy import load_file + import numpy as np + from flax import nnx + + if pt_state_dict is None: + max_logging.log(f"Loading transformer safetensors from: {safetensors_path}") + if os.path.isdir(safetensors_path): + st_files = sorted(glob.glob(os.path.join(safetensors_path, "*.safetensors"))) + else: + st_files = [safetensors_path] + pt_state_dict = {} + for st_file in st_files: + pt_state_dict.update(load_file(st_file)) + + flat_state = dict(nnx.to_flat_state(nnx_state)) + target_dtype = dtype if dtype is not None else jnp.bfloat16 + + def convert_and_transpose_tensor(tensor, transpose=False, is_norm=False): + if transpose and len(tensor.shape) == 2: + tensor = tensor.T + leaf_dtype = jnp.float32 if is_norm else target_dtype + return jnp.array(tensor, dtype=leaf_dtype) + + def set_val(var, tensor): + if hasattr(var, "set_value"): + var.set_value(tensor) + elif hasattr(var, "value"): + var.value = tensor + return var + + # Global layers + global_mappings = [ + ("context_embedder.weight", ("context_embedder", "kernel"), True), + ("x_embedder.weight", ("x_embedder", "kernel"), True), + ("double_stream_modulation_img.linear.weight", ("double_stream_modulation_img", "kernel"), True), + ("double_stream_modulation_txt.linear.weight", ("double_stream_modulation_txt", "kernel"), True), + ("single_stream_modulation.linear.weight", ("single_stream_modulation", "kernel"), True), + ("proj_out.weight", ("proj_out", "kernel"), True), + ("norm_out.linear.weight", ("norm_out", "linear", "kernel"), True), + ] + for pt_key, nnx_key, transpose in global_mappings: + if pt_key in pt_state_dict and nnx_key in flat_state: + set_val(flat_state[nnx_key], convert_and_transpose_tensor(pt_state_dict.pop(pt_key), transpose=transpose)) + + # Timestep / Guidance / Text projections + embedder_mappings = [ + ("time_guidance_embed.timestep_embedder.linear_1", ("time_text_embed", "timestep_embedder", "linear_1")), + ("time_guidance_embed.timestep_embedder.linear_2", ("time_text_embed", "timestep_embedder", "linear_2")), + ("time_guidance_embed.guidance_embedder.linear_1", ("time_text_embed", "guidance_embedder", "linear_1")), + ("time_guidance_embed.guidance_embedder.linear_2", ("time_text_embed", "guidance_embedder", "linear_2")), + ("time_guidance_embed.text_embedder.linear_1", ("time_text_embed", "pooled_embedder", "linear_1")), + ("time_guidance_embed.text_embedder.linear_2", ("time_text_embed", "pooled_embedder", "linear_2")), + ] + for pt_prefix, nnx_prefix in embedder_mappings: + if f"{pt_prefix}.weight" in pt_state_dict and (*nnx_prefix, "kernel") in flat_state: + set_val( + flat_state[(*nnx_prefix, "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(f"{pt_prefix}.weight"), transpose=True), + ) + if f"{pt_prefix}.bias" in pt_state_dict and (*nnx_prefix, "bias") in flat_state: + set_val(flat_state[(*nnx_prefix, "bias")], convert_and_transpose_tensor(pt_state_dict.pop(f"{pt_prefix}.bias"))) + + # Double blocks + for block_idx in range(num_double_layers): + prefix = f"transformer_blocks.{block_idx}." + to_q = pt_state_dict.pop(prefix + "attn.to_q.weight").T + to_k = pt_state_dict.pop(prefix + "attn.to_k.weight").T + to_v = pt_state_dict.pop(prefix + "attn.to_v.weight").T + set_val( + flat_state[("double_blocks", block_idx, "attn", "i_qkv", "kernel")], + jnp.array(np.concatenate([to_q, to_k, to_v], axis=1), dtype=target_dtype), + ) + + add_q = pt_state_dict.pop(prefix + "attn.add_q_proj.weight").T + add_k = pt_state_dict.pop(prefix + "attn.add_k_proj.weight").T + add_v = pt_state_dict.pop(prefix + "attn.add_v_proj.weight").T + set_val( + flat_state[("double_blocks", block_idx, "attn", "e_qkv", "kernel")], + jnp.array(np.concatenate([add_q, add_k, add_v], axis=1), dtype=target_dtype), + ) + + set_val( + flat_state[("double_blocks", block_idx, "attn", "i_proj", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "attn.to_out.0.weight"), transpose=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "attn", "e_proj", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "attn.to_add_out.weight"), transpose=True), + ) + + set_val( + flat_state[("double_blocks", block_idx, "attn", "query_norm", "scale")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "attn.norm_q.weight"), is_norm=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "attn", "key_norm", "scale")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "attn.norm_k.weight"), is_norm=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "attn", "encoder_query_norm", "scale")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "attn.norm_added_q.weight"), is_norm=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "attn", "encoder_key_norm", "scale")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "attn.norm_added_k.weight"), is_norm=True), + ) + + set_val( + flat_state[("double_blocks", block_idx, "ff", "linear_in", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "ff.linear_in.weight"), transpose=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "ff", "linear_out", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "ff.linear_out.weight"), transpose=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "ff_context", "linear_in", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "ff_context.linear_in.weight"), transpose=True), + ) + set_val( + flat_state[("double_blocks", block_idx, "ff_context", "linear_out", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(prefix + "ff_context.linear_out.weight"), transpose=True), + ) + + # Single blocks + for block_idx in range(num_single_layers): + s_prefix = f"single_transformer_blocks.{block_idx}." + set_val( + flat_state[("single_blocks", block_idx, "linear1", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(s_prefix + "attn.to_qkv_mlp_proj.weight"), transpose=True), + ) + set_val( + flat_state[("single_blocks", block_idx, "linear2", "kernel")], + convert_and_transpose_tensor(pt_state_dict.pop(s_prefix + "attn.to_out.weight"), transpose=True), + ) + set_val( + flat_state[("single_blocks", block_idx, "attn", "query_norm", "scale")], + convert_and_transpose_tensor(pt_state_dict.pop(s_prefix + "attn.norm_q.weight"), is_norm=True), + ) + set_val( + flat_state[("single_blocks", block_idx, "attn", "key_norm", "scale")], + convert_and_transpose_tensor(pt_state_dict.pop(s_prefix + "attn.norm_k.weight"), is_norm=True), + ) + + for path, var in flat_state.items(): + val = var.get_value() if hasattr(var, "get_value") else getattr(var, "value", var) + if isinstance(val, jax.ShapeDtypeStruct): + set_val(var, jnp.zeros(val.shape, dtype=val.dtype)) + + del pt_state_dict + gc.collect() + max_logging.log("NNX Weight conversion complete & verified!") + return nnx.from_flat_state(flat_state) + + +def patchify_latents(latents): + """Patchifies latents: (B, C, H, W) -> (B, C*4, H//2, W//2).""" + import jax.numpy as jnp + + batch_size, num_channels, height, width = latents.shape + latents = latents.reshape((batch_size, num_channels, height // 2, 2, width // 2, 2)) + latents = jnp.transpose(latents, (0, 1, 3, 5, 2, 4)) + latents = latents.reshape((batch_size, num_channels * 4, height // 2, width // 2)) + return latents + + +def prepare_multi_image_ids(image_latents_list, scale=10): + """Generates 4D position IDs (T, H, W, L) for a sequence of reference image latents. + + For the k-th image, T = scale * (k + 1). + image_latents_list: list of arrays with shape (1, C, H, W) or (C, H, W). + Returns: array of shape (1, total_tokens, 4). + """ + import jax.numpy as jnp + + all_ids = [] + for idx, latent in enumerate(image_latents_list): + if latent.ndim == 4: + latent = latent[0] + _, h, w = latent.shape + t_val = scale * (idx + 1) + t = jnp.full((h * w, 1), t_val, dtype=jnp.int32) + h_grid, w_grid = jnp.meshgrid(jnp.arange(h, dtype=jnp.int32), jnp.arange(w, dtype=jnp.int32), indexing="ij") + h_coords = h_grid.reshape(-1, 1) + w_coords = w_grid.reshape(-1, 1) + l_coords = jnp.zeros((h * w, 1), dtype=jnp.int32) + coords = jnp.concatenate([t, h_coords, w_coords, l_coords], axis=-1) + all_ids.append(coords) + combined = jnp.concatenate(all_ids, axis=0) + return jnp.expand_dims(combined, axis=0) + + +def prepare_image_latents(vae, images, bn_mean, bn_std, scale=10): + """Encodes, patchifies, normalizes, packs, and generates 4D RoPE IDs for a list of reference images. + + images: list of arrays of shape (1, 3, H_k, W_k) or (3, H_k, W_k) in range [-1, 1]. + Returns: + image_latents_concat: shape (1, total_ref_tokens, 128) + image_latent_ids: shape (1, total_ref_tokens, 4) + """ + import jax.numpy as jnp + from einops import rearrange + + norm_latents = [] + for img in images: + if img.ndim == 3: + img = jnp.expand_dims(img, axis=0) + raw_latents = vae.encode(img) # (1, 32, H/8, W/8) + patchified = patchify_latents(raw_latents) # (1, 128, H/16, W/16) + normalized = (patchified - bn_mean) / bn_std + norm_latents.append(normalized) + + image_latent_ids = prepare_multi_image_ids(norm_latents, scale=scale) + + packed_latents = [] + for latent in norm_latents: + packed = rearrange(latent, "b c h w -> b (h w) c") + packed_latents.append(packed) + + image_latents_concat = jnp.concatenate(packed_latents, axis=1) + return image_latents_concat, image_latent_ids + + +def load_and_convert_vae_weights(safetensors_path, jax_params, dtype=None, pt_state_dict=None): """Loads VAE weights from safetensors via zero-copy safetensors.numpy, maps them to JAX, and extracts BN stats.""" from safetensors.numpy import load_file import flax import jax.numpy as jnp - max_logging.log(f"Loading VAE weights from: {safetensors_path}") - pt_state_dict = load_file(safetensors_path) - - def get_pytorch_weight_tensor(key): - return pt_state_dict[key] + if pt_state_dict is None: + max_logging.log(f"Loading VAE weights from: {safetensors_path}") + pt_state_dict = load_file(safetensors_path) # Unfreeze JAX params so we can load the weights jax_params = flax.core.unfreeze(jax_params) - # Map weights - max_logging.log("Mapping VAE decoder weights to JAX parameters...") - - # post_quant_conv - jax_params["post_quant_conv"]["kernel"] = jnp.array( - get_pytorch_weight_tensor("post_quant_conv.weight").transpose(2, 3, 1, 0) - ) - jax_params["post_quant_conv"]["bias"] = jnp.array(get_pytorch_weight_tensor("post_quant_conv.bias")) + first_leaf = jax.tree_util.tree_leaves(jax_params)[0] + target_dtype = dtype if dtype is not None else first_leaf.dtype + + def get_pytorch_weight_tensor(key, dtype_val=target_dtype): + tensor = pt_state_dict[key] + is_norm = any(kw in key.lower() for kw in ("norm", "layernorm", "rmsnorm", "groupnorm")) + leaf_dtype = jnp.float32 if is_norm else dtype_val + return jnp.array(tensor, dtype=leaf_dtype) + + # 1. Map VAE Encoder Weights + if "encoder" in jax_params: + max_logging.log("Mapping VAE encoder weights to JAX parameters...") + enc_jax = jax_params["encoder"] + + if "encoder.conv_in.weight" in pt_state_dict: + enc_jax["conv_in"]["kernel"] = jnp.array(get_pytorch_weight_tensor("encoder.conv_in.weight").transpose(2, 3, 1, 0)) + enc_jax["conv_in"]["bias"] = jnp.array(get_pytorch_weight_tensor("encoder.conv_in.bias")) + + for b_idx in range(4): + down_block_pt = f"encoder.down_blocks.{b_idx}" + down_block_jax = enc_jax[f"down_blocks_{b_idx}"] + + for r_idx in range(2): + res_pt = f"{down_block_pt}.resnets.{r_idx}" + res_jax = down_block_jax[f"resnets_{r_idx}"] + + res_jax["norm1"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm1.weight")) + res_jax["norm1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm1.bias")) + res_jax["conv1"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv1.weight").transpose(2, 3, 1, 0)) + res_jax["conv1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv1.bias")) + + res_jax["norm2"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm2.weight")) + res_jax["norm2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm2.bias")) + res_jax["conv2"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv2.weight").transpose(2, 3, 1, 0)) + res_jax["conv2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv2.bias")) + + shortcut_key = f"{res_pt}.conv_shortcut.weight" + if shortcut_key in pt_state_dict: + res_jax["conv_shortcut"]["kernel"] = jnp.array(get_pytorch_weight_tensor(shortcut_key).transpose(2, 3, 1, 0)) + res_jax["conv_shortcut"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv_shortcut.bias")) + + if b_idx < 3: + downsampler_pt = f"{down_block_pt}.downsamplers.0" + downsampler_jax = down_block_jax["downsamplers_0"] + downsampler_jax["conv"]["kernel"] = jnp.array( + get_pytorch_weight_tensor(f"{downsampler_pt}.conv.weight").transpose(2, 3, 1, 0) + ) + downsampler_jax["conv"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{downsampler_pt}.conv.bias")) - # decoder.conv_in - jax_params["decoder"]["conv_in"]["kernel"] = jnp.array( - get_pytorch_weight_tensor("decoder.conv_in.weight").transpose(2, 3, 1, 0) - ) - jax_params["decoder"]["conv_in"]["bias"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_in.bias")) - - # decoder.mid_block - # resnets - for idx in [0, 1]: - res_jax = jax_params["decoder"]["mid_block"][f"resnets_{idx}"] - res_pt_prefix = f"decoder.mid_block.resnets.{idx}" - - res_jax["norm1"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm1.weight")) - res_jax["norm1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm1.bias")) - res_jax["conv1"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.conv1.weight").transpose(2, 3, 1, 0)) - res_jax["conv1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.conv1.bias")) - - res_jax["norm2"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm2.weight")) - res_jax["norm2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm2.bias")) - res_jax["conv2"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.conv2.weight").transpose(2, 3, 1, 0)) - res_jax["conv2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.conv2.bias")) - - # attentions - attn_pt_prefix = "decoder.mid_block.attentions.0" - attn_jax = jax_params["decoder"]["mid_block"]["attentions_0"] - - attn_jax["group_norm"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.group_norm.weight")) - attn_jax["group_norm"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.group_norm.bias")) - - attn_jax["query"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_q.weight").T) - attn_jax["query"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_q.bias")) - attn_jax["key"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_k.weight").T) - attn_jax["key"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_k.bias")) - attn_jax["value"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_v.weight").T) - attn_jax["value"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_v.bias")) - - attn_jax["proj_attn"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_out.0.weight").T) - attn_jax["proj_attn"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_out.0.bias")) - - # decoder.up_blocks - for b_idx in range(4): - up_block_jax = jax_params["decoder"][f"up_blocks_{b_idx}"] - up_block_pt = f"decoder.up_blocks.{b_idx}" - - for r_idx in range(3): - res_jax = up_block_jax[f"resnets_{r_idx}"] - res_pt = f"{up_block_pt}.resnets.{r_idx}" + for r_idx in range(2): + res_pt = f"encoder.mid_block.resnets.{r_idx}" + res_jax = enc_jax["mid_block"][f"resnets_{r_idx}"] res_jax["norm1"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm1.weight")) res_jax["norm1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm1.bias")) @@ -644,27 +878,112 @@ def get_pytorch_weight_tensor(key): res_jax["conv2"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv2.weight").transpose(2, 3, 1, 0)) res_jax["conv2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv2.bias")) - shortcut_key = f"{res_pt}.conv_shortcut.weight" - if shortcut_key in pt_state_dict: - res_jax["conv_shortcut"]["kernel"] = jnp.array(get_pytorch_weight_tensor(shortcut_key).transpose(2, 3, 1, 0)) - res_jax["conv_shortcut"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv_shortcut.bias")) + attn_enc_pt = "encoder.mid_block.attentions.0" + attn_enc_jax = enc_jax["mid_block"]["attentions_0"] + attn_enc_jax["group_norm"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.group_norm.weight")) + attn_enc_jax["group_norm"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.group_norm.bias")) + attn_enc_jax["query"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_q.weight").T) + attn_enc_jax["query"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_q.bias")) + attn_enc_jax["key"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_k.weight").T) + attn_enc_jax["key"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_k.bias")) + attn_enc_jax["value"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_v.weight").T) + attn_enc_jax["value"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_v.bias")) + attn_enc_jax["proj_attn"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_out.0.weight").T) + attn_enc_jax["proj_attn"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_enc_pt}.to_out.0.bias")) + + enc_jax["conv_norm_out"]["scale"] = jnp.array(get_pytorch_weight_tensor("encoder.conv_norm_out.weight")) + enc_jax["conv_norm_out"]["bias"] = jnp.array(get_pytorch_weight_tensor("encoder.conv_norm_out.bias")) + enc_jax["conv_out"]["kernel"] = jnp.array(get_pytorch_weight_tensor("encoder.conv_out.weight").transpose(2, 3, 1, 0)) + enc_jax["conv_out"]["bias"] = jnp.array(get_pytorch_weight_tensor("encoder.conv_out.bias")) + + if "quant_conv" in jax_params and "quant_conv.weight" in pt_state_dict: + jax_params["quant_conv"]["kernel"] = jnp.array(get_pytorch_weight_tensor("quant_conv.weight").transpose(2, 3, 1, 0)) + jax_params["quant_conv"]["bias"] = jnp.array(get_pytorch_weight_tensor("quant_conv.bias")) + + # 2. Map VAE Decoder Weights + max_logging.log("Mapping VAE decoder weights to JAX parameters...") + + if "post_quant_conv" in jax_params and "post_quant_conv.weight" in pt_state_dict: + jax_params["post_quant_conv"]["kernel"] = jnp.array( + get_pytorch_weight_tensor("post_quant_conv.weight").transpose(2, 3, 1, 0) + ) + jax_params["post_quant_conv"]["bias"] = jnp.array(get_pytorch_weight_tensor("post_quant_conv.bias")) + + if "decoder" in jax_params: + dec_jax = jax_params["decoder"] + dec_jax["conv_in"]["kernel"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_in.weight").transpose(2, 3, 1, 0)) + dec_jax["conv_in"]["bias"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_in.bias")) - if b_idx < 3: - upsampler_jax = up_block_jax["upsamplers_0"] - upsampler_pt = f"{up_block_pt}.upsamplers.0" + for idx in [0, 1]: + res_jax = dec_jax["mid_block"][f"resnets_{idx}"] + res_pt_prefix = f"decoder.mid_block.resnets.{idx}" - upsampler_jax["conv"]["kernel"] = jnp.array( - get_pytorch_weight_tensor(f"{upsampler_pt}.conv.weight").transpose(2, 3, 1, 0) + res_jax["norm1"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm1.weight")) + res_jax["norm1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm1.bias")) + res_jax["conv1"]["kernel"] = jnp.array( + get_pytorch_weight_tensor(f"{res_pt_prefix}.conv1.weight").transpose(2, 3, 1, 0) ) - upsampler_jax["conv"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{upsampler_pt}.conv.bias")) + res_jax["conv1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.conv1.bias")) - # decoder.conv_norm_out & conv_out - jax_params["decoder"]["conv_norm_out"]["scale"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_norm_out.weight")) - jax_params["decoder"]["conv_norm_out"]["bias"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_norm_out.bias")) - jax_params["decoder"]["conv_out"]["kernel"] = jnp.array( - get_pytorch_weight_tensor("decoder.conv_out.weight").transpose(2, 3, 1, 0) - ) - jax_params["decoder"]["conv_out"]["bias"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_out.bias")) + res_jax["norm2"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm2.weight")) + res_jax["norm2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.norm2.bias")) + res_jax["conv2"]["kernel"] = jnp.array( + get_pytorch_weight_tensor(f"{res_pt_prefix}.conv2.weight").transpose(2, 3, 1, 0) + ) + res_jax["conv2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt_prefix}.conv2.bias")) + + attn_pt_prefix = "decoder.mid_block.attentions.0" + attn_jax = dec_jax["mid_block"]["attentions_0"] + + attn_jax["group_norm"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.group_norm.weight")) + attn_jax["group_norm"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.group_norm.bias")) + + attn_jax["query"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_q.weight").T) + attn_jax["query"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_q.bias")) + attn_jax["key"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_k.weight").T) + attn_jax["key"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_k.bias")) + attn_jax["value"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_v.weight").T) + attn_jax["value"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_v.bias")) + + attn_jax["proj_attn"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_out.0.weight").T) + attn_jax["proj_attn"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{attn_pt_prefix}.to_out.0.bias")) + + for b_idx in range(4): + up_block_jax = dec_jax[f"up_blocks_{b_idx}"] + up_block_pt = f"decoder.up_blocks.{b_idx}" + + for r_idx in range(3): + res_jax = up_block_jax[f"resnets_{r_idx}"] + res_pt = f"{up_block_pt}.resnets.{r_idx}" + + res_jax["norm1"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm1.weight")) + res_jax["norm1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm1.bias")) + res_jax["conv1"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv1.weight").transpose(2, 3, 1, 0)) + res_jax["conv1"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv1.bias")) + + res_jax["norm2"]["scale"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm2.weight")) + res_jax["norm2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.norm2.bias")) + res_jax["conv2"]["kernel"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv2.weight").transpose(2, 3, 1, 0)) + res_jax["conv2"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv2.bias")) + + shortcut_key = f"{res_pt}.conv_shortcut.weight" + if shortcut_key in pt_state_dict: + res_jax["conv_shortcut"]["kernel"] = jnp.array(get_pytorch_weight_tensor(shortcut_key).transpose(2, 3, 1, 0)) + res_jax["conv_shortcut"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{res_pt}.conv_shortcut.bias")) + + if b_idx < 3: + upsampler_jax = up_block_jax["upsamplers_0"] + upsampler_pt = f"{up_block_pt}.upsamplers.0" + + upsampler_jax["conv"]["kernel"] = jnp.array( + get_pytorch_weight_tensor(f"{upsampler_pt}.conv.weight").transpose(2, 3, 1, 0) + ) + upsampler_jax["conv"]["bias"] = jnp.array(get_pytorch_weight_tensor(f"{upsampler_pt}.conv.bias")) + + dec_jax["conv_norm_out"]["scale"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_norm_out.weight")) + dec_jax["conv_norm_out"]["bias"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_norm_out.bias")) + dec_jax["conv_out"]["kernel"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_out.weight").transpose(2, 3, 1, 0)) + dec_jax["conv_out"]["bias"] = jnp.array(get_pytorch_weight_tensor("decoder.conv_out.bias")) jax_params = jax.tree_util.tree_map( lambda leaf: jnp.zeros(leaf.shape, dtype=leaf.dtype) if isinstance(leaf, jax.ShapeDtypeStruct) else leaf, jax_params diff --git a/src/maxdiffusion/models/flux/vae/__init__.py b/src/maxdiffusion/models/flux/vae/__init__.py new file mode 100644 index 000000000..e69de29bb diff --git a/src/maxdiffusion/models/flux/vae/autoencoder_kl_flux2_nnx.py b/src/maxdiffusion/models/flux/vae/autoencoder_kl_flux2_nnx.py new file mode 100644 index 000000000..01adabd1f --- /dev/null +++ b/src/maxdiffusion/models/flux/vae/autoencoder_kl_flux2_nnx.py @@ -0,0 +1,799 @@ +""" +Copyright 2026 Google LLC + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + https://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +""" + +import math +from typing import Optional, Tuple + +import jax +import jax.numpy as jnp +from flax import nnx + + +class NNXUpsample2D(nnx.Module): + """2D Nearest-neighbor Upsample + Conv layer in NNX.""" + + def __init__( + self, + in_channels: int, + out_channels: Optional[int] = None, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + out_channels = out_channels or in_channels + self.conv = nnx.Conv( + in_features=in_channels, + out_features=out_channels, + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + batch, height, width, channels = x.shape + x = jnp.broadcast_to(x[:, :, None, :, None, :], (batch, height, 2, width, 2, channels)) + x = jnp.reshape(x, (batch, height * 2, width * 2, channels)) + return self.conv(x) + + +class NNXDownsample2D(nnx.Module): + """2D Downsample layer with asymmetric padding in NNX.""" + + def __init__( + self, + in_channels: int, + out_channels: Optional[int] = None, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + out_channels = out_channels or in_channels + self.conv = nnx.Conv( + in_features=in_channels, + out_features=out_channels, + kernel_size=(3, 3), + strides=(2, 2), + padding="VALID", + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + pad_width = ((0, 0), (0, 1), (0, 1), (0, 0)) + x = jnp.pad(x, pad_width) + return self.conv(x) + + +class NNXResnetBlock2D(nnx.Module): + """2D ResNet Block with GroupNorm and SiLU activations in NNX.""" + + def __init__( + self, + in_channels: int, + out_channels: Optional[int] = None, + groups: int = 32, + use_conv_shortcut: Optional[bool] = None, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + out_channels = out_channels or in_channels + self.in_channels = in_channels + self.out_channels = out_channels + + self.norm1 = nnx.GroupNorm( + num_groups=groups, + num_features=in_channels, + epsilon=1e-6, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.conv1 = nnx.Conv( + in_features=in_channels, + out_features=out_channels, + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.norm2 = nnx.GroupNorm( + num_groups=groups, + num_features=out_channels, + epsilon=1e-6, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.conv2 = nnx.Conv( + in_features=out_channels, + out_features=out_channels, + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + use_shortcut = (in_channels != out_channels) if use_conv_shortcut is None else use_conv_shortcut + if use_shortcut: + self.conv_shortcut = nnx.Conv( + in_features=in_channels, + out_features=out_channels, + kernel_size=(1, 1), + strides=(1, 1), + padding="VALID", + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + else: + self.conv_shortcut = None + + def __call__(self, x: jax.Array) -> jax.Array: + residual = self.conv_shortcut(x) if self.conv_shortcut is not None else x + h = self.norm1(x) + h = nnx.silu(h) + h = self.conv1(h) + h = self.norm2(h) + h = nnx.silu(h) + h = self.conv2(h) + return h + residual + + +class NNXAttentionBlock(nnx.Module): + """Self-Attention block with GroupNorm in NNX.""" + + def __init__( + self, + channels: int, + groups: int = 32, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + self.channels = channels + self.group_norm = nnx.GroupNorm( + num_groups=groups, + num_features=channels, + epsilon=1e-6, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.to_q = nnx.Linear( + in_features=channels, + out_features=channels, + use_bias=True, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.to_k = nnx.Linear( + in_features=channels, + out_features=channels, + use_bias=True, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.to_v = nnx.Linear( + in_features=channels, + out_features=channels, + use_bias=True, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.to_out = nnx.Linear( + in_features=channels, + out_features=channels, + use_bias=True, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + residual = x + b, h, w, c = x.shape + h_states = self.group_norm(x) + h_flat = h_states.reshape((b, h * w, c)) + + q = self.to_q(h_flat) + k = self.to_k(h_flat) + v = self.to_v(h_flat) + + scale = 1.0 / math.sqrt(c) + attn_weights = jnp.einsum("bqc,bkc->bqk", q * scale, k) + attn_weights = jax.nn.softmax(attn_weights, axis=-1) + + out = jnp.einsum("bqk,bkc->bqc", attn_weights, v) + out = self.to_out(out) + out = out.reshape((b, h, w, c)) + return out + residual + + +class NNXUNetMidBlock2D(nnx.Module): + """Mid-Block module in NNX with resnets and attention.""" + + def __init__( + self, + in_channels: int, + groups: int = 32, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + self.resnets_0 = NNXResnetBlock2D( + in_channels=in_channels, + out_channels=in_channels, + groups=groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.attentions_0 = NNXAttentionBlock( + channels=in_channels, + groups=groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.resnets_1 = NNXResnetBlock2D( + in_channels=in_channels, + out_channels=in_channels, + groups=groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + x = self.resnets_0(x) + x = self.attentions_0(x) + x = self.resnets_1(x) + return x + + +class NNXDownEncoderBlock2D(nnx.Module): + """Down-Encoder block containing ResNet layers and an optional Downsampler in NNX.""" + + def __init__( + self, + in_channels: int, + out_channels: int, + num_layers: int = 2, + groups: int = 32, + add_downsample: bool = True, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + resnets = [] + for i in range(num_layers): + in_ch = in_channels if i == 0 else out_channels + resnets.append( + NNXResnetBlock2D( + in_channels=in_ch, + out_channels=out_channels, + groups=groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + ) + self.resnets = nnx.List(resnets) + + if add_downsample: + self.downsamplers_0 = NNXDownsample2D( + in_channels=out_channels, + out_channels=out_channels, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + else: + self.downsamplers_0 = None + + def __call__(self, x: jax.Array) -> jax.Array: + for resnet in self.resnets: + x = resnet(x) + if self.downsamplers_0 is not None: + x = self.downsamplers_0(x) + return x + + +class NNXUpDecoderBlock2D(nnx.Module): + """Up-Decoder block containing ResNet layers and an optional Upsampler in NNX.""" + + def __init__( + self, + in_channels: int, + out_channels: int, + num_layers: int = 3, + groups: int = 32, + add_upsample: bool = True, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + resnets = [] + for i in range(num_layers): + in_ch = in_channels if i == 0 else out_channels + resnets.append( + NNXResnetBlock2D( + in_channels=in_ch, + out_channels=out_channels, + groups=groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + ) + self.resnets = nnx.List(resnets) + + if add_upsample: + self.upsamplers_0 = NNXUpsample2D( + in_channels=out_channels, + out_channels=out_channels, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + else: + self.upsamplers_0 = None + + def __call__(self, x: jax.Array) -> jax.Array: + for resnet in self.resnets: + x = resnet(x) + if self.upsamplers_0 is not None: + x = self.upsamplers_0(x) + return x + + +class NNXEncoder(nnx.Module): + """Complete VAE Encoder in NNX.""" + + def __init__( + self, + in_channels: int = 3, + out_channels: int = 32, + block_out_channels: Tuple[int, ...] = (128, 256, 512, 512), + layers_per_block: int = 2, + norm_num_groups: int = 32, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + self.conv_in = nnx.Conv( + in_features=in_channels, + out_features=block_out_channels[0], + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + down_blocks = [] + output_ch = block_out_channels[0] + for i, ch in enumerate(block_out_channels): + input_ch = output_ch + output_ch = ch + is_final = i == len(block_out_channels) - 1 + down_blocks.append( + NNXDownEncoderBlock2D( + in_channels=input_ch, + out_channels=output_ch, + num_layers=layers_per_block, + groups=norm_num_groups, + add_downsample=not is_final, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + ) + self.down_blocks = nnx.List(down_blocks) + + self.mid_block = NNXUNetMidBlock2D( + in_channels=block_out_channels[-1], + groups=norm_num_groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + self.conv_norm_out = nnx.GroupNorm( + num_groups=norm_num_groups, + num_features=block_out_channels[-1], + epsilon=1e-6, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.conv_out = nnx.Conv( + in_features=block_out_channels[-1], + out_features=2 * out_channels, # double_z for Gaussian distribution moments + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + x = self.conv_in(x) + for block in self.down_blocks: + x = block(x) + x = self.mid_block(x) + x = self.conv_norm_out(x) + x = nnx.silu(x) + x = self.conv_out(x) + return x + + +class NNXDecoder(nnx.Module): + """Complete VAE Decoder in NNX.""" + + def __init__( + self, + in_channels: int = 32, + out_channels: int = 3, + block_out_channels: Tuple[int, ...] = (128, 256, 512, 512), + layers_per_block: int = 3, + norm_num_groups: int = 32, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + reversed_channels = list(reversed(block_out_channels)) + self.conv_in = nnx.Conv( + in_features=in_channels, + out_features=reversed_channels[0], + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + self.mid_block = NNXUNetMidBlock2D( + in_channels=reversed_channels[0], + groups=norm_num_groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + up_blocks = [] + output_ch = reversed_channels[0] + for i, ch in enumerate(reversed_channels): + input_ch = output_ch + output_ch = ch + is_final = i == len(reversed_channels) - 1 + up_blocks.append( + NNXUpDecoderBlock2D( + in_channels=input_ch, + out_channels=output_ch, + num_layers=layers_per_block, + groups=norm_num_groups, + add_upsample=not is_final, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + ) + self.up_blocks = nnx.List(up_blocks) + + self.conv_norm_out = nnx.GroupNorm( + num_groups=norm_num_groups, + num_features=reversed_channels[-1], + epsilon=1e-6, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.conv_out = nnx.Conv( + in_features=reversed_channels[-1], + out_features=out_channels, + kernel_size=(3, 3), + strides=(1, 1), + padding=((1, 1), (1, 1)), + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def __call__(self, x: jax.Array) -> jax.Array: + x = self.conv_in(x) + x = self.mid_block(x) + for block in self.up_blocks: + x = block(x) + x = self.conv_norm_out(x) + x = nnx.silu(x) + x = self.conv_out(x) + return x + + +class NNXAutoencoderKLFlux2(nnx.Module): + """Full FLUX.2-Klein Variational Autoencoder (VAE) in Flax NNX.""" + + def __init__( + self, + in_channels: int = 3, + out_channels: int = 3, + latent_channels: int = 32, + block_out_channels: Tuple[int, ...] = (128, 256, 512, 512), + layers_per_block: int = 2, + norm_num_groups: int = 32, + rngs: Optional[nnx.Rngs] = None, + dtype: jnp.dtype = jnp.float32, + param_dtype: jnp.dtype = jnp.float32, + ): + rngs = rngs or nnx.Rngs(0) + self.latent_channels = latent_channels + self.dtype = dtype + + self.encoder = NNXEncoder( + in_channels=in_channels, + out_channels=latent_channels, + block_out_channels=block_out_channels, + layers_per_block=layers_per_block, + norm_num_groups=norm_num_groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.quant_conv = nnx.Conv( + in_features=2 * latent_channels, + out_features=2 * latent_channels, + kernel_size=(1, 1), + strides=(1, 1), + padding="VALID", + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.post_quant_conv = nnx.Conv( + in_features=latent_channels, + out_features=latent_channels, + kernel_size=(1, 1), + strides=(1, 1), + padding="VALID", + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + self.decoder = NNXDecoder( + in_channels=latent_channels, + out_channels=out_channels, + block_out_channels=block_out_channels, + layers_per_block=layers_per_block + 1, # 3 resnet blocks in decoder + norm_num_groups=norm_num_groups, + rngs=rngs, + dtype=dtype, + param_dtype=param_dtype, + ) + + def encode(self, sample: jax.Array) -> jax.Array: + """Encodes image tensor of shape (B, 3, H, W) to mode latents of shape (B, 32, H/8, W/8).""" + # Transpose to channels last (B, H, W, 3) + x = jnp.transpose(sample, (0, 2, 3, 1)) + h = self.encoder(x) + moments = self.quant_conv(h) + # Extract mean / mode (first latent_channels) + mean, _ = jnp.split(moments, 2, axis=-1) # (B, H/8, W/8, 32) + # Transpose back to (B, 32, H/8, W/8) + return jnp.transpose(mean, (0, 3, 1, 2)) + + def decode(self, latents: jax.Array) -> jax.Array: + """Decodes latent tensor of shape (B, 32, H/8, W/8) to image tensor of shape (B, 3, H, W).""" + # Transpose to channels last (B, H/8, W/8, 32) + z = jnp.transpose(latents, (0, 2, 3, 1)) + h = self.post_quant_conv(z) + img = self.decoder(h) + # Transpose back to (B, 3, H, W) + return jnp.transpose(img, (0, 3, 1, 2)) + + +def load_and_convert_flux2klein_nnx_vae_weights( + safetensors_path: str, + nnx_vae: NNXAutoencoderKLFlux2, + dtype: Optional[jnp.dtype] = None, + pt_state_dict: Optional[dict] = None, +): + """Directly loads and maps PyTorch safetensors into NNXAutoencoderKLFlux2 State.""" + from safetensors.numpy import load_file + + if pt_state_dict is None: + pt_state_dict = load_file(safetensors_path) + + target_dtype = dtype if dtype is not None else jnp.float32 + + def get_pt_tensor(key, is_norm=False): + tensor = pt_state_dict[key] + leaf_dtype = jnp.float32 if is_norm else target_dtype + return jnp.array(tensor, dtype=leaf_dtype) + + def get_conv_kernel(key): + return jnp.array(pt_state_dict[key].transpose(2, 3, 1, 0), dtype=target_dtype) + + def get_linear_kernel(key): + return jnp.array(pt_state_dict[key].T, dtype=target_dtype) + + flat_state = dict(nnx.to_flat_state(nnx.state(nnx_vae, nnx.Param))) + + def set_val(var, val): + var[...] = val + + # ========================================================================= + # 1. ENCODER + # ========================================================================= + set_val(flat_state[("encoder", "conv_in", "kernel")], get_conv_kernel("encoder.conv_in.weight")) + set_val(flat_state[("encoder", "conv_in", "bias")], get_pt_tensor("encoder.conv_in.bias")) + + for b_idx in range(4): + down_block_pt = f"encoder.down_blocks.{b_idx}" + for r_idx in range(2): + res_pt = f"{down_block_pt}.resnets.{r_idx}" + res_path = ("encoder", "down_blocks", b_idx, "resnets", r_idx) + + set_val(flat_state[res_path + ("norm1", "scale")], get_pt_tensor(f"{res_pt}.norm1.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm1", "bias")], get_pt_tensor(f"{res_pt}.norm1.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv1", "kernel")], get_conv_kernel(f"{res_pt}.conv1.weight")) + set_val(flat_state[res_path + ("conv1", "bias")], get_pt_tensor(f"{res_pt}.conv1.bias")) + + set_val(flat_state[res_path + ("norm2", "scale")], get_pt_tensor(f"{res_pt}.norm2.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm2", "bias")], get_pt_tensor(f"{res_pt}.norm2.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv2", "kernel")], get_conv_kernel(f"{res_pt}.conv2.weight")) + set_val(flat_state[res_path + ("conv2", "bias")], get_pt_tensor(f"{res_pt}.conv2.bias")) + + shortcut_key = f"{res_pt}.conv_shortcut.weight" + if shortcut_key in pt_state_dict: + set_val(flat_state[res_path + ("conv_shortcut", "kernel")], get_conv_kernel(shortcut_key)) + set_val(flat_state[res_path + ("conv_shortcut", "bias")], get_pt_tensor(f"{res_pt}.conv_shortcut.bias")) + + if b_idx < 3: + ds_pt = f"{down_block_pt}.downsamplers.0.conv" + ds_path = ("encoder", "down_blocks", b_idx, "downsamplers_0", "conv") + set_val(flat_state[ds_path + ("kernel",)], get_conv_kernel(f"{ds_pt}.weight")) + set_val(flat_state[ds_path + ("bias",)], get_pt_tensor(f"{ds_pt}.bias")) + + # Encoder Mid Block + for r_idx in [0, 1]: + res_pt = f"encoder.mid_block.resnets.{r_idx}" + res_path = ("encoder", "mid_block", f"resnets_{r_idx}") + set_val(flat_state[res_path + ("norm1", "scale")], get_pt_tensor(f"{res_pt}.norm1.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm1", "bias")], get_pt_tensor(f"{res_pt}.norm1.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv1", "kernel")], get_conv_kernel(f"{res_pt}.conv1.weight")) + set_val(flat_state[res_path + ("conv1", "bias")], get_pt_tensor(f"{res_pt}.conv1.bias")) + set_val(flat_state[res_path + ("norm2", "scale")], get_pt_tensor(f"{res_pt}.norm2.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm2", "bias")], get_pt_tensor(f"{res_pt}.norm2.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv2", "kernel")], get_conv_kernel(f"{res_pt}.conv2.weight")) + set_val(flat_state[res_path + ("conv2", "bias")], get_pt_tensor(f"{res_pt}.conv2.bias")) + + attn_pt = "encoder.mid_block.attentions.0" + attn_path = ("encoder", "mid_block", "attentions_0") + set_val(flat_state[attn_path + ("group_norm", "scale")], get_pt_tensor(f"{attn_pt}.group_norm.weight", is_norm=True)) + set_val(flat_state[attn_path + ("group_norm", "bias")], get_pt_tensor(f"{attn_pt}.group_norm.bias", is_norm=True)) + set_val(flat_state[attn_path + ("to_q", "kernel")], get_linear_kernel(f"{attn_pt}.to_q.weight")) + set_val(flat_state[attn_path + ("to_q", "bias")], get_pt_tensor(f"{attn_pt}.to_q.bias")) + set_val(flat_state[attn_path + ("to_k", "kernel")], get_linear_kernel(f"{attn_pt}.to_k.weight")) + set_val(flat_state[attn_path + ("to_k", "bias")], get_pt_tensor(f"{attn_pt}.to_k.bias")) + set_val(flat_state[attn_path + ("to_v", "kernel")], get_linear_kernel(f"{attn_pt}.to_v.weight")) + set_val(flat_state[attn_path + ("to_v", "bias")], get_pt_tensor(f"{attn_pt}.to_v.bias")) + set_val(flat_state[attn_path + ("to_out", "kernel")], get_linear_kernel(f"{attn_pt}.to_out.0.weight")) + set_val(flat_state[attn_path + ("to_out", "bias")], get_pt_tensor(f"{attn_pt}.to_out.0.bias")) + + set_val(flat_state[("encoder", "conv_norm_out", "scale")], get_pt_tensor("encoder.conv_norm_out.weight", is_norm=True)) + set_val(flat_state[("encoder", "conv_norm_out", "bias")], get_pt_tensor("encoder.conv_norm_out.bias", is_norm=True)) + set_val(flat_state[("encoder", "conv_out", "kernel")], get_conv_kernel("encoder.conv_out.weight")) + set_val(flat_state[("encoder", "conv_out", "bias")], get_pt_tensor("encoder.conv_out.bias")) + + # ========================================================================= + # 2. QUANT CONV & POST QUANT CONV + # ========================================================================= + set_val(flat_state[("quant_conv", "kernel")], get_conv_kernel("quant_conv.weight")) + set_val(flat_state[("quant_conv", "bias")], get_pt_tensor("quant_conv.bias")) + set_val(flat_state[("post_quant_conv", "kernel")], get_conv_kernel("post_quant_conv.weight")) + set_val(flat_state[("post_quant_conv", "bias")], get_pt_tensor("post_quant_conv.bias")) + + # ========================================================================= + # 3. DECODER + # ========================================================================= + set_val(flat_state[("decoder", "conv_in", "kernel")], get_conv_kernel("decoder.conv_in.weight")) + set_val(flat_state[("decoder", "conv_in", "bias")], get_pt_tensor("decoder.conv_in.bias")) + + for r_idx in [0, 1]: + res_pt = f"decoder.mid_block.resnets.{r_idx}" + res_path = ("decoder", "mid_block", f"resnets_{r_idx}") + set_val(flat_state[res_path + ("norm1", "scale")], get_pt_tensor(f"{res_pt}.norm1.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm1", "bias")], get_pt_tensor(f"{res_pt}.norm1.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv1", "kernel")], get_conv_kernel(f"{res_pt}.conv1.weight")) + set_val(flat_state[res_path + ("conv1", "bias")], get_pt_tensor(f"{res_pt}.conv1.bias")) + set_val(flat_state[res_path + ("norm2", "scale")], get_pt_tensor(f"{res_pt}.norm2.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm2", "bias")], get_pt_tensor(f"{res_pt}.norm2.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv2", "kernel")], get_conv_kernel(f"{res_pt}.conv2.weight")) + set_val(flat_state[res_path + ("conv2", "bias")], get_pt_tensor(f"{res_pt}.conv2.bias")) + + dec_attn_pt = "decoder.mid_block.attentions.0" + dec_attn_path = ("decoder", "mid_block", "attentions_0") + set_val( + flat_state[dec_attn_path + ("group_norm", "scale")], get_pt_tensor(f"{dec_attn_pt}.group_norm.weight", is_norm=True) + ) + set_val(flat_state[dec_attn_path + ("group_norm", "bias")], get_pt_tensor(f"{dec_attn_pt}.group_norm.bias", is_norm=True)) + set_val(flat_state[dec_attn_path + ("to_q", "kernel")], get_linear_kernel(f"{dec_attn_pt}.to_q.weight")) + set_val(flat_state[dec_attn_path + ("to_q", "bias")], get_pt_tensor(f"{dec_attn_pt}.to_q.bias")) + set_val(flat_state[dec_attn_path + ("to_k", "kernel")], get_linear_kernel(f"{dec_attn_pt}.to_k.weight")) + set_val(flat_state[dec_attn_path + ("to_k", "bias")], get_pt_tensor(f"{dec_attn_pt}.to_k.bias")) + set_val(flat_state[dec_attn_path + ("to_v", "kernel")], get_linear_kernel(f"{dec_attn_pt}.to_v.weight")) + set_val(flat_state[dec_attn_path + ("to_v", "bias")], get_pt_tensor(f"{dec_attn_pt}.to_v.bias")) + set_val(flat_state[dec_attn_path + ("to_out", "kernel")], get_linear_kernel(f"{dec_attn_pt}.to_out.0.weight")) + set_val(flat_state[dec_attn_path + ("to_out", "bias")], get_pt_tensor(f"{dec_attn_pt}.to_out.0.bias")) + + for b_idx in range(4): + up_block_pt = f"decoder.up_blocks.{b_idx}" + for r_idx in range(3): + res_pt = f"{up_block_pt}.resnets.{r_idx}" + res_path = ("decoder", "up_blocks", b_idx, "resnets", r_idx) + + set_val(flat_state[res_path + ("norm1", "scale")], get_pt_tensor(f"{res_pt}.norm1.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm1", "bias")], get_pt_tensor(f"{res_pt}.norm1.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv1", "kernel")], get_conv_kernel(f"{res_pt}.conv1.weight")) + set_val(flat_state[res_path + ("conv1", "bias")], get_pt_tensor(f"{res_pt}.conv1.bias")) + + set_val(flat_state[res_path + ("norm2", "scale")], get_pt_tensor(f"{res_pt}.norm2.weight", is_norm=True)) + set_val(flat_state[res_path + ("norm2", "bias")], get_pt_tensor(f"{res_pt}.norm2.bias", is_norm=True)) + set_val(flat_state[res_path + ("conv2", "kernel")], get_conv_kernel(f"{res_pt}.conv2.weight")) + set_val(flat_state[res_path + ("conv2", "bias")], get_pt_tensor(f"{res_pt}.conv2.bias")) + + shortcut_key = f"{res_pt}.conv_shortcut.weight" + if shortcut_key in pt_state_dict: + set_val(flat_state[res_path + ("conv_shortcut", "kernel")], get_conv_kernel(shortcut_key)) + set_val(flat_state[res_path + ("conv_shortcut", "bias")], get_pt_tensor(f"{res_pt}.conv_shortcut.bias")) + + if b_idx < 3: + ups_pt = f"{up_block_pt}.upsamplers.0.conv" + ups_path = ("decoder", "up_blocks", b_idx, "upsamplers_0", "conv") + set_val(flat_state[ups_path + ("kernel",)], get_conv_kernel(f"{ups_pt}.weight")) + set_val(flat_state[ups_path + ("bias",)], get_pt_tensor(f"{ups_pt}.bias")) + + set_val(flat_state[("decoder", "conv_norm_out", "scale")], get_pt_tensor("decoder.conv_norm_out.weight", is_norm=True)) + set_val(flat_state[("decoder", "conv_norm_out", "bias")], get_pt_tensor("decoder.conv_norm_out.bias", is_norm=True)) + set_val(flat_state[("decoder", "conv_out", "kernel")], get_conv_kernel("decoder.conv_out.weight")) + set_val(flat_state[("decoder", "conv_out", "bias")], get_pt_tensor("decoder.conv_out.bias")) + + # Update nnx_vae state + nnx.update(nnx_vae, nnx.from_flat_state(flat_state)) + + # Extract Batch Normalization running stats + bn_mean = jnp.array(get_pt_tensor("bn.running_mean")).reshape(1, -1, 1, 1) + bn_var = jnp.array(get_pt_tensor("bn.running_var")).reshape(1, -1, 1, 1) + batch_norm_eps = 0.0001 + bn_std = jnp.sqrt(bn_var + batch_norm_eps) + + return bn_mean, bn_std diff --git a/src/maxdiffusion/models/normalization_flax.py b/src/maxdiffusion/models/normalization_flax.py index abe63f1db..9f6463754 100644 --- a/src/maxdiffusion/models/normalization_flax.py +++ b/src/maxdiffusion/models/normalization_flax.py @@ -187,11 +187,13 @@ def __init__( rngs: nnx.Rngs, embedding_dim: int, eps: float = 1e-6, + scale_shift_order: str = "shift_scale", dtype: jnp.dtype = jnp.float32, weights_dtype: jnp.dtype = jnp.float32, ): self.embedding_dim = embedding_dim self.eps = eps + self.scale_shift_order = scale_shift_order self.dtype = dtype self.layer_norm = nnx.LayerNorm( num_features=embedding_dim, epsilon=eps, use_bias=False, use_scale=False, dtype=dtype, rngs=rngs @@ -199,7 +201,8 @@ def __init__( self.linear = nnx.Linear( in_features=embedding_dim, out_features=embedding_dim * 2, - use_bias=True, + use_bias=False, + kernel_init=nnx.with_partitioning(nnx.initializers.lecun_normal(), ("embed", "mlp")), dtype=dtype, param_dtype=weights_dtype, rngs=rngs, @@ -207,7 +210,12 @@ def __init__( def __call__(self, x: jax.Array, conditioning_embedding: jax.Array) -> jax.Array: emb = self.linear(jax.nn.silu(conditioning_embedding)) - scale, shift = jnp.split(emb, 2, axis=-1) + if self.scale_shift_order == "shift_scale": + shift, scale = jnp.split(emb, 2, axis=-1) + else: + scale, shift = jnp.split(emb, 2, axis=-1) + shift = nn.with_logical_constraint(shift, ("activation_batch", "activation_embed")) + scale = nn.with_logical_constraint(scale, ("activation_batch", "activation_embed")) x_norm = self.layer_norm(x) return (1.0 + scale[:, None, :]) * x_norm + shift[:, None, :] diff --git a/src/maxdiffusion/models/qwen3_flax.py b/src/maxdiffusion/models/qwen3_flax.py index b3ca43003..af0a82344 100644 --- a/src/maxdiffusion/models/qwen3_flax.py +++ b/src/maxdiffusion/models/qwen3_flax.py @@ -15,7 +15,7 @@ """ import math -from typing import Any, List, Optional, Tuple +from typing import Any, Dict, List, Optional, Tuple from flax import nnx import flax.linen as nn import jax @@ -41,6 +41,13 @@ def __init__( rope_theta: float = 1000000.0, max_position_embeddings: int = 40960, dtype=jnp.float32, + attention_kernel: str = "dot_product", + flash_block_sizes: Optional[Dict[str, int]] = None, + mesh: Optional[jax.sharding.Mesh] = None, + ulysses_shards: int = -1, + ulysses_attention_chunks: int = 1, + max_layer_to_run: Optional[int] = 27, + is_causal: bool = True, ): self.vocab_size = vocab_size self.hidden_size = hidden_size @@ -53,6 +60,13 @@ def __init__( self.rope_theta = rope_theta self.max_position_embeddings = max_position_embeddings self.dtype = dtype + self.attention_kernel = attention_kernel + self.flash_block_sizes = flash_block_sizes + self.mesh = mesh + self.ulysses_shards = ulysses_shards + self.ulysses_attention_chunks = ulysses_attention_chunks + self.max_layer_to_run = max_layer_to_run + self.is_causal = is_causal # ----------------------------------------------------------------------------- @@ -154,6 +168,9 @@ def rotate_half(x): return q_rot, k_rot +from maxdiffusion.models.attention_flax import AttentionOp + + # ----------------------------------------------------------------------------- # Self Attention (Grouped Query Attention) # ----------------------------------------------------------------------------- @@ -247,33 +264,53 @@ def __call__( k = jnp.repeat(k, gqa_ratio, axis=-2) v = jnp.repeat(v, gqa_ratio, axis=-2) - # 5. Transpose to (batch, num_heads, seq_len, head_dim) for attention - q = jnp.transpose(q, (0, 2, 1, 3)) - k = jnp.transpose(k, (0, 2, 1, 3)) - v = jnp.transpose(v, (0, 2, 1, 3)) - - # 6. Compute attention logits in float32 - q_f = q.astype(jnp.float32) - k_f = k.astype(jnp.float32) - v_f = v.astype(jnp.float32) - - scores = jnp.matmul(q_f, jnp.transpose(k_f, (0, 1, 3, 2))) / math.sqrt(self.config.head_dim) - - # 7. Apply causal attention mask - causal_mask = jnp.tril(jnp.ones((seq_len, seq_len), dtype=jnp.bool_)) - scores = jnp.where(causal_mask, scores, -1e4) - - # 8. Apply padding attention mask if provided - if attention_mask is not None: - p_mask = attention_mask[:, jnp.newaxis, jnp.newaxis, :].astype(jnp.bool_) - scores = jnp.where(p_mask, scores, -1e4) - - # 9. Softmax & Weighted Sum in float32 - probs = jax.nn.softmax(scores, axis=-1) - out = jnp.matmul(probs, v_f).astype(self.config.dtype) + if self.config.attention_kernel != "dot_product": + scale = self.config.head_dim**-0.5 + q_3d = q.reshape((batch_size, seq_len, self.config.num_attention_heads * self.config.head_dim)) + k_3d = k.reshape((batch_size, seq_len, self.config.num_attention_heads * self.config.head_dim)) + v_3d = v.reshape((batch_size, seq_len, self.config.num_attention_heads * self.config.head_dim)) + attn_op = AttentionOp( + mesh=self.config.mesh, + attention_kernel=self.config.attention_kernel, + scale=scale, + heads=self.config.num_attention_heads, + dim_head=self.config.head_dim, + flash_min_seq_length=128, + flash_block_sizes=self.config.flash_block_sizes, + dtype=self.config.dtype, + ulysses_shards=self.config.ulysses_shards, + ulysses_attention_chunks=self.config.ulysses_attention_chunks, + is_causal=getattr(self.config, "is_causal", True), + ) + out = attn_op.apply_attention(q_3d, k_3d, v_3d, attention_mask=attention_mask) + else: + # 5. Transpose to (batch, num_heads, seq_len, head_dim) for attention + q = jnp.transpose(q, (0, 2, 1, 3)) + k = jnp.transpose(k, (0, 2, 1, 3)) + v = jnp.transpose(v, (0, 2, 1, 3)) + + # 6. Compute attention logits in float32 + q_f = q.astype(jnp.float32) + k_f = k.astype(jnp.float32) + v_f = v.astype(jnp.float32) + + scores = jnp.matmul(q_f, jnp.transpose(k_f, (0, 1, 3, 2))) / math.sqrt(self.config.head_dim) + + # 7. Apply causal attention mask if configured + if getattr(self.config, "is_causal", True): + causal_mask = jnp.tril(jnp.ones((seq_len, seq_len), dtype=jnp.bool_)) + scores = jnp.where(causal_mask, scores, -1e4) + + # 8. Apply padding attention mask if provided + if attention_mask is not None: + p_mask = attention_mask[:, jnp.newaxis, jnp.newaxis, :].astype(jnp.bool_) + scores = jnp.where(p_mask, scores, -1e4) + + # 9. Softmax & Weighted Sum in float32 + probs = jax.nn.softmax(scores, axis=-1) + out = jnp.matmul(probs, v_f).astype(self.config.dtype) + out = jnp.transpose(out, (0, 2, 1, 3)).reshape((batch_size, seq_len, -1)) - # 10. Reshape back and project out: (batch, seq_len, hidden_size) - out = jnp.transpose(out, (0, 2, 1, 3)).reshape((batch_size, seq_len, -1)) return o_proj(out) @@ -380,7 +417,12 @@ def __call__( ) # 3. Stacked Decoder Layers - for i in range(self.config.num_hidden_layers): + num_layers_to_exec = ( + self.config.num_hidden_layers + if self.config.max_layer_to_run is None + else min(self.config.num_hidden_layers, self.config.max_layer_to_run + 1) + ) + for i in range(num_layers_to_exec): layer = FlaxQwen3DecoderLayer( config=self.config, name=f"layers_{i}", diff --git a/src/maxdiffusion/models/resnet_flax.py b/src/maxdiffusion/models/resnet_flax.py index 79ddcb30e..8371a4432 100644 --- a/src/maxdiffusion/models/resnet_flax.py +++ b/src/maxdiffusion/models/resnet_flax.py @@ -57,9 +57,8 @@ def setup(self): @nn.compact def __call__(self, hidden_states): batch, height, width, channels = hidden_states.shape - hidden_states = jax.image.resize( - hidden_states, shape=(batch, height * 2, width * 2, channels), method="nearest", precision=self.precision - ) + hidden_states = jnp.broadcast_to(hidden_states[:, :, None, :, None, :], (batch, height, 2, width, 2, channels)) + hidden_states = jnp.reshape(hidden_states, (batch, height * 2, width * 2, channels)) hidden_states = nn.with_logical_constraint(hidden_states, ("conv_batch", "height", "keep_2", "out_channels")) diff --git a/src/maxdiffusion/models/vae_flax.py b/src/maxdiffusion/models/vae_flax.py index 72adcbe79..af13327bf 100644 --- a/src/maxdiffusion/models/vae_flax.py +++ b/src/maxdiffusion/models/vae_flax.py @@ -87,11 +87,8 @@ def setup(self): def __call__(self, hidden_states): batch, height, width, channels = hidden_states.shape - hidden_states = jax.image.resize( - hidden_states, - shape=(batch, height * 2, width * 2, channels), - method="nearest", - ) + hidden_states = jnp.broadcast_to(hidden_states[:, :, None, :, None, :], (batch, height, 2, width, 2, channels)) + hidden_states = jnp.reshape(hidden_states, (batch, height * 2, width * 2, channels)) hidden_states = self.conv(hidden_states) return hidden_states diff --git a/src/maxdiffusion/pipelines/flux/__init__.py b/src/maxdiffusion/pipelines/flux/__init__.py index 39ea05b57..c94ebd8a3 100644 --- a/src/maxdiffusion/pipelines/flux/__init__.py +++ b/src/maxdiffusion/pipelines/flux/__init__.py @@ -19,3 +19,6 @@ from .flux_pipeline import ( FluxPipeline, ) +from .flux2klein_pipeline import ( + FlaxFlux2KleinPipeline, +) diff --git a/src/maxdiffusion/pipelines/flux/flux2klein_pipeline.py b/src/maxdiffusion/pipelines/flux/flux2klein_pipeline.py index 634ec8d9e..8ff868a0b 100644 --- a/src/maxdiffusion/pipelines/flux/flux2klein_pipeline.py +++ b/src/maxdiffusion/pipelines/flux/flux2klein_pipeline.py @@ -27,18 +27,23 @@ import numpy as np from flax.linen import partitioning as nn_partitioning +from flax import nnx from maxdiffusion import max_logging from maxdiffusion.max_utils import device_put_replicated from ..pipeline_flax_utils import FlaxDiffusionPipeline -from ...models.flux.transformers.transformer_flux_flax import Flux2KleinTransformer2DModel -from ...models.vae_flax import FlaxAutoencoderKL +from ...models.flux.transformers.transformer_flux_flax import ( + Flux2KleinTransformer2DModel, + NNXFlux2KleinTransformer2DModel, +) +from ...models.vae_flax import FlaxAutoencoderKL, FlaxDecoderOutput from ...models.qwen3_flax import FlaxQwen3Model from ...schedulers.scheduling_flow_match_flax import FlaxFlowMatchScheduler, compute_empirical_mu from ...models.flux.util import ( pack_latents, - unpack_latents, + patchify_latents, prepare_latent_image_ids, + prepare_multi_image_ids, prepare_text_ids, ) @@ -51,7 +56,7 @@ class FlaxFlux2KleinPipeline(FlaxDiffusionPipeline): def __init__( self, - transformer: Flux2KleinTransformer2DModel, + transformer: Union[Flux2KleinTransformer2DModel, NNXFlux2KleinTransformer2DModel], vae: FlaxAutoencoderKL, text_encoder: FlaxQwen3Model, tokenizer, @@ -69,11 +74,40 @@ def __init__( scheduler=scheduler, ) self._config = config + max_layer = getattr(config, "text_encoder_max_layer", 27) + if max_layer is not None and max_layer < 27: + raise ValueError( + f"Invalid configuration `text_encoder_max_layer={max_layer}`. " + f"FLUX.2-Klein requires extracting intermediate prompt embeddings from Qwen3 layers 9, 18, and 27, " + f"so `text_encoder_max_layer` must be >= 27." + ) self.mesh = mesh + self.tokenizer = tokenizer + if self.tokenizer is None: + tokenizer_path = getattr(config, "tokenizer_model_name_or_path", None) or getattr( + config, "pretrained_model_name_or_path", "" + ) + hf_home = os.environ.get("HF_HOME", os.path.expanduser("~/.cache/huggingface")) + repo_cache = os.path.join( + hf_home, + "hub", + f"models--{getattr(config, 'pretrained_model_name_or_path', '').replace('/', '--')}", + "snapshots", + ) + if os.path.exists(repo_cache) and os.listdir(repo_cache): + tokenizer_path = os.path.join(repo_cache, os.listdir(repo_cache)[0]) + + from transformers import Qwen2TokenizerFast + + try: + self.tokenizer = Qwen2TokenizerFast.from_pretrained(tokenizer_path, local_files_only=True) + except Exception: + self.tokenizer = Qwen2TokenizerFast.from_pretrained(tokenizer_path, subfolder="tokenizer", local_files_only=True) # JIT compilation cache self._jitted_qwen3_forward = None self._jitted_transformer_step = None + self._jitted_vae_encode = None self._jitted_vae_decode = None def _setup_jit_functions(self): @@ -82,29 +116,304 @@ def _setup_jit_functions(self): @jax.jit def qwen3_forward(q_params, ids, mask): - return self.text_encoder.apply({"params": q_params}, input_ids=ids, attention_mask=mask) + _, all_hidden_states = self.text_encoder.apply({"params": q_params}, input_ids=ids, attention_mask=mask) + h_9 = all_hidden_states[9] + h_18 = all_hidden_states[18] + h_27 = all_hidden_states[27] + out = jnp.stack([h_9, h_18, h_27], axis=1) + prompt_embeds = jnp.transpose(out, (0, 2, 1, 3)).reshape((ids.shape[0], ids.shape[1], -1)) + context_spec = P(None, "context") if "context" in self.mesh.axis_names and self.mesh.shape["context"] > 1 else P() + prompt_embeds = jax.lax.with_sharding_constraint(prompt_embeds, jax.sharding.NamedSharding(self.mesh, context_spec)) + return prompt_embeds + + if isinstance(self.vae, nnx.Module): + v_graph, _, v_rest = nnx.split(self.vae, nnx.Param, ...) + + @jax.jit + def vae_encode(v_params, img): + merged = nnx.merge(v_graph, v_params, v_rest) + return merged.encode(img) + + @jax.jit(static_argnums=(4, 5), donate_argnums=(1,)) + def vae_decode(v_params, latents_packed, vae_bn_mean, vae_bn_std, height, width): + batch_size_val = latents_packed.shape[0] + h_latent = height // 8 + w_latent = width // 8 + + vae_bn_mean_seq = vae_bn_mean.reshape(1, 1, 128) + vae_bn_std_seq = vae_bn_std.reshape(1, 1, 128) + + latents_bn = latents_packed * vae_bn_std_seq + vae_bn_mean_seq + latents_unpacked = jnp.reshape(latents_bn, (batch_size_val, h_latent // 2, w_latent // 2, 32, 2, 2)) + latents_unpacked = jnp.transpose(latents_unpacked, (0, 3, 1, 4, 2, 5)) + latents_unpacked = jnp.reshape(latents_unpacked, (batch_size_val, 32, h_latent, w_latent)) + + merged = nnx.merge(v_graph, v_params, v_rest) + res = merged.decode(latents_unpacked) + return FlaxDecoderOutput(sample=res) - @jax.jit - def transformer_step(t_params, latents, img_ids, prompt_embeds, txt_ids, vec, timestep, guidance): - return self.transformer.apply( - {"params": t_params}, - hidden_states=latents, - img_ids=img_ids, - encoder_hidden_states=prompt_embeds, - txt_ids=txt_ids, - pooled_projections=vec, - timestep=timestep, - guidance=guidance, - ) + else: - @jax.jit - def vae_decode(v_params, latents_unpatched): - return self.vae.apply({"params": v_params}, latents=latents_unpatched, method=self.vae.decode) + @jax.jit + def vae_encode(v_params, img): + # FlaxAutoencoderKL expects (B, 3, H, W) + res = self.vae.apply({"params": v_params}, sample=img, method=self.vae.encode) + moments = res.latent_dist.mode() + return jnp.transpose(moments, (0, 3, 1, 2)) + + @jax.jit(static_argnums=(4, 5), donate_argnums=(1,)) + def vae_decode(v_params, latents_packed, vae_bn_mean, vae_bn_std, height, width): + batch_size_val = latents_packed.shape[0] + h_latent = height // 8 + w_latent = width // 8 + + vae_bn_mean_seq = vae_bn_mean.reshape(1, 1, 128) + vae_bn_std_seq = vae_bn_std.reshape(1, 1, 128) + + latents_bn = latents_packed * vae_bn_std_seq + vae_bn_mean_seq + latents_unpacked = jnp.reshape(latents_bn, (batch_size_val, h_latent // 2, w_latent // 2, 32, 2, 2)) + latents_unpacked = jnp.transpose(latents_unpacked, (0, 3, 1, 4, 2, 5)) + latents_unpacked = jnp.reshape(latents_unpacked, (batch_size_val, 32, h_latent, w_latent)) + + res = self.vae.apply({"params": v_params}, latents=latents_unpacked, method=self.vae.decode) + return FlaxDecoderOutput(sample=res.sample) + + if isinstance(self.transformer, nnx.Module): + g, nnx_state, r = nnx.split(self.transformer, nnx.Param, ...) + + @jax.jit + def transformer_step(t_params, latents, img_ids, prompt_embeds, txt_ids, vec, timestep, guidance): + nnx_merged = nnx.merge(g, t_params, r) + return nnx_merged( + hidden_states=latents, + encoder_hidden_states=prompt_embeds, + pooled_projections=vec, + timestep=timestep, + img_ids=img_ids, + txt_ids=txt_ids, + guidance=guidance, + return_dict=True, + ) + + @jax.jit(static_argnums=(9,)) + def fused_denoise_loop( + t_params, latents, img_ids, prompt_embeds, txt_ids, vec, timesteps, sigmas, guidance, target_len=None + ): + sigmas_padded = jnp.concatenate([sigmas, jnp.array([0.0], dtype=sigmas.dtype)]) + nnx_merged = nnx.merge(g, t_params, r) + + def scan_body(cur_latents, step_idx): + t_val = timesteps[step_idx] + t_vec = jnp.broadcast_to(t_val / 1000.0, (cur_latents.shape[0],)) + model_output = nnx_merged( + hidden_states=cur_latents, + img_ids=img_ids, + encoder_hidden_states=prompt_embeds, + txt_ids=txt_ids, + pooled_projections=vec, + timestep=t_vec, + guidance=guidance, + return_dict=True, + ) + sigma = sigmas_padded[step_idx] + sigma_next = sigmas_padded[step_idx + 1] + dt = sigma_next - sigma + v = model_output.sample + if target_len is not None and cur_latents.shape[1] > target_len: + target_latents = cur_latents[:, :target_len, :] + v_target = v[:, :target_len, :] + next_target = target_latents + v_target * dt + prev_sample = jnp.concatenate([next_target, cur_latents[:, target_len:, :]], axis=1) + else: + prev_sample = cur_latents + v * dt + return prev_sample, None + + steps = jnp.arange(timesteps.shape[0]) + final_latents, _ = jax.lax.scan(scan_body, latents, steps) + return final_latents + + else: + + @jax.jit + def transformer_step(t_params, latents, img_ids, prompt_embeds, txt_ids, vec, timestep, guidance): + return self.transformer.apply( + {"params": t_params}, + hidden_states=latents, + img_ids=img_ids, + encoder_hidden_states=prompt_embeds, + txt_ids=txt_ids, + pooled_projections=vec, + timestep=timestep, + guidance=guidance, + ) + + @jax.jit(static_argnums=(9,)) + def fused_denoise_loop( + t_params, latents, img_ids, prompt_embeds, txt_ids, vec, timesteps, sigmas, guidance, target_len=None + ): + sigmas_padded = jnp.concatenate([sigmas, jnp.array([0.0], dtype=sigmas.dtype)]) + + def scan_body(cur_latents, step_idx): + t_val = timesteps[step_idx] + t_vec = jnp.broadcast_to(t_val / 1000.0, (cur_latents.shape[0],)) + model_output = self.transformer.apply( + {"params": t_params}, + hidden_states=cur_latents, + img_ids=img_ids, + encoder_hidden_states=prompt_embeds, + txt_ids=txt_ids, + pooled_projections=vec, + timestep=t_vec, + guidance=guidance, + ) + sigma = sigmas_padded[step_idx] + sigma_next = sigmas_padded[step_idx + 1] + dt = sigma_next - sigma + v = model_output.sample + if target_len is not None and cur_latents.shape[1] > target_len: + target_latents = cur_latents[:, :target_len, :] + v_target = v[:, :target_len, :] + next_target = target_latents + v_target * dt + prev_sample = jnp.concatenate([next_target, cur_latents[:, target_len:, :]], axis=1) + else: + prev_sample = cur_latents + v * dt + return prev_sample, None + + steps = jnp.arange(timesteps.shape[0]) + final_latents, _ = jax.lax.scan(scan_body, latents, steps) + return final_latents self._jitted_qwen3_forward = qwen3_forward self._jitted_transformer_step = transformer_step + self._jitted_fused_denoise_loop = fused_denoise_loop + self._jitted_vae_encode = vae_encode self._jitted_vae_decode = vae_decode + def _get_dynamic_batch_sharding(self): + """Dynamically infers the batch dimension sharding specification from self.mesh.""" + batch_axes = [axis for axis in ("data", "fsdp") if axis in self.mesh.axis_names and self.mesh.shape[axis] > 1] + spec = P(tuple(batch_axes)) if batch_axes else P(None) + return jax.sharding.NamedSharding(self.mesh, spec) + + def compile_aot_async( + self, + params, + vae_params, + qwen3_params, + vae_bn_mean, + vae_bn_std, + batch_size=1, + height=1024, + width=1024, + images=None, + image=None, + num_conditioning_images=0, + ): + """Triggers AOT compilation for Qwen3, Flux Transformer, and VAE concurrently using ThreadPoolExecutor.""" + self._setup_jit_functions() + max_logging.log("🚀 Pre-compiling XLA graphs for Qwen3, Flux Transformer, and VAE concurrently...") + from concurrent.futures import ThreadPoolExecutor + + if images is None and image is not None: + images = image if isinstance(image, list) else [image] + + if images is not None and len(images) > 0: + num_conditioning_images = len(images) + + seq_len_img = (height // 16) * (width // 16) + total_img_len = (1 + num_conditioning_images) * seq_len_img + seq_len_txt = self._config.max_sequence_length + + dummy_ids = jnp.zeros((batch_size, seq_len_txt), dtype=jnp.int32) + dummy_mask = jnp.ones((batch_size, seq_len_txt), dtype=jnp.int32) + + dummy_latents = jnp.zeros((batch_size, total_img_len, 128), dtype=jnp.float32) + dummy_img_ids = jnp.zeros((batch_size, total_img_len, 4), dtype=jnp.int32) + dummy_prompt_embeds = jnp.zeros((batch_size, seq_len_txt, self.transformer.joint_attention_dim), dtype=jnp.bfloat16) + dummy_txt_ids = jnp.zeros((batch_size, seq_len_txt, 4), dtype=jnp.float32) + dummy_t_vec = jnp.zeros((batch_size,), dtype=jnp.float32) + + dummy_target_latents = jnp.zeros((batch_size, seq_len_img, 128), dtype=jnp.float32) + dummy_bn_mean = jnp.array(vae_bn_mean, dtype=jnp.float32) + dummy_bn_std = jnp.array(vae_bn_std, dtype=jnp.float32) + + data_sharding = self._get_dynamic_batch_sharding() + replicated_sharding = jax.sharding.NamedSharding(self.mesh, P()) + context_sharding = jax.sharding.NamedSharding(self.mesh, P(None, "context")) + + def put_data_on_devices(x, sharding): + if isinstance(x, jax.Array) and hasattr(x, "sharding") and not x.sharding.is_fully_addressable: + return x + if hasattr(sharding, "is_fully_addressable") and sharding.is_fully_addressable: + return jax.device_put(x, sharding) + return device_put_replicated(x, sharding) + + dummy_ids = put_data_on_devices(dummy_ids, data_sharding) + dummy_mask = put_data_on_devices(dummy_mask, data_sharding) + dummy_latents = put_data_on_devices(dummy_latents, data_sharding) + dummy_img_ids = put_data_on_devices(dummy_img_ids, data_sharding) + dummy_prompt_embeds = put_data_on_devices(dummy_prompt_embeds, context_sharding) + dummy_txt_ids = put_data_on_devices(dummy_txt_ids, data_sharding) + dummy_t_vec = put_data_on_devices(dummy_t_vec, data_sharding) + dummy_target_latents = put_data_on_devices(dummy_target_latents, data_sharding) + dummy_bn_mean = put_data_on_devices(dummy_bn_mean, replicated_sharding) + dummy_bn_std = put_data_on_devices(dummy_bn_std, replicated_sharding) + + def compile_qwen3(): + t0 = time.perf_counter() + with self.mesh, nn_partitioning.axis_rules(self._config.logical_axis_rules): + self._jitted_qwen3_forward.lower(qwen3_params, dummy_ids, dummy_mask).compile() + max_logging.log(f" -> [AOT COMPILED] Qwen3 Text Encoder in {time.perf_counter() - t0:.2f}s") + + num_steps = getattr(self._config, "num_inference_steps", 4) + dummy_timesteps = put_data_on_devices(jnp.zeros((num_steps,), dtype=jnp.float32), replicated_sharding) + dummy_sigmas = put_data_on_devices(jnp.zeros((num_steps + 1,), dtype=jnp.float32), replicated_sharding) + + def compile_transformer(): + t0 = time.perf_counter() + with self.mesh, nn_partitioning.axis_rules(self._config.logical_axis_rules): + self._jitted_fused_denoise_loop.lower( + params, + dummy_latents, + dummy_img_ids, + dummy_prompt_embeds, + dummy_txt_ids, + None, + dummy_timesteps, + dummy_sigmas, + None, + seq_len_img, + ).compile() + max_logging.log(f" -> [AOT COMPILED] Fused Flux Transformer Denoise Scan in {time.perf_counter() - t0:.2f}s") + + def compile_vae(): + t0 = time.perf_counter() + with self.mesh, nn_partitioning.axis_rules(self._config.logical_axis_rules): + self._jitted_vae_decode.lower(vae_params, dummy_target_latents, dummy_bn_mean, dummy_bn_std, height, width).compile() + max_logging.log(f" -> [AOT COMPILED] VAE Decoder in {time.perf_counter() - t0:.2f}s") + + def compile_vae_encode(): + t0 = time.perf_counter() + dummy_rgb = jnp.zeros((1, 3, height, width), dtype=jnp.float32) + with self.mesh, nn_partitioning.axis_rules(self._config.logical_axis_rules): + self._jitted_vae_encode.lower(vae_params, dummy_rgb).compile() + max_logging.log(f" -> [AOT COMPILED] VAE Encoder in {time.perf_counter() - t0:.2f}s") + + t_start = time.perf_counter() + with ThreadPoolExecutor(max_workers=4) as executor: + futures = [ + executor.submit(compile_qwen3), + executor.submit(compile_transformer), + executor.submit(compile_vae), + ] + if num_conditioning_images > 0 or (images is not None and len(images) > 0): + futures.append(executor.submit(compile_vae_encode)) + for future in futures: + future.result() + aot_duration = time.perf_counter() - t_start + max_logging.log(f"⚡ [AOT CONCURRENT COMPILATION COMPLETE] Total AOT compile time: {aot_duration:.2f}s") + return aot_duration + def _prepare_latents(self, config, batch_size, height, width): num_channels_latents = 32 latent_height = height // 8 @@ -144,15 +453,22 @@ def __call__( width: int = 1024, num_inference_steps: int = 4, batch_size: int = 1, + images: Optional[List[Any]] = None, + image: Optional[Union[Any, List[Any]]] = None, use_latents: bool = False, latents: Optional[Any] = None, measure_time: bool = False, + warmup: bool = False, output_dir: str = "output/", output_name: str = "flux2klein_generated_image.png", + profile_target: Optional[str] = None, ): # 1. Setup JIT functions self._setup_jit_functions() + if images is None and image is not None: + images = image if isinstance(image, list) else [image] + # 2. Setup prompts and inputs if isinstance(prompt, str): prompts = [prompt] * batch_size @@ -170,6 +486,8 @@ def __call__( if C == 32: max_logging.log(" [PIPELINE] Unpacked 32-channel latents detected. Packing using pack_latents...") latents_jax = pack_latents(latents_jax) + elif C == 128: + latents_jax = jnp.transpose(jnp.reshape(latents_jax, (B, C, H * W)), (0, 2, 1)) else: latents_jax = jnp.transpose(jnp.reshape(latents_jax, (B, C, H * W)), (0, 2, 1)) else: @@ -179,7 +497,9 @@ def __call__( # RoPE position IDs txt_ids_val = prepare_text_ids(batch_size, seq_len_txt) - img_ids_val = prepare_latent_image_ids(batch_size, height // 16, width // 16) + target_img_ids_val = prepare_latent_image_ids(batch_size, height // 16, width // 16) + t_pipeline_start = time.perf_counter() + trace = {} # Scheduler mu = compute_empirical_mu(seq_len_img, num_inference_steps) @@ -192,67 +512,150 @@ def __call__( sigmas=sigmas_custom, ) - trace = {} - with self.mesh, nn_partitioning.axis_rules(self._config.logical_axis_rules): proc_id = jax.process_index() proc_cnt = jax.process_count() host_prefix = f"[HOST {proc_id}/{proc_cnt}] " + # Shard pipeline batch inputs across data axis ("data") for SPMD multi-host execution + data_sharding = jax.sharding.NamedSharding(self.mesh, P("data")) + + def put_data_on_devices(x, sharding): + if isinstance(x, jax.Array) and hasattr(x, "sharding") and not x.sharding.is_fully_addressable: + return x + if hasattr(sharding, "is_fully_addressable") and sharding.is_fully_addressable: + return jax.device_put(x, sharding) + return device_put_replicated(x, sharding) + # --------------------------------------------------------------------- - # PHASE A: Encode Prompt (Qwen3) + # PHASE 0: Encode Reference Images (VAE) # --------------------------------------------------------------------- - print(f"{host_prefix} [PHASE A] Encoding {len(prompts)} prompt(s) using JAX Qwen3 on TPU...", flush=True) - t0 = time.perf_counter() - - try: - # Resolve tokenizer path from config - tokenizer_path = self._config.tokenizer_model_name_or_path - hf_home = os.environ.get("HF_HOME", os.path.expanduser("~/.cache/huggingface")) - repo_cache = os.path.join( - hf_home, "hub", f"models--{self._config.pretrained_model_name_or_path.replace('/', '--')}", "snapshots" - ) - if os.path.exists(repo_cache) and os.listdir(repo_cache): - tokenizer_path = os.path.join(repo_cache, os.listdir(repo_cache)[0]) + if images is not None and len(images) > 0: + t0_vae_enc_start = time.perf_counter() + trace["start_to_vae_encode"] = t0_vae_enc_start - t_pipeline_start + max_logging.log(f"{host_prefix} [PHASE 0] Encoding {len(images)} reference image(s) using JAX VAE encoder on TPU...") + norm_ref_latents = [] + packed_ref_latents = [] + bn_mean_arr = jnp.array(vae_bn_mean, dtype=jnp.float32) + bn_std_arr = jnp.array(vae_bn_std, dtype=jnp.float32) + + for img in images: + if isinstance(img, Image.Image): + img = img.convert("RGB").resize((width, height), Image.Resampling.BICUBIC) + arr = np.array(img, dtype=np.float32) / 127.5 - 1.0 + arr = np.transpose(arr, (2, 0, 1)) + img_tensor = jnp.expand_dims(jnp.array(arr), axis=0) + elif isinstance(img, np.ndarray): + if img.ndim == 3: + img = np.expand_dims(img, axis=0) + if img.shape[-1] == 3: + img = np.transpose(img, (0, 3, 1, 2)) + if np.issubdtype(img.dtype, np.integer): + img = img.astype(np.float32) / 127.5 - 1.0 + elif np.issubdtype(img.dtype, np.floating): + if img.max() > 1.0: + img = img / 127.5 - 1.0 + elif img.min() >= 0.0: + img = img * 2.0 - 1.0 + img_tensor = jnp.array(img, dtype=np.float32) + elif isinstance(img, jnp.ndarray): + if img.ndim == 3: + img = jnp.expand_dims(img, axis=0) + if img.shape[-1] == 3: + img = jnp.transpose(img, (0, 3, 1, 2)) + if jnp.issubdtype(img.dtype, jnp.integer): + img = img.astype(jnp.float32) / 127.5 - 1.0 + elif jnp.issubdtype(img.dtype, jnp.floating): + if img.max() > 1.0: + img = img / 127.5 - 1.0 + elif img.min() >= 0.0: + img = img * 2.0 - 1.0 + img_tensor = img + else: + raise ValueError(f"Unsupported image type: {type(img)}") + + raw_ref_latents = self._jitted_vae_encode(vae_params, img_tensor) + raw_ref_latents.block_until_ready() + patchified_ref = patchify_latents(raw_ref_latents) + normalized_ref = (patchified_ref - bn_mean_arr) / bn_std_arr + norm_ref_latents.append(normalized_ref) + + packed = jnp.transpose( + jnp.reshape(normalized_ref, (normalized_ref.shape[0], normalized_ref.shape[1], -1)), (0, 2, 1) + ) + if packed.shape[0] == 1 and batch_size > 1: + packed = jnp.repeat(packed, batch_size, axis=0) + packed_ref_latents.append(packed) + + ref_img_ids_val = prepare_multi_image_ids(norm_ref_latents, scale=10) + if ref_img_ids_val.shape[0] == 1 and batch_size > 1: + ref_img_ids_val = jnp.repeat(ref_img_ids_val, batch_size, axis=0) + img_ids_val = jnp.concatenate([target_img_ids_val, ref_img_ids_val], axis=1) + latents_jax = jnp.concatenate([latents_jax] + packed_ref_latents, axis=1) + max_logging.log(f" [PIPELINE] Joint latents shape: {latents_jax.shape}, Joint img_ids shape: {img_ids_val.shape}") + + t0_vae_enc_end = time.perf_counter() + trace["vae_encode"] = t0_vae_enc_end - t0_vae_enc_start + trace["image_encoding"] = trace["vae_encode"] + max_logging.log(f" -> [TIMING] Reference Image Encoding (VAE): {trace['vae_encode']:.4f} seconds ⏱️") + else: + img_ids_val = target_img_ids_val + trace["vae_encode"] = 0.0 + trace["image_encoding"] = 0.0 + + t0_qwen3_start = time.perf_counter() + if trace.get("vae_encode", 0.0) > 0: + trace["vae_encode_to_qwen3"] = t0_qwen3_start - t0_vae_enc_end + max_logging.log(f" -> [TIMING] VAE Encode to Qwen3 Overhead: {trace['vae_encode_to_qwen3']:.4f} seconds ⏱️") + else: + trace["start_to_qwen3"] = t0_qwen3_start - t_pipeline_start + max_logging.log(f" -> [TIMING] Start to Qwen3: {trace['start_to_qwen3']:.4f} seconds ⏱️") - from transformers import Qwen2TokenizerFast + # --------------------------------------------------------------------- + # PHASE A: Encode Prompt (Qwen3) + # --------------------------------------------------------------------- + if not prompts: + raise ValueError("Prompt must be provided to FlaxFlux2KleinPipeline") + if isinstance(prompts, str): + prompts = [prompts] - try: - tokenizer = Qwen2TokenizerFast.from_pretrained(tokenizer_path, local_files_only=True) - except Exception: - tokenizer = Qwen2TokenizerFast.from_pretrained(tokenizer_path, subfolder="tokenizer", local_files_only=True) + max_logging.log(f"{host_prefix} [PHASE A] Encoding {len(prompts)} prompt(s) using JAX Qwen3 on TPU...") + try: # Tokenize using deterministic explicit template string (version-agnostic across transformers versions) templated_texts = [ f"<|im_start|>user\n{p}<|im_end|>\n<|im_start|>assistant\n\n\n\n\n" for p in prompts ] - inputs = tokenizer( + inputs = self.tokenizer( templated_texts, return_tensors="np", padding="max_length", truncation=True, max_length=seq_len_txt ) prompt_ids = jnp.array(inputs["input_ids"]) prompt_mask = jnp.array(inputs["attention_mask"]) - # Run Text Encoding - hidden_states, all_hidden_states = self._jitted_qwen3_forward(qwen3_params, prompt_ids, prompt_mask) - - # Stack layers 9, 18, 27 to form prompt embeddings - h_9 = all_hidden_states[9] - h_18 = all_hidden_states[18] - h_27 = all_hidden_states[27] - out = jnp.stack([h_9, h_18, h_27], axis=1) - # Transpose shape to [B, seq_len, 3*hidden_size] - prompt_embeds_jax = jnp.transpose(out, (0, 2, 1, 3)).reshape((batch_size, seq_len_txt, -1)) + # Run Text Encoding with sharded input arrays matching compile_aot_async + prompt_ids = put_data_on_devices(prompt_ids, data_sharding) + prompt_mask = put_data_on_devices(prompt_mask, data_sharding) + do_prof_qwen3 = profile_target in ("all", "qwen3") + if do_prof_qwen3: + tb_dir = getattr(self._config, "tensorboard_dir", "/tmp") + jax.profiler.start_trace(os.path.join(tb_dir, "profile_qwen3")) + with jax.named_scope("qwen3_text_encoder"): + prompt_embeds_jax = self._jitted_qwen3_forward(qwen3_params, prompt_ids, prompt_mask) prompt_embeds_jax.block_until_ready() + if do_prof_qwen3: + jax.profiler.stop_trace() except Exception as e: - print(f"❌ {host_prefix} EXCEPTION IN PHASE A (QWEN3 ENCODING): {e}", flush=True) + max_logging.log(f"❌ {host_prefix} EXCEPTION IN PHASE A (QWEN3 ENCODING): {e}") import traceback traceback.print_exc() sys.stdout.flush() raise e - trace["prompt_encoding"] = time.perf_counter() - t0 - max_logging.log(f" -> [TIMING] Prompt Encoding (Qwen3): {trace['prompt_encoding']:.4f} seconds ⏱️") + t0_qwen3_end = time.perf_counter() + trace["qwen3_encoding"] = t0_qwen3_end - t0_qwen3_start + trace["prompt_encoding"] = trace["qwen3_encoding"] + max_logging.log(f" -> [TIMING] Prompt Encoding (Qwen3): {trace['qwen3_encoding']:.4f} seconds ⏱️") proc_id = jax.process_index() proc_cnt = jax.process_count() @@ -260,69 +663,69 @@ def __call__( # Stage Sync 1: Phase A Complete multihost_utils.sync_global_devices("phase_a_complete") - print(f"{host_prefix} Passed Phase A Sync Barrier (phase_a_complete) successfully! ✅", flush=True) - - # Shard pipeline batch inputs across data axis ("data") for SPMD multi-host execution - data_sharding = jax.sharding.NamedSharding(self.mesh, P("data")) - - def put_data_on_devices(x, sharding): - if isinstance(x, jax.Array) and hasattr(x, "sharding") and not x.sharding.is_fully_addressable: - return x - if hasattr(sharding, "is_fully_addressable") and sharding.is_fully_addressable: - return jax.device_put(x, sharding) - return device_put_replicated(x, sharding) + max_logging.log(f"{host_prefix} Passed Phase A Sync Barrier (phase_a_complete) successfully! ✅") latents_jax = put_data_on_devices(latents_jax, data_sharding) - prompt_embeds_jax = put_data_on_devices(prompt_embeds_jax, data_sharding) + context_spec = P(None, "context") if "context" in self.mesh.axis_names and self.mesh.shape["context"] > 1 else P() + context_sharding = jax.sharding.NamedSharding(self.mesh, context_spec) + prompt_embeds_jax = put_data_on_devices(prompt_embeds_jax, context_sharding) txt_ids_val = put_data_on_devices(txt_ids_val, data_sharding) img_ids_val = put_data_on_devices(img_ids_val, data_sharding) - print( + max_logging.log( f"{host_prefix} DIAGNOSTIC TENSORS BEFORE PHASE B:\n" f" latents_jax: shape={latents_jax.shape}, dtype={latents_jax.dtype}, sharding={getattr(latents_jax, 'sharding', None)}\n" f" prompt_embeds_jax: shape={prompt_embeds_jax.shape}, dtype={prompt_embeds_jax.dtype}, sharding={getattr(prompt_embeds_jax, 'sharding', None)}\n" f" txt_ids_val: shape={txt_ids_val.shape}, dtype={txt_ids_val.dtype}, sharding={getattr(txt_ids_val, 'sharding', None)}\n" - f" img_ids_val: shape={img_ids_val.shape}, dtype={img_ids_val.dtype}, sharding={getattr(img_ids_val, 'sharding', None)}", - flush=True, + f" img_ids_val: shape={img_ids_val.shape}, dtype={img_ids_val.dtype}, sharding={getattr(img_ids_val, 'sharding', None)}" ) # Stage Sync 2: Pre-Phase B Start multihost_utils.sync_global_devices("pre_phase_b_start") - print(f"{host_prefix} Passed Pre-Phase B Sync Barrier (pre_phase_b_start) successfully! ✅", flush=True) + max_logging.log(f"{host_prefix} Passed Pre-Phase B Sync Barrier (pre_phase_b_start) successfully! ✅") + + t0_denoise_start = time.perf_counter() + trace["qwen3_to_denoise"] = t0_denoise_start - t0_qwen3_end + max_logging.log(f" -> [TIMING] Qwen3 to Denoising Overhead: {trace['qwen3_to_denoise']:.4f} seconds ⏱️") # --------------------------------------------------------------------- # PHASE B: Denoising Loop (Flux Transformer - Standalone Step JIT) # --------------------------------------------------------------------- - print( - f"{host_prefix} [PHASE B] Running {num_inference_steps}-step E2E Denoising Loop on a batch of {batch_size} images...", - flush=True, + steps_to_run = num_inference_steps + max_logging.log( + f"{host_prefix} [PHASE B] Running fused {steps_to_run}-step E2E Denoising Loop Scan on a batch of {batch_size} images (warmup={warmup})..." ) - t0 = time.perf_counter() try: guidance_vec_val = None vec_val = None - - for step_idx in range(num_inference_steps): - timestep = scheduler_state.timesteps[step_idx] - t_vec = jnp.full((batch_size,), timestep / 1000.0, dtype=latents_jax.dtype) - - model_output = self._jitted_transformer_step( - params, latents_jax, img_ids_val, prompt_embeds_jax, txt_ids_val, vec_val, t_vec, guidance_vec_val + replicated_sharding = jax.sharding.NamedSharding(self.mesh, P()) + timesteps_device = put_data_on_devices(scheduler_state.timesteps, replicated_sharding) + sigmas_device = put_data_on_devices(scheduler_state.sigmas, replicated_sharding) + + do_prof_denoise = profile_target in ("all", "denoise") + if do_prof_denoise: + tb_dir = getattr(self._config, "tensorboard_dir", "/tmp") + jax.profiler.start_trace(os.path.join(tb_dir, "profile_denoise")) + with jax.named_scope("fused_flux_denoise_loop"): + latents_jax = self._jitted_fused_denoise_loop( + params, + latents_jax, + img_ids_val, + prompt_embeds_jax, + txt_ids_val, + vec_val, + timesteps_device, + sigmas_device, + guidance_vec_val, + seq_len_img, ) + latents_jax.block_until_ready() + if do_prof_denoise: + jax.profiler.stop_trace() - prev_sample, _ = self.scheduler.step( - state=scheduler_state, - model_output=model_output.sample, - timestep=scheduler_state.timesteps[step_idx], - sample=latents_jax, - return_dict=False, - ) - latents_jax = prev_sample - - latents_jax.block_until_ready() except Exception as e: - print(f"❌ {host_prefix} EXCEPTION IN DENOISE LOOP: {e}", flush=True) + max_logging.log(f"❌ {host_prefix} EXCEPTION IN DENOISE LOOP: {e}") import traceback traceback.print_exc() @@ -331,32 +734,45 @@ def put_data_on_devices(x, sharding): # Stage Sync 3: Phase B Complete multihost_utils.sync_global_devices("phase_b_complete") - print(f"{host_prefix} Passed Phase B Sync Barrier (phase_b_complete) successfully! ✅", flush=True) + max_logging.log(f"{host_prefix} Passed Phase B Sync Barrier (phase_b_complete) successfully! ✅") - trace["denoise_loop"] = time.perf_counter() - t0 + t0_denoise_end = time.perf_counter() + trace["denoise_loop"] = t0_denoise_end - t0_denoise_start max_logging.log(f" -> [TIMING] Denoising Loop (Flux): {trace['denoise_loop']:.4f} seconds ⏱️") # --------------------------------------------------------------------- # PHASE C: Decode Latents (VAE Decoder) # --------------------------------------------------------------------- max_logging.log("[PHASE C] Decoding final latents to RGB image using JAX VAE decoder on TPU...") - t0 = time.perf_counter() - - # Apply Channel-wise Batch Normalization Scaling in packed sequence format (denormalize) - vae_bn_mean_seq = vae_bn_mean.reshape(1, 1, 128) - vae_bn_std_seq = vae_bn_std.reshape(1, 1, 128) - latents_bn = latents_jax * vae_bn_std_seq + vae_bn_mean_seq - - # Unpack packed latents back to spatial grid - latents_unpacked = unpack_latents(latents_bn, batch_size, 32, height, width) - # Decode VAE latents to RGB pixels - decoded_out = self._jitted_vae_decode(vae_params, latents_unpacked) - # VAE output is in decoded_out.sample + # Slice target latents from joint latents if reference images were present + if latents_jax.shape[1] > seq_len_img: + latents_jax = latents_jax[:, :seq_len_img, :] + + # Decode VAE latents to RGB pixels using fused JIT vae_decode + data_sharding = self._get_dynamic_batch_sharding() + replicated_sharding = jax.sharding.NamedSharding(self.mesh, P()) + latents_jax = put_data_on_devices(latents_jax, data_sharding) + vae_bn_mean_jax = put_data_on_devices(jnp.array(vae_bn_mean, dtype=jnp.float32), replicated_sharding) + vae_bn_std_jax = put_data_on_devices(jnp.array(vae_bn_std, dtype=jnp.float32), replicated_sharding) + + t0_vae_start = time.perf_counter() + trace["denoise_to_vae"] = t0_vae_start - t0_denoise_end + max_logging.log(f" -> [TIMING] Denoising to VAE Overhead: {trace['denoise_to_vae']:.4f} seconds ⏱️") + + do_prof_vae = profile_target in ("all", "vae") + if do_prof_vae: + tb_dir = getattr(self._config, "tensorboard_dir", "/tmp") + jax.profiler.start_trace(os.path.join(tb_dir, "profile_vae")) + with jax.named_scope("vae_decoder"): + decoded_out = self._jitted_vae_decode(vae_params, latents_jax, vae_bn_mean_jax, vae_bn_std_jax, height, width) images_rgb = decoded_out.sample images_rgb.block_until_ready() + if do_prof_vae: + jax.profiler.stop_trace() - trace["vae_decode"] = time.perf_counter() - t0 + t0_vae_end = time.perf_counter() + trace["vae_decode"] = t0_vae_end - t0_vae_start max_logging.log(f" -> [TIMING] VAE Decoding: {trace['vae_decode']:.4f} seconds ⏱️") # --------------------------------------------------------------------- @@ -364,15 +780,15 @@ def put_data_on_devices(x, sharding): # --------------------------------------------------------------------- max_logging.log("Postprocessing and saving generated images...") saved_paths = [] - # Clamp pixels and scale to [0, 255] - images_rgb = jnp.clip((images_rgb + 1.0) / 2.0, 0.0, 1.0) + # Perform pixel scaling, clamping, and uint8 conversion directly on TPU hardware + images_uint8 = jnp.clip((images_rgb + 1.0) * 127.5, 0.0, 255.0).astype(jnp.uint8) if jax.process_count() > 1: - images_numpy = multihost_utils.process_allgather(images_rgb, tiled=True) + images_numpy = multihost_utils.process_allgather(images_uint8, tiled=True) else: - images_numpy = np.array(images_rgb) + images_numpy = np.array(images_uint8) for b_idx in range(batch_size): - image_np = np.array(images_numpy[b_idx] * 255.0, dtype=np.uint8) + image_np = np.array(images_numpy[b_idx]) # Transpose channel dimension if shape is (C, H, W) instead of (H, W, C) if image_np.shape[0] == 3: image_np = image_np.transpose(1, 2, 0) @@ -386,8 +802,15 @@ def put_data_on_devices(x, sharding): batch_output_name = output_name output_png_path = os.path.join(output_dir, batch_output_name) - img.save(output_png_path) + img.save(output_png_path, format="PNG", compress_level=1) max_logging.log(f" -> Saved image: {output_png_path} | Prompt: '{prompts[b_idx]}'") saved_paths.append(output_png_path) + t0_save_end = time.perf_counter() + trace["image_saving"] = t0_save_end - t0_vae_end + trace["e2e_pipeline_total"] = t0_save_end - t_pipeline_start + + max_logging.log(f" -> [TIMING] Image Saving: {trace['image_saving']:.4f} seconds ⏱️") + max_logging.log(f" -> [TIMING] E2E Pipeline Total: {trace['e2e_pipeline_total']:.4f} seconds ⏱️") + return saved_paths, trace diff --git a/src/maxdiffusion/schedulers/scheduling_flow_match_flax.py b/src/maxdiffusion/schedulers/scheduling_flow_match_flax.py index 8e9f38ff4..f649b08da 100644 --- a/src/maxdiffusion/schedulers/scheduling_flow_match_flax.py +++ b/src/maxdiffusion/schedulers/scheduling_flow_match_flax.py @@ -300,6 +300,7 @@ def step( sample: jnp.ndarray, to_final: bool = False, return_dict: bool = True, + step_index: Optional[int] = None, ) -> Union[FlaxFlowMatchSchedulerOutput, Tuple]: """ Propagates the sample with the flow matching scheduler. @@ -317,12 +318,17 @@ def step( Whether this is the final step. return_dict (`bool`): Whether to return a `FlaxFlowMatchSchedulerOutput` object. + step_index (`Optional[int]`): + Optional direct step index to bypass dynamic _find_timestep_id calculation. Returns: `FlaxFlowMatchSchedulerOutput` or `tuple`: A tuple (`prev_sample`, `state`) or a `FlaxFlowMatchSchedulerOutput` object containing the previous sample and the updated state. """ - timestep_id = self._find_timestep_id(state, timestep) + if step_index is not None: + timestep_id = step_index + else: + timestep_id = self._find_timestep_id(state, timestep) sigma = state.sigmas[timestep_id] def get_next_sigma(): diff --git a/src/maxdiffusion/tests/edit_flux2klein_e2e_test.py b/src/maxdiffusion/tests/edit_flux2klein_e2e_test.py new file mode 100644 index 000000000..18a4ff074 --- /dev/null +++ b/src/maxdiffusion/tests/edit_flux2klein_e2e_test.py @@ -0,0 +1,349 @@ +""" +Copyright 2026 Google LLC + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + https://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +""" + +import os +import gc +import unittest +import pytest +import numpy as np +from PIL import Image +from skimage.metrics import structural_similarity as ssim +import torch + +import jax +import jax.numpy as jnp +from flax import nnx +from jax.sharding import Mesh +from transformers import AutoConfig, Qwen2TokenizerFast + +from maxdiffusion import pyconfig +from maxdiffusion.max_utils import create_device_mesh +from maxdiffusion.models.flux.transformers.transformer_flux_flax import NNXFlux2KleinTransformer2DModel +from maxdiffusion.models.flux.vae.autoencoder_kl_flux2_nnx import ( + NNXAutoencoderKLFlux2, + load_and_convert_flux2klein_nnx_vae_weights, +) +from maxdiffusion.models.flux.util import load_and_convert_flux_klein_nnx_weights +from maxdiffusion.models.qwen3_flax import FlaxQwen3Model, FlaxQwen3Config +from maxdiffusion.models.qwen3_utils import load_and_convert_qwen3_weights +from maxdiffusion.schedulers.scheduling_flow_match_flax import FlaxFlowMatchScheduler +from maxdiffusion.pipelines.flux.flux2klein_pipeline import FlaxFlux2KleinPipeline + +IN_GITHUB_ACTIONS = os.getenv("GITHUB_ACTIONS") == "true" +THIS_DIR = os.path.dirname(os.path.abspath(__file__)) +PROMPT = "a vibrant artistic painting combining the dog, car, mountain, and fruit bowl in surreal neon lighting" + + +class TestFlux2KleinImageEditE2EParity(unittest.TestCase): + """End-to-End Parity Test between PyTorch Diffusers CPU and MaxDiffusion TPU.""" + + def setUp(self): + jax.config.update("jax_default_matmul_precision", "highest") + jax.config.update("jax_use_shardy_partitioner", True) + + if "FLUX2_KLEIN_4B_MODEL_PATH" in os.environ: + self.model_dir = os.environ["FLUX2_KLEIN_4B_MODEL_PATH"] + else: + hf_home = os.environ.get("HF_HOME", os.path.expanduser("~/.cache/huggingface")) + candidates = [ + os.path.join(hf_home, "hub/models--black-forest-labs--FLUX.2-klein-4B/snapshots"), + os.path.join(hf_home, "hub/models--black-forest-labs--FLUX.2-klein-4b/snapshots"), + "/mnt/hyperdisk_weights/hub/models--black-forest-labs--FLUX.2-klein-4B/snapshots", + "/mnt/data/models/flux2klein-4b", + ] + self.model_dir = None + for c in candidates: + if os.path.exists(c): + if "snapshots" in c: + snaps = os.listdir(c) + if snaps: + self.model_dir = os.path.join(c, snaps[0]) + else: + self.model_dir = c + if self.model_dir: + self.transformer_path = os.path.join(self.model_dir, "transformer") + self.vae_path = os.path.join(self.model_dir, "vae", "diffusion_pytorch_model.safetensors") + self.text_encoder_path = os.path.join(self.model_dir, "text_encoder") + self.tokenizer_path = os.path.join(self.model_dir, "tokenizer") + if os.path.exists(self.transformer_path) and os.path.exists(self.vae_path): + break + if self.model_dir is None: + self.model_dir = "black-forest-labs/FLUX.2-klein-4B" + + if hasattr(self, "model_dir") and self.model_dir and not hasattr(self, "transformer_path"): + self.transformer_path = os.path.join(self.model_dir, "transformer") + self.vae_path = os.path.join(self.model_dir, "vae", "diffusion_pytorch_model.safetensors") + self.text_encoder_path = os.path.join(self.model_dir, "text_encoder") + self.tokenizer_path = os.path.join(self.model_dir, "tokenizer") + + self.output_dir = "/tmp/e2e_parity" + os.makedirs(self.output_dir, exist_ok=True) + + # Resolve reference images + ref_dir = os.path.join(THIS_DIR, "images", "flux2klein") + self.ref_images = [] + if os.path.exists(ref_dir): + for i in range(4): + p = os.path.join(ref_dir, f"ref_image_{i}.png") + if os.path.exists(p): + self.ref_images.append(Image.open(p).convert("RGB")) + + if len(self.ref_images) < 4: + # Generate synthetic test reference images if not present + colors = [(255, 0, 0), (0, 255, 0), (0, 0, 255), (255, 255, 0)] + for i, c in enumerate(colors): + arr = np.full((512, 512, 3), c, dtype=np.uint8) + self.ref_images.append(Image.fromarray(arr)) + + @pytest.mark.skipif(IN_GITHUB_ACTIONS, reason="Don't run on Github Actions (requires TPU and full weights)") + def test_e2e_image_edit_parity_vs_diffusers(self): + """Generates an image edit on PyTorch Diffusers CPU and MaxDiffusion TPU and asserts SSIM >= 0.75.""" + from diffusers import Flux2KleinPipeline as DiffusersFlux2KleinPipeline + + print("\n" + "=" * 80) + print("🚀 [STEP 1/3] Running Reference PyTorch Diffusers CPU Pipeline...") + print("=" * 80) + + diffusers_pipe = DiffusersFlux2KleinPipeline.from_pretrained(self.model_dir, torch_dtype=torch.bfloat16) + diffusers_pipe.to("cpu") + + # Generate initial noise latents deterministically on CPU (4D tensor for Diffusers prepare_latents) + gen = torch.Generator(device="cpu").manual_seed(42) + raw_latents_pt = torch.randn( + (1, 128, 512 // 16, 512 // 16), + generator=gen, + dtype=torch.bfloat16, + device="cpu", + ) + + with torch.no_grad(): + diffusers_out = diffusers_pipe( + prompt=PROMPT, + image=self.ref_images, + height=512, + width=512, + num_inference_steps=4, + latents=raw_latents_pt, + guidance_scale=1.0, + ) + + diffusers_image = diffusers_out.images[0] + diffusers_img_path = os.path.join(self.output_dir, "diffusers_cpu_output.png") + diffusers_image.save(diffusers_img_path) + print(f" -> Saved PyTorch Diffusers output to: {diffusers_img_path}") + + # Free PyTorch pipeline memory before TPU run + del diffusers_pipe + gc.collect() + + print("\n" + "=" * 80) + print("🚀 [STEP 2/3] Running MaxDiffusion Unified FlaxFlux2KleinPipeline on TPU...") + print("=" * 80) + + # 1. Device mesh setup + active_devices = jax.devices() + active_device_count = len(active_devices) + + pyconfig._config = None + pyconfig.config = None + config_path = os.path.join(THIS_DIR, "..", "configs", "base_flux2klein.yml") + args = [ + None, + config_path, + "run_name=e2e_parity_test", + f"output_dir={self.output_dir}", + f"per_device_batch_size={1.0 / active_device_count}", + "height=512", + "width=512", + "seed=42", + "weights_dtype=bfloat16", + "activations_dtype=bfloat16", + "precision=DEFAULT", + "text_encoder_attention=dot_product", + ] + pyconfig.initialize(args) + config = pyconfig.config + + if active_device_count > 1: + pyconfig._config.keys["ici_tensor_parallelism"] = active_device_count + pyconfig._config.keys["ici_data_parallelism"] = 1 + pyconfig._config.keys["ici_fsdp_parallelism"] = 1 + pyconfig._config.keys["ici_context_parallelism"] = 1 + + devices_array = create_device_mesh(config, devices=active_devices) + mesh = Mesh(devices_array, config.mesh_axes) + + # 2. Load NNX Transformer + print(" -> Loading NNX Transformer weights...") + rngs = nnx.Rngs(0) + transformer = NNXFlux2KleinTransformer2DModel( + rngs=rngs, + patch_size=1, + in_channels=128, + num_layers=5, + num_single_layers=20, + attention_head_dim=128, + num_attention_heads=24, + joint_attention_dim=7680, + pooled_projection_dim=None, + guidance_embeds=False, + axes_dim=(32, 32, 32, 32), + scale_shift_order="scale_shift", + dtype=jnp.bfloat16, + weights_dtype=jnp.bfloat16, + ) + t_state = load_and_convert_flux_klein_nnx_weights( + self.transformer_path, + nnx.state(transformer, nnx.Param), + num_double_layers=5, + num_single_layers=20, + dtype=jnp.bfloat16, + ) + nnx.update(transformer, t_state) + + # 3. Load NNX VAE + print(" -> Loading NNX VAE weights...") + nnx_vae = NNXAutoencoderKLFlux2(dtype=jnp.bfloat16, param_dtype=jnp.bfloat16) + bn_mean, bn_std = load_and_convert_flux2klein_nnx_vae_weights(self.vae_path, nnx_vae, dtype=jnp.bfloat16) + + # 4. Load Qwen3 + print(" -> Loading Qwen3 weights...") + pt_config = AutoConfig.from_pretrained(self.text_encoder_path) + qwen3_config = FlaxQwen3Config( + vocab_size=pt_config.vocab_size, + hidden_size=pt_config.hidden_size, + intermediate_size=pt_config.intermediate_size, + num_hidden_layers=pt_config.num_hidden_layers, + num_attention_heads=pt_config.num_attention_heads, + num_key_value_heads=pt_config.num_key_value_heads, + max_position_embeddings=pt_config.max_position_embeddings, + rms_norm_eps=pt_config.rms_norm_eps, + rope_theta=pt_config.rope_theta, + dtype=jnp.bfloat16, + max_layer_to_run=27, + ) + text_encoder = FlaxQwen3Model(config=qwen3_config) + abstract_q_vars = text_encoder.init( + jax.random.PRNGKey(0), jnp.zeros((1, 512), dtype=jnp.int32), jnp.zeros((1, 512), dtype=jnp.int32) + ) + q_params = load_and_convert_qwen3_weights(self.text_encoder_path, abstract_q_vars["params"], qwen3_config) + + tokenizer = Qwen2TokenizerFast.from_pretrained(self.tokenizer_path) + scheduler = FlaxFlowMatchScheduler( + num_train_timesteps=1000, + shift=1.0, + sigma_max=1.0, + sigma_min=0.001, + inverse_timesteps=False, + extra_one_step=False, + reverse_sigmas=False, + use_dynamic_shifting=True, + time_shift_type="exponential", + ) + + # 5. Place parameters on TPU HBM + t_params = nnx.state(transformer, nnx.Param) + v_params = nnx.state(nnx_vae, nnx.Param) + + t_params = jax.device_put(t_params) + v_params = jax.device_put(v_params) + q_params = jax.device_put(q_params) + + # 6. Instantiate Unified FlaxFlux2KleinPipeline + pipeline = FlaxFlux2KleinPipeline( + transformer=transformer, + vae=nnx_vae, + text_encoder=text_encoder, + tokenizer=tokenizer, + scheduler=scheduler, + config=config, + mesh=mesh, + ) + + # 7. AOT Compile async + pipeline.compile_aot_async( + params=t_params, + vae_params=v_params, + qwen3_params=q_params, + vae_bn_mean=bn_mean, + vae_bn_std=bn_std, + batch_size=1, + height=512, + width=512, + images=self.ref_images, + ) + + # Convert PyTorch initial noise latents to JAX array (shape: 1, 32, 64, 64) + initial_latents_jax = jnp.array(raw_latents_pt.detach().float().cpu().numpy()) + + # 8. Run pipeline + print(f" -> Running FlaxFlux2KleinPipeline with {len(self.ref_images)} reference images on TPU...") + pipeline( + prompt=PROMPT, + params=t_params, + vae_params=v_params, + qwen3_params=q_params, + vae_bn_mean=bn_mean, + vae_bn_std=bn_std, + transformer_shardings=None, + vae_shardings=None, + qwen3_shardings=None, + height=512, + width=512, + num_inference_steps=4, + batch_size=1, + images=self.ref_images, + use_latents=True, + latents=initial_latents_jax, + output_dir=self.output_dir, + output_name="maxdiffusion_tpu_output.png", + ) + + maxdiff_img_path = os.path.join(self.output_dir, "maxdiffusion_tpu_output.png") + self.assertTrue(os.path.exists(maxdiff_img_path), "MaxDiffusion output image was not saved!") + maxdiff_image = Image.open(maxdiff_img_path).convert("RGB") + + print("\n" + "=" * 80) + print("📊 [STEP 3/3] Evaluating End-to-End Parity (SSIM & PSNR)...") + print("=" * 80) + + diffusers_arr = np.array(diffusers_image).astype(np.uint8) + maxdiff_arr = np.array(maxdiff_image).astype(np.uint8) + + self.assertEqual(diffusers_arr.shape, maxdiff_arr.shape) + + ssim_val = ssim(diffusers_arr, maxdiff_arr, channel_axis=-1, data_range=255) + mse = np.mean((diffusers_arr.astype(np.float64) - maxdiff_arr.astype(np.float64)) ** 2) + psnr_val = 10.0 * np.log10(255.0**2 / (mse + 1e-10)) + + print(f" -> SSIM (Diffusers CPU vs MaxDiffusion TPU): {ssim_val:.6f}") + print(f" -> PSNR (Diffusers CPU vs MaxDiffusion TPU): {psnr_val:.2f} dB") + print(f" -> MSE: {mse:.4f}") + + # Create side-by-side comparison image + side_by_side = Image.new("RGB", (1024, 512)) + side_by_side.paste(diffusers_image, (0, 0)) + side_by_side.paste(maxdiff_image, (512, 0)) + comparison_path = os.path.join(self.output_dir, "e2e_parity_diffusers_vs_maxdiffusion.png") + side_by_side.save(comparison_path) + print(f" -> Saved side-by-side comparison to: {comparison_path}") + + self.assertGreaterEqual(ssim_val, 0.75, f"SSIM score {ssim_val:.4f} is below target threshold 0.75!") + print("🎉 END-TO-END PARITY TEST PASSED! MaxDiffusion matches Diffusers reference!") + + +if __name__ == "__main__": + unittest.main() diff --git a/src/maxdiffusion/tests/generate_flux2klein_smoke_test.py b/src/maxdiffusion/tests/generate_flux2klein_smoke_test.py index 24362d35d..0265acb8b 100644 --- a/src/maxdiffusion/tests/generate_flux2klein_smoke_test.py +++ b/src/maxdiffusion/tests/generate_flux2klein_smoke_test.py @@ -17,6 +17,7 @@ import os import unittest import pytest +import jax import numpy as np from PIL import Image @@ -40,7 +41,7 @@ def test_flux2klein_4b_smoke(self): self.assertTrue(os.path.exists(ref_path), f"Reference image not found: {ref_path}") base_image = np.array(Image.open(ref_path)).astype(np.uint8) - output_dir = "/mnt/data/smoke_test_4b" if os.path.exists("/mnt/data") else "/tmp/smoke_test_4b" + output_dir = "/tmp/smoke_test_4b" os.makedirs(output_dir, exist_ok=True) out_path = os.path.join(output_dir, "flux2klein_generated_image.png") if os.path.exists(out_path): @@ -54,22 +55,24 @@ def test_flux2klein_4b_smoke(self): "run_name=smoke_test_4b", f"output_dir={output_dir}", "jax_cache_dir=/tmp/cache_dir", - "skip_jax_distributed_system=True", f"prompt={PROMPT}", "height=512", "width=512", - "batch_size=1", + f"per_device_batch_size={1.0 / jax.device_count()}", "seed=42", - "ici_fsdp_parallelism=-1", "weights_dtype=bfloat16", "activations_dtype=bfloat16", "precision=DEFAULT", + "num_reps=5", + "text_encoder_attention=dot_product", ] generate_flux2klein.main(args) - self.assertTrue(os.path.exists(out_path), "Smoke test 4B failed to produce output image!") - test_image = np.array(Image.open(out_path)).astype(np.uint8) + rep_out_path = os.path.join(output_dir, "rep_1_flux2klein_generated_image.png") + final_out_path = rep_out_path if os.path.exists(rep_out_path) else out_path + self.assertTrue(os.path.exists(final_out_path), "Smoke test 4B failed to produce output image!") + test_image = np.array(Image.open(final_out_path)).astype(np.uint8) self.assertEqual(base_image.shape, test_image.shape) ssim_compare = ssim(base_image, test_image, channel_axis=-1, data_range=255) @@ -83,7 +86,7 @@ def test_flux2klein_9b_smoke(self): self.assertTrue(os.path.exists(ref_path), f"Reference image not found: {ref_path}") base_image = np.array(Image.open(ref_path)).astype(np.uint8) - output_dir = "/mnt/data/smoke_test_9b" if os.path.exists("/mnt/data") else "/tmp/smoke_test_9b" + output_dir = "/tmp/smoke_test_9b" os.makedirs(output_dir, exist_ok=True) out_path = os.path.join(output_dir, "flux2klein_generated_image.png") if os.path.exists(out_path): @@ -97,27 +100,127 @@ def test_flux2klein_9b_smoke(self): "run_name=smoke_test_9b", f"output_dir={output_dir}", "jax_cache_dir=/tmp/cache_dir", - "skip_jax_distributed_system=True", f"prompt={PROMPT}", "height=512", "width=512", - "batch_size=1", + f"per_device_batch_size={1.0 / jax.device_count()}", "seed=42", - "ici_fsdp_parallelism=-1", "weights_dtype=bfloat16", "activations_dtype=bfloat16", "precision=DEFAULT", + "num_reps=5", + "text_encoder_attention=dot_product", ] generate_flux2klein.main(args) - self.assertTrue(os.path.exists(out_path), "Smoke test 9B failed to produce output image!") - test_image = np.array(Image.open(out_path)).astype(np.uint8) + rep_out_path = os.path.join(output_dir, "rep_1_flux2klein_generated_image.png") + final_out_path = rep_out_path if os.path.exists(rep_out_path) else out_path + self.assertTrue(os.path.exists(final_out_path), "Smoke test 9B failed to produce output image!") + test_image = np.array(Image.open(final_out_path)).astype(np.uint8) self.assertEqual(base_image.shape, test_image.shape) ssim_compare = ssim(base_image, test_image, channel_axis=-1, data_range=255) print(f"\n[SMOKE TEST 9B] SSIM Score: {ssim_compare:.6f}") - self.assertGreaterEqual(ssim_compare, 0.80) + self.assertGreaterEqual(ssim_compare, 0.8) + + @pytest.mark.skipif(IN_GITHUB_ACTIONS, reason="Don't run smoke tests on Github Actions (requires TPU HBM)") + def test_flux2klein_4b_image_edit_smoke(self): + """End-to-end smoke test for Flux.2-klein-4B image editing at 512x512.""" + ref_gold_path = os.path.join(THIS_DIR, "images", "ref_flux2klein_4b_image_edit.png") + self.assertTrue(os.path.exists(ref_gold_path), f"Golden reference image not found: {ref_gold_path}") + base_image = np.array(Image.open(ref_gold_path)).astype(np.uint8) + + input_img_path = os.path.join(THIS_DIR, "images", "ref_flux2klein_4b.png") + self.assertTrue(os.path.exists(input_img_path), f"Input reference image not found: {input_img_path}") + + output_dir = "/tmp/smoke_test_image_edit_4b" + os.makedirs(output_dir, exist_ok=True) + out_path = os.path.join(output_dir, "flux2klein_generated_image.png") + if os.path.exists(out_path): + os.remove(out_path) + + pyconfig._config = None + pyconfig.config = None + args = [ + None, + os.path.join(THIS_DIR, "..", "configs", "base_flux2klein.yml"), + "run_name=smoke_test_image_edit_4b", + f"output_dir={output_dir}", + "jax_cache_dir=/tmp/cache_dir", + f"image_paths=['{input_img_path}']", + "prompt=change the lighting to evening", + "height=512", + "width=512", + f"per_device_batch_size={1.0 / jax.device_count()}", + "seed=42", + "weights_dtype=bfloat16", + "activations_dtype=bfloat16", + "precision=DEFAULT", + "num_reps=5", + "text_encoder_attention=dot_product", + ] + + generate_flux2klein.main(args) + + rep_out_path = os.path.join(output_dir, "rep_1_flux2klein_generated_image.png") + final_out_path = rep_out_path if os.path.exists(rep_out_path) else out_path + self.assertTrue(os.path.exists(final_out_path), "Smoke test 4B image edit failed to produce output image!") + test_image = np.array(Image.open(final_out_path)).astype(np.uint8) + + self.assertEqual(base_image.shape, test_image.shape) + ssim_compare = ssim(base_image, test_image, channel_axis=-1, data_range=255) + print(f"\n[SMOKE TEST 4B IMAGE EDIT] SSIM Score: {ssim_compare:.6f}") + self.assertGreaterEqual(ssim_compare, 0.95) + + @pytest.mark.skipif(IN_GITHUB_ACTIONS, reason="Don't run smoke tests on Github Actions (requires TPU HBM)") + def test_flux2klein_9b_image_edit_smoke(self): + """End-to-end smoke test for Flux.2-klein-9B image editing at 512x512.""" + ref_gold_path = os.path.join(THIS_DIR, "images", "ref_flux2klein_9b_image_edit.png") + self.assertTrue(os.path.exists(ref_gold_path), f"Golden reference image not found: {ref_gold_path}") + base_image = np.array(Image.open(ref_gold_path)).astype(np.uint8) + + input_img_path = os.path.join(THIS_DIR, "images", "ref_flux2klein_4b.png") + self.assertTrue(os.path.exists(input_img_path), f"Input reference image not found: {input_img_path}") + + output_dir = "/tmp/smoke_test_image_edit_9b" + os.makedirs(output_dir, exist_ok=True) + out_path = os.path.join(output_dir, "flux2klein_generated_image.png") + if os.path.exists(out_path): + os.remove(out_path) + + pyconfig._config = None + pyconfig.config = None + args = [ + None, + os.path.join(THIS_DIR, "..", "configs", "base_flux2klein_9B.yml"), + "run_name=smoke_test_image_edit_9b", + f"output_dir={output_dir}", + "jax_cache_dir=/tmp/cache_dir", + f"image_paths=['{input_img_path}']", + "prompt=change the lighting to evening", + "height=512", + "width=512", + f"per_device_batch_size={1.0 / jax.device_count()}", + "seed=42", + "weights_dtype=bfloat16", + "activations_dtype=bfloat16", + "precision=DEFAULT", + "num_reps=5", + "text_encoder_attention=dot_product", + ] + + generate_flux2klein.main(args) + + rep_out_path = os.path.join(output_dir, "rep_1_flux2klein_generated_image.png") + final_out_path = rep_out_path if os.path.exists(rep_out_path) else out_path + self.assertTrue(os.path.exists(final_out_path), "Smoke test 9B image edit failed to produce output image!") + test_image = np.array(Image.open(final_out_path)).astype(np.uint8) + + self.assertEqual(base_image.shape, test_image.shape) + ssim_compare = ssim(base_image, test_image, channel_axis=-1, data_range=255) + print(f"\n[SMOKE TEST 9B IMAGE EDIT] SSIM Score: {ssim_compare:.6f}") + self.assertGreaterEqual(ssim_compare, 0.95) if __name__ == "__main__": diff --git a/src/maxdiffusion/tests/images/flux2klein/ref_image_0.png b/src/maxdiffusion/tests/images/flux2klein/ref_image_0.png new file mode 100644 index 000000000..476ba5984 Binary files /dev/null and b/src/maxdiffusion/tests/images/flux2klein/ref_image_0.png differ diff --git a/src/maxdiffusion/tests/images/flux2klein/ref_image_1.png b/src/maxdiffusion/tests/images/flux2klein/ref_image_1.png new file mode 100644 index 000000000..d14acbe2a Binary files /dev/null and b/src/maxdiffusion/tests/images/flux2klein/ref_image_1.png differ diff --git a/src/maxdiffusion/tests/images/flux2klein/ref_image_2.png b/src/maxdiffusion/tests/images/flux2klein/ref_image_2.png new file mode 100644 index 000000000..cb379a6e8 Binary files /dev/null and b/src/maxdiffusion/tests/images/flux2klein/ref_image_2.png differ diff --git a/src/maxdiffusion/tests/images/flux2klein/ref_image_3.png b/src/maxdiffusion/tests/images/flux2klein/ref_image_3.png new file mode 100644 index 000000000..bffdbe23c Binary files /dev/null and b/src/maxdiffusion/tests/images/flux2klein/ref_image_3.png differ diff --git a/src/maxdiffusion/tests/images/ref_flux2klein_4b.png b/src/maxdiffusion/tests/images/ref_flux2klein_4b.png index 0eba6a072..e6a30c408 100644 Binary files a/src/maxdiffusion/tests/images/ref_flux2klein_4b.png and b/src/maxdiffusion/tests/images/ref_flux2klein_4b.png differ diff --git a/src/maxdiffusion/tests/images/ref_flux2klein_4b_image_edit.png b/src/maxdiffusion/tests/images/ref_flux2klein_4b_image_edit.png new file mode 100644 index 000000000..f24599b58 Binary files /dev/null and b/src/maxdiffusion/tests/images/ref_flux2klein_4b_image_edit.png differ diff --git a/src/maxdiffusion/tests/images/ref_flux2klein_9b.png b/src/maxdiffusion/tests/images/ref_flux2klein_9b.png index 594464a8f..704d0fee8 100644 Binary files a/src/maxdiffusion/tests/images/ref_flux2klein_9b.png and b/src/maxdiffusion/tests/images/ref_flux2klein_9b.png differ diff --git a/src/maxdiffusion/tests/images/ref_flux2klein_9b_image_edit.png b/src/maxdiffusion/tests/images/ref_flux2klein_9b_image_edit.png new file mode 100644 index 000000000..7d938fb17 Binary files /dev/null and b/src/maxdiffusion/tests/images/ref_flux2klein_9b_image_edit.png differ diff --git a/src/maxdiffusion/tests/nnx_flux2klein_test.py b/src/maxdiffusion/tests/nnx_flux2klein_test.py index 5ae88d8eb..9eb9d1f41 100644 --- a/src/maxdiffusion/tests/nnx_flux2klein_test.py +++ b/src/maxdiffusion/tests/nnx_flux2klein_test.py @@ -23,7 +23,7 @@ import jax.numpy as jnp from flax import nnx -from maxdiffusion.models.flux.transformers.transformer_flux_flax import NNXFluxTransformer2DModel +from maxdiffusion.models.flux.transformers.transformer_flux_flax import NNXFlux2KleinTransformer2DModel from maxdiffusion.models.qwen3_flax import FlaxQwen3Config, NNXFlaxQwen3Model from maxdiffusion.models.vae_flax import NNXFlaxAutoencoderKL from maxdiffusion.models.embeddings_flax import NNXCombinedTimestepGuidanceTextProjEmbeddings @@ -83,9 +83,9 @@ def test_nnx_vae_decoder_forward(self): def test_nnx_flux_transformer_forward(self): rngs = nnx.Rngs(0) - transformer = NNXFluxTransformer2DModel( + transformer = NNXFlux2KleinTransformer2DModel( rngs=rngs, - in_channels=16, + in_channels=128, num_layers=1, num_single_layers=2, attention_head_dim=128, @@ -93,15 +93,16 @@ def test_nnx_flux_transformer_forward(self): joint_attention_dim=128, pooled_projection_dim=128, guidance_embeds=True, - axes_dim=(16, 56, 56), + axes_dim=(32, 32, 32, 32), + theta=2000.0, ) - hidden_states = jnp.ones((1, 64, 16)) + hidden_states = jnp.ones((1, 64, 128)) encoder_hidden_states = jnp.ones((1, 16, 128)) pooled_projections = jnp.ones((1, 128)) timestep = jnp.array([100.0]) guidance = jnp.array([3.5]) - img_ids = jnp.zeros((64, 3)) - txt_ids = jnp.zeros((16, 3)) + img_ids = jnp.zeros((64, 4)) + txt_ids = jnp.zeros((16, 4)) output = transformer( hidden_states=hidden_states, @@ -111,8 +112,9 @@ def test_nnx_flux_transformer_forward(self): img_ids=img_ids, txt_ids=txt_ids, guidance=guidance, - ) - self.assertEqual(output.shape, (1, 64, 16)) + return_dict=False, + )[0] + self.assertEqual(output.shape, (1, 64, 128)) if __name__ == "__main__":