finished hdr writeup

Cargo.lock

@@ -2228,9 +2228,9 @@ dependencies = [
 
 [[package]]
 name = "once_cell"
-version = "1.15.0"
+version = "1.18.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "e82dad04139b71a90c080c8463fe0dc7902db5192d939bd0950f074d014339e1"
+checksum = "dd8b5dd2ae5ed71462c540258bedcb51965123ad7e7ccf4b9a8cafaa4a63576d"
 
 [[package]]
 name = "openssl"
@@ -4401,7 +4401,7 @@ version = "0.48.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "677d2418bec65e3338edb076e806bc1ec15693c5d0104683f2efe857f61056a9"
 dependencies = [
- "windows-targets 0.48.0",
+ "windows-targets 0.48.5",
 ]
 
 [[package]]
@@ -4421,17 +4421,17 @@ dependencies = [
 
 [[package]]
 name = "windows-targets"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "7b1eb6f0cd7c80c79759c929114ef071b87354ce476d9d94271031c0497adfd5"
+checksum = "9a2fa6e2155d7247be68c096456083145c183cbbbc2764150dda45a87197940c"
 dependencies = [
- "windows_aarch64_gnullvm 0.48.0",
- "windows_aarch64_msvc 0.48.0",
- "windows_i686_gnu 0.48.0",
- "windows_i686_msvc 0.48.0",
- "windows_x86_64_gnu 0.48.0",
- "windows_x86_64_gnullvm 0.48.0",
- "windows_x86_64_msvc 0.48.0",
+ "windows_aarch64_gnullvm 0.48.5",
+ "windows_aarch64_msvc 0.48.5",
+ "windows_i686_gnu 0.48.5",
+ "windows_i686_msvc 0.48.5",
+ "windows_x86_64_gnu 0.48.5",
+ "windows_x86_64_gnullvm 0.48.5",
+ "windows_x86_64_msvc 0.48.5",
 ]
 
 [[package]]
@@ -4442,9 +4442,9 @@ checksum = "597a5118570b68bc08d8d59125332c54f1ba9d9adeedeef5b99b02ba2b0698f8"
 
 [[package]]
 name = "windows_aarch64_gnullvm"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "91ae572e1b79dba883e0d315474df7305d12f569b400fcf90581b06062f7e1bc"
+checksum = "2b38e32f0abccf9987a4e3079dfb67dcd799fb61361e53e2882c3cbaf0d905d8"
 
 [[package]]
 name = "windows_aarch64_msvc"
@@ -4466,9 +4466,9 @@ checksum = "e08e8864a60f06ef0d0ff4ba04124db8b0fb3be5776a5cd47641e942e58c4d43"
 
 [[package]]
 name = "windows_aarch64_msvc"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "b2ef27e0d7bdfcfc7b868b317c1d32c641a6fe4629c171b8928c7b08d98d7cf3"
+checksum = "dc35310971f3b2dbbf3f0690a219f40e2d9afcf64f9ab7cc1be722937c26b4bc"
 
 [[package]]
 name = "windows_i686_gnu"
@@ -4490,9 +4490,9 @@ checksum = "c61d927d8da41da96a81f029489353e68739737d3beca43145c8afec9a31a84f"
 
 [[package]]
 name = "windows_i686_gnu"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "622a1962a7db830d6fd0a69683c80a18fda201879f0f447f065a3b7467daa241"
+checksum = "a75915e7def60c94dcef72200b9a8e58e5091744960da64ec734a6c6e9b3743e"
 
 [[package]]
 name = "windows_i686_msvc"
@@ -4514,9 +4514,9 @@ checksum = "44d840b6ec649f480a41c8d80f9c65108b92d89345dd94027bfe06ac444d1060"
 
 [[package]]
 name = "windows_i686_msvc"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "4542c6e364ce21bf45d69fdd2a8e455fa38d316158cfd43b3ac1c5b1b19f8e00"
+checksum = "8f55c233f70c4b27f66c523580f78f1004e8b5a8b659e05a4eb49d4166cca406"
 
 [[package]]
 name = "windows_x86_64_gnu"
@@ -4538,9 +4538,9 @@ checksum = "8de912b8b8feb55c064867cf047dda097f92d51efad5b491dfb98f6bbb70cb36"
 
 [[package]]
 name = "windows_x86_64_gnu"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "ca2b8a661f7628cbd23440e50b05d705db3686f894fc9580820623656af974b1"
+checksum = "53d40abd2583d23e4718fddf1ebec84dbff8381c07cae67ff7768bbf19c6718e"
 
 [[package]]
 name = "windows_x86_64_gnullvm"
@@ -4550,9 +4550,9 @@ checksum = "26d41b46a36d453748aedef1486d5c7a85db22e56aff34643984ea85514e94a3"
 
 [[package]]
 name = "windows_x86_64_gnullvm"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "7896dbc1f41e08872e9d5e8f8baa8fdd2677f29468c4e156210174edc7f7b953"
+checksum = "0b7b52767868a23d5bab768e390dc5f5c55825b6d30b86c844ff2dc7414044cc"
 
 [[package]]
 name = "windows_x86_64_msvc"
@@ -4574,9 +4574,9 @@ checksum = "9aec5da331524158c6d1a4ac0ab1541149c0b9505fde06423b02f5ef0106b9f0"
 
 [[package]]
 name = "windows_x86_64_msvc"
-version = "0.48.0"
+version = "0.48.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "1a515f5799fe4961cb532f983ce2b23082366b898e52ffbce459c86f67c8378a"
+checksum = "ed94fce61571a4006852b7389a063ab983c02eb1bb37b47f8272ce92d06d9538"
 
 [[package]]
 name = "winit"

code/intermediate/tutorial13-hdr/Cargo.toml

@@ -29,7 +29,6 @@ features = ["png", "jpeg", "hdr"]
 reqwest = { version = "0.11" }
 console_error_panic_hook = "0.1"
 console_log = "1.0"
-wgpu = { version = "0.17", features = ["webgl"]}
 wasm-bindgen = "0.2"
 wasm-bindgen-futures = "0.4"
 web-sys = { version = "0.3", features = [

code/intermediate/tutorial13-hdr/src/equirectangular.wgsl

@@ -1,4 +1,4 @@
-const PI: f32 = 3.14159;
+const PI: f32 = 3.1415926535897932384626433832795;
 
 struct Face {
     forward: vec3<f32>,
@@ -12,15 +12,22 @@ var src: texture_2d<f32>;
 
 @group(0)
 @binding(1)
-var output: texture_storage_2d_array<rgba32float, write>;
+var dst: texture_storage_2d_array<rgba32float, write>;
 
 
 @compute
-@workgroup_size(1, 1, 1)
+@workgroup_size(16, 16, 1)
 fn compute_equirect_to_cubemap(
     @builtin(global_invocation_id)
     gid: vec3<u32>,
 ) {
+    // If texture size is not divisible by 32 we
+    // need to make sure we don't try to write to
+    // pixels that don't exist.
+    if gid.x >= u32(textureDimensions(dst).x) {
+        return;
+    }
+
     var FACES: array<Face, 6> = array(
         // FACES +X
         Face(
@@ -60,13 +67,12 @@ fn compute_equirect_to_cubemap(
         ),
     );
 
-    let face = FACES[gid.z];
-
     // Get texture coords relative to cubemap face
-    let dst_dimensions = vec2<f32>(textureDimensions(output));
+    let dst_dimensions = vec2<f32>(textureDimensions(dst));
     let cube_uv = vec2<f32>(gid.xy) / dst_dimensions * 2.0 - 1.0;
 
     // Get spherical coordinate from cube_uv
+    let face = FACES[gid.z];
     let spherical = normalize(face.forward + face.right * cube_uv.x + face.up * cube_uv.y);
 
     // Get coordinate on the equirectangular texture
@@ -77,8 +83,5 @@ fn compute_equirect_to_cubemap(
     // We use textureLoad() as textureSample() is not allowed in compute shaders
     var sample = textureLoad(src, eq_pixel, 0);
 
-    // sample = vec4(cube_uv * 0.5 + 0.5, 0.0, 1.0);
-    // sample = vec4(spherical * 0.5 + 0.5, 1.0);
-
-    textureStore(output, gid.xy, gid.z, sample);
+    textureStore(dst, gid.xy, gid.z, sample);
 }

code/intermediate/tutorial13-hdr/src/lib.rs

@@ -16,6 +16,8 @@ mod hdr;
 mod model;
 mod resources;
 mod texture;
+
+#[cfg(feature = "debug")]
 mod debug;
 
 use model::{DrawLight, DrawModel, Vertex};
@@ -52,8 +54,6 @@ impl CameraUniform {
         self.view = view.into();
         self.view_proj = view_proj.into();
         self.inv_proj = proj.invert().unwrap().into();
-        // self.inv_proj = proj.transpose().into();
-        // self.inv_view = view.invert().unwrap().into();
         self.inv_view = view.transpose().into();
     }
 }
@@ -176,6 +176,7 @@ struct State {
     hdr: hdr::HdrPipeline,
     environment_bind_group: wgpu::BindGroup,
     sky_pipeline: wgpu::RenderPipeline,
+    #[cfg(feature = "debug")]
     debug: debug::Debug,
 }
 
@@ -244,6 +245,10 @@ impl State {
         // The instance is a handle to our GPU
         // BackendBit::PRIMARY => Vulkan + Metal + DX12 + Browser WebGPU
         let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
+            // UPDATED
+            #[cfg(target_arch="wasm32")]
+            backends: wgpu::Backends::BROWSER_WEBGPU,
+            #[cfg(not(target_arch="wasm32"))]
             backends: wgpu::Backends::all(),
             dx12_shader_compiler: Default::default(),
         });
@@ -267,7 +272,8 @@ impl State {
                 &wgpu::DeviceDescriptor {
                     label: None,
                     // UPDATED!
-                    features: wgpu::Features::all_webgpu_mask(),
+                    features: wgpu::Features::empty(),
+                    // UPDATED!
                     limits: wgpu::Limits::downlevel_defaults(),
                 },
                 None, // Trace path
@@ -599,6 +605,7 @@ impl State {
             )
         };
 
+        #[cfg(feature = "debug")]
         let debug = debug::Debug::new(&device, &camera_bind_group_layout, surface_format);
 
         Ok(Self {
@@ -630,6 +637,8 @@ impl State {
             hdr,
             environment_bind_group,
             sky_pipeline,
+
+            #[cfg(feature = "debug")]
             debug,
         })
     }
@@ -768,6 +777,7 @@ impl State {
         // Apply tonemapping
         self.hdr.process(&mut encoder, &view);
 
+        #[cfg(feature = "debug")]
         {
             let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
                 label: Some("Debug"),
@@ -792,7 +802,7 @@ impl State {
 }
 
 #[cfg_attr(target_arch = "wasm32", wasm_bindgen(start))]
-pub async fn run() -> anyhow::Result<()> {
+pub async fn run() {
     cfg_if::cfg_if! {
         if #[cfg(target_arch = "wasm32")] {
             std::panic::set_hook(Box::new(console_error_panic_hook::hook));
@@ -828,7 +838,7 @@ pub async fn run() -> anyhow::Result<()> {
             .expect("Couldn't append canvas to document body.");
     }
 
-    let mut state = State::new(window).await?; // NEW!
+    let mut state = State::new(window).await.unwrap(); // NEW!
     let mut last_render_time = instant::Instant::now();
     event_loop.run(move |event, _, control_flow| {
         *control_flow = ControlFlow::Poll;

code/intermediate/tutorial13-hdr/src/light.wgsl

@@ -5,6 +5,7 @@ struct Camera {
     view: mat4x4<f32>,
     view_proj: mat4x4<f32>,
     inv_proj: mat4x4<f32>,
+    inv_view: mat4x4<f32>,
 }
 @group(0) @binding(0)
 var<uniform> camera: Camera;

code/intermediate/tutorial13-hdr/src/main.rs

@@ -1,5 +1,5 @@
 use tutorial13_hdr::run;
 
 fn main() {
-    pollster::block_on(run()).unwrap();
+    pollster::block_on(run());
 }

code/intermediate/tutorial13-hdr/src/resources.rs

@@ -358,9 +358,10 @@ impl HdrLoader {
         let mut encoder = device.create_command_encoder(&Default::default());
         let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label });
 
+        let num_workgroups = (dst_size + 15) / 16;
         pass.set_pipeline(&self.equirect_to_cubemap);
         pass.set_bind_group(0, &bind_group, &[]);
-        pass.dispatch_workgroups(dst_size, dst_size, 6);
+        pass.dispatch_workgroups(num_workgroups, num_workgroups, 6);
 
         drop(pass);
 

code/intermediate/tutorial13-hdr/src/shader.wgsl

@@ -133,7 +133,8 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     let reflection = textureSample(env_map, env_sampler, world_reflect).rgb;
     let shininess = 0.1;
 
-    let result = (diffuse_color + specular_color) * object_color.xyz + reflection * shininess;
+    // let result = (diffuse_color + specular_color) * object_color.xyz + reflection * shininess;
+    let result = reflection;
 
     return vec4<f32>(result, object_color.a);
 }

code/intermediate/tutorial13-hdr/src/sky.wgsl

@@ -29,9 +29,8 @@ fn vs_main(
         (id >> 1u) & 1u,
     ));
     var out: VertexOutput;
-    // out.clip_position = vec4(uv * vec2(4.0, -4.0) + vec2(-1.0, 1.0), 0.0, 1.0);
     out.clip_position = vec4(uv * 4.0 - 1.0, 1.0, 1.0);
-    out.frag_position = vec4(uv * 4.0 - 1.0, 1.0, 1.0);
+    out.frag_position = out.clip_position;
     return out;
 }
 
@@ -40,10 +39,7 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     let view_pos_homogeneous = camera.inv_proj * in.clip_position;
     let view_ray_direction = view_pos_homogeneous.xyz / view_pos_homogeneous.w;
     var ray_direction = normalize((camera.inv_view * vec4(view_ray_direction, 0.0)).xyz);
-    // ray_direction.z *= -1.0;
 
-    // let sample = vec4(ray_direction, 1.0);
     let sample = textureSample(env_map, env_sampler, ray_direction);
-    // let sample = in.clip_position;
     return sample;
 }

code/intermediate/tutorial13-hdr/src/texture.rs

@@ -181,15 +181,6 @@ impl Texture {
     }
 }
 
-pub enum CubeSide {
-    PosX = 0,
-    NegX = 1,
-    PosY = 2,
-    NegY = 3,
-    PosZ = 4,
-    NegZ = 5,
-}
-
 pub struct CubeTexture {
     texture: wgpu::Texture,
     sampler: wgpu::Sampler,

docs/intermediate/tutorial13-hdr/equirectangular.svg

@@ -0,0 +1,118 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<!DOCTYPE svg PUBLIC "-//W3C//DTD SVG 1.1//EN" "http://www.w3.org/Graphics/SVG/1.1/DTD/svg11.dtd">
+<svg version="1.2" width="102.67mm" height="55.04mm" viewBox="5492 2852 10267 5504" preserveAspectRatio="xMidYMid" fill-rule="evenodd" stroke-width="28.222" stroke-linejoin="round" xmlns="http://www.w3.org/2000/svg" xmlns:ooo="http://xml.openoffice.org/svg/export" xmlns:xlink="http://www.w3.org/1999/xlink" xmlns:presentation="http://sun.com/xmlns/staroffice/presentation" xmlns:smil="http://www.w3.org/2001/SMIL20/" xmlns:anim="urn:oasis:names:tc:opendocument:xmlns:animation:1.0" xml:space="preserve">
+ <defs>
+  <font id="EmbeddedFont_1" horiz-adv-x="2048">
+   <font-face font-family="Liberation Sans embedded" units-per-em="2048" font-weight="normal" font-style="normal" ascent="1852" descent="423"/>
+   <missing-glyph horiz-adv-x="2048" d="M 0,0 L 2047,0 2047,2047 0,2047 0,0 Z"/>
+   <glyph unicode="u" horiz-adv-x="874" d="M 314,1082 L 314,396 C 314,325 321,269 335,230 349,191 371,162 402,145 433,128 478,119 537,119 624,119 692,149 742,208 792,267 817,350 817,455 L 817,1082 997,1082 997,231 C 997,105 999,28 1003,0 L 833,0 C 832,3 832,12 831,27 830,42 830,59 829,78 828,97 826,132 825,185 L 822,185 C 781,110 733,58 679,27 624,-4 557,-20 476,-20 357,-20 271,10 216,69 161,128 133,225 133,361 L 133,1082 314,1082 Z"/>
+   <glyph unicode="t" horiz-adv-x="531" d="M 554,8 C 495,-8 434,-16 372,-16 228,-16 156,66 156,229 L 156,951 31,951 31,1082 163,1082 216,1324 336,1324 336,1082 536,1082 536,951 336,951 336,268 C 336,216 345,180 362,159 379,138 408,127 450,127 474,127 509,132 554,141 L 554,8 Z"/>
+   <glyph unicode="r" horiz-adv-x="530" d="M 142,0 L 142,830 C 142,906 140,990 136,1082 L 306,1082 C 311,959 314,886 314,861 L 318,861 C 347,954 380,1017 417,1051 454,1085 507,1102 575,1102 599,1102 623,1099 648,1092 L 648,927 C 624,934 592,937 552,937 477,937 420,905 381,841 342,776 322,684 322,564 L 322,0 142,0 Z"/>
+   <glyph unicode="q" horiz-adv-x="927" d="M 484,-20 C 347,-20 246,26 182,119 118,212 86,351 86,536 86,913 219,1102 484,1102 566,1102 634,1088 687,1059 740,1030 785,981 821,914 L 823,914 C 823,934 824,969 827,1018 830,1067 832,1093 835,1096 L 1008,1096 C 1003,1057 1001,958 1001,801 L 1001,-425 821,-425 821,14 825,178 823,178 C 787,107 743,56 690,26 637,-5 569,-20 484,-20 Z M 821,554 C 821,695 798,799 752,867 706,935 633,969 532,969 441,969 375,935 335,867 295,799 275,691 275,542 275,391 295,282 336,217 376,152 441,119 530,119 632,119 706,155 752,228 798,301 821,409 821,554 Z"/>
+   <glyph unicode="o" horiz-adv-x="980" d="M 1053,542 C 1053,353 1011,212 928,119 845,26 724,-20 565,-20 407,-20 288,28 207,125 126,221 86,360 86,542 86,915 248,1102 571,1102 736,1102 858,1057 936,966 1014,875 1053,733 1053,542 Z M 864,542 C 864,691 842,800 798,868 753,935 679,969 574,969 469,969 393,935 346,866 299,797 275,689 275,542 275,399 298,292 345,221 391,149 464,113 563,113 671,113 748,148 795,217 841,286 864,395 864,542 Z"/>
+   <glyph unicode="n" horiz-adv-x="874" d="M 825,0 L 825,686 C 825,757 818,813 804,852 790,891 768,920 737,937 706,954 661,963 602,963 515,963 447,933 397,874 347,815 322,732 322,627 L 322,0 142,0 142,851 C 142,977 140,1054 136,1082 L 306,1082 C 307,1079 307,1070 308,1055 309,1040 310,1024 311,1005 312,986 313,950 314,897 L 317,897 C 358,972 406,1025 461,1056 515,1087 582,1102 663,1102 782,1102 869,1073 924,1014 979,955 1006,857 1006,721 L 1006,0 825,0 Z"/>
+   <glyph unicode="m" horiz-adv-x="1457" d="M 768,0 L 768,686 C 768,791 754,863 725,903 696,943 645,963 570,963 493,963 433,934 388,875 343,816 321,734 321,627 L 321,0 142,0 142,851 C 142,977 140,1054 136,1082 L 306,1082 C 307,1079 307,1070 308,1055 309,1040 310,1024 311,1005 312,986 313,950 314,897 L 317,897 C 356,974 400,1027 450,1057 500,1087 561,1102 633,1102 715,1102 780,1086 828,1053 875,1020 908,968 927,897 L 930,897 C 967,970 1013,1022 1066,1054 1119,1086 1183,1102 1258,1102 1367,1102 1447,1072 1497,1013 1546,954 1571,856 1571,721 L 1571,0 1393,0 1393,686 C 1393,791 1379,863 1350,903 1321,943 1270,963 1195,963 1116,963 1055,934 1012,876 968,817 946,734 946,627 L 946,0 768,0 Z"/>
+   <glyph unicode="l" horiz-adv-x="187" d="M 138,0 L 138,1484 318,1484 318,0 138,0 Z"/>
+   <glyph unicode="i" horiz-adv-x="187" d="M 137,1312 L 137,1484 317,1484 317,1312 137,1312 Z M 137,0 L 137,1082 317,1082 317,0 137,0 Z"/>
+   <glyph unicode="h" horiz-adv-x="874" d="M 317,897 C 356,968 402,1020 457,1053 511,1086 580,1102 663,1102 780,1102 867,1073 923,1015 978,956 1006,858 1006,721 L 1006,0 825,0 825,686 C 825,762 818,819 804,856 790,893 767,920 735,937 703,954 659,963 602,963 517,963 450,934 399,875 348,816 322,737 322,638 L 322,0 142,0 142,1484 322,1484 322,1098 C 322,1057 321,1015 319,972 316,929 315,904 314,897 L 317,897 Z"/>
+   <glyph unicode="e" horiz-adv-x="980" d="M 276,503 C 276,379 302,283 353,216 404,149 479,115 578,115 656,115 719,131 766,162 813,193 844,233 861,281 L 1019,236 C 954,65 807,-20 578,-20 418,-20 296,28 213,123 129,218 87,360 87,548 87,727 129,864 213,959 296,1054 416,1102 571,1102 889,1102 1048,910 1048,527 L 1048,503 276,503 Z M 862,641 C 852,755 823,838 775,891 727,943 658,969 568,969 481,969 412,940 361,882 310,823 282,743 278,641 L 862,641 Z"/>
+   <glyph unicode="d" horiz-adv-x="927" d="M 821,174 C 788,105 744,55 689,25 634,-5 565,-20 484,-20 347,-20 247,26 183,118 118,210 86,349 86,536 86,913 219,1102 484,1102 566,1102 634,1087 689,1057 744,1027 788,979 821,914 L 823,914 821,1035 821,1484 1001,1484 1001,223 C 1001,110 1003,36 1007,0 L 835,0 C 833,11 831,35 829,74 826,113 825,146 825,174 L 821,174 Z M 275,542 C 275,391 295,282 335,217 375,152 440,119 530,119 632,119 706,154 752,225 798,296 821,405 821,554 821,697 798,802 752,869 706,936 633,969 532,969 441,969 376,936 336,869 295,802 275,693 275,542 Z"/>
+   <glyph unicode="a" horiz-adv-x="1060" d="M 414,-20 C 305,-20 224,9 169,66 114,123 87,202 87,302 87,414 124,500 198,560 271,620 390,652 554,656 L 797,660 797,719 C 797,807 778,870 741,908 704,946 645,965 565,965 484,965 426,951 389,924 352,897 330,853 323,793 L 135,810 C 166,1005 310,1102 569,1102 705,1102 807,1071 876,1009 945,946 979,856 979,738 L 979,272 C 979,219 986,179 1000,152 1014,125 1041,111 1080,111 1097,111 1117,113 1139,118 L 1139,6 C 1094,-5 1047,-10 1000,-10 933,-10 885,8 855,43 824,78 807,132 803,207 L 797,207 C 751,124 698,66 637,32 576,-3 501,-20 414,-20 Z M 455,115 C 521,115 580,130 631,160 682,190 723,231 753,284 782,336 797,390 797,445 L 797,534 600,530 C 515,529 451,520 408,504 364,488 330,463 307,430 284,397 272,353 272,299 272,240 288,195 320,163 351,131 396,115 455,115 Z"/>
+   <glyph unicode="S" horiz-adv-x="1192" d="M 1272,389 C 1272,259 1221,158 1120,87 1018,16 875,-20 690,-20 347,-20 148,99 93,338 L 278,375 C 299,290 345,228 414,189 483,149 578,129 697,129 820,129 916,150 983,193 1050,235 1083,297 1083,379 1083,425 1073,462 1052,491 1031,520 1001,543 963,562 925,581 880,596 827,609 774,622 716,635 652,650 541,675 456,699 399,724 341,749 295,776 262,807 229,837 203,872 186,913 168,954 159,1000 159,1053 159,1174 205,1267 298,1332 390,1397 522,1430 694,1430 854,1430 976,1406 1061,1357 1146,1308 1205,1224 1239,1106 L 1051,1073 C 1030,1148 991,1202 933,1236 875,1269 795,1286 692,1286 579,1286 493,1267 434,1230 375,1193 345,1137 345,1063 345,1020 357,984 380,956 403,927 436,903 479,884 522,864 609,840 738,811 781,801 825,791 868,781 911,770 952,758 991,744 1030,729 1067,712 1102,693 1136,674 1166,650 1191,622 1216,594 1236,561 1251,523 1265,485 1272,440 1272,389 Z"/>
+   <glyph unicode="P" horiz-adv-x="1112" d="M 1258,985 C 1258,852 1215,746 1128,667 1041,588 922,549 773,549 L 359,549 359,0 168,0 168,1409 761,1409 C 919,1409 1041,1372 1128,1298 1215,1224 1258,1120 1258,985 Z M 1066,983 C 1066,1165 957,1256 738,1256 L 359,1256 359,700 746,700 C 959,700 1066,794 1066,983 Z"/>
+   <glyph unicode="N" horiz-adv-x="1165" d="M 1082,0 L 328,1200 333,1103 338,936 338,0 168,0 168,1409 390,1409 1152,201 C 1144,332 1140,426 1140,485 L 1140,1409 1312,1409 1312,0 1082,0 Z"/>
+   <glyph unicode="M" horiz-adv-x="1377" d="M 1366,0 L 1366,940 C 1366,1044 1369,1144 1375,1240 1342,1121 1313,1027 1287,960 L 923,0 789,0 420,960 364,1130 331,1240 334,1129 338,940 338,0 168,0 168,1409 419,1409 794,432 C 807,393 820,351 833,306 845,261 853,228 857,208 862,235 874,275 891,330 908,384 919,418 925,432 L 1293,1409 1538,1409 1538,0 1366,0 Z"/>
+   <glyph unicode="E" horiz-adv-x="1138" d="M 168,0 L 168,1409 1237,1409 1237,1253 359,1253 359,801 1177,801 1177,647 359,647 359,156 1278,156 1278,0 168,0 Z"/>
+   <glyph unicode=" " horiz-adv-x="556"/>
+  </font>
+ </defs>
+ <defs class="EmbeddedBulletChars">
+  <g id="bullet-char-template-57356" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 580,1141 L 1163,571 580,0 -4,571 580,1141 Z"/>
+  </g>
+  <g id="bullet-char-template-57354" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 8,1128 L 1137,1128 1137,0 8,0 8,1128 Z"/>
+  </g>
+  <g id="bullet-char-template-10146" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 174,0 L 602,739 174,1481 1456,739 174,0 Z M 1358,739 L 309,1346 659,739 1358,739 Z"/>
+  </g>
+  <g id="bullet-char-template-10132" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 2015,739 L 1276,0 717,0 1260,543 174,543 174,936 1260,936 717,1481 1274,1481 2015,739 Z"/>
+  </g>
+  <g id="bullet-char-template-10007" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 0,-2 C -7,14 -16,27 -25,37 L 356,567 C 262,823 215,952 215,954 215,979 228,992 255,992 264,992 276,990 289,987 310,991 331,999 354,1012 L 381,999 492,748 772,1049 836,1024 860,1049 C 881,1039 901,1025 922,1006 886,937 835,863 770,784 769,783 710,716 594,584 L 774,223 C 774,196 753,168 711,139 L 727,119 C 717,90 699,76 672,76 641,76 570,178 457,381 L 164,-76 C 142,-110 111,-127 72,-127 30,-127 9,-110 8,-76 1,-67 -2,-52 -2,-32 -2,-23 -1,-13 0,-2 Z"/>
+  </g>
+  <g id="bullet-char-template-10004" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 285,-33 C 182,-33 111,30 74,156 52,228 41,333 41,471 41,549 55,616 82,672 116,743 169,778 240,778 293,778 328,747 346,684 L 369,508 C 377,444 397,411 428,410 L 1163,1116 C 1174,1127 1196,1133 1229,1133 1271,1133 1292,1118 1292,1087 L 1292,965 C 1292,929 1282,901 1262,881 L 442,47 C 390,-6 338,-33 285,-33 Z"/>
+  </g>
+  <g id="bullet-char-template-9679" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 813,0 C 632,0 489,54 383,161 276,268 223,411 223,592 223,773 276,916 383,1023 489,1130 632,1184 813,1184 992,1184 1136,1130 1245,1023 1353,916 1407,772 1407,592 1407,412 1353,268 1245,161 1136,54 992,0 813,0 Z"/>
+  </g>
+  <g id="bullet-char-template-8226" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 346,457 C 273,457 209,483 155,535 101,586 74,649 74,723 74,796 101,859 155,911 209,963 273,989 346,989 419,989 480,963 531,910 582,859 608,796 608,723 608,648 583,586 532,535 482,483 420,457 346,457 Z"/>
+  </g>
+  <g id="bullet-char-template-8211" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M -4,459 L 1135,459 1135,606 -4,606 -4,459 Z"/>
+  </g>
+  <g id="bullet-char-template-61548" transform="scale(0.00048828125,-0.00048828125)">
+   <path d="M 173,740 C 173,903 231,1043 346,1159 462,1274 601,1332 765,1332 928,1332 1067,1274 1183,1159 1299,1043 1357,903 1357,740 1357,577 1299,437 1183,322 1067,206 928,148 765,148 601,148 462,206 346,322 231,437 173,577 173,740 Z"/>
+  </g>
+ </defs>
+ <g class="Page">
+  <g class="com.sun.star.drawing.CustomShape">
+   <g id="id3">
+    <rect class="BoundingBox" stroke="none" fill="none" x="5544" y="2903" width="10163" height="5400"/>
+    <path fill="rgb(222,230,239)" stroke="none" d="M 10625,8301 L 5545,8301 5545,2904 15705,2904 15705,8301 10625,8301 Z"/>
+    <path fill="none" stroke="rgb(222,230,239)" d="M 10625,8301 L 5545,8301 5545,2904 15705,2904 15705,8301 10625,8301 Z"/>
+   </g>
+  </g>
+  <g class="com.sun.star.drawing.LineShape">
+   <g id="id4">
+    <rect class="BoundingBox" stroke="none" fill="none" x="10572" y="2852" width="107" height="5504"/>
+    <path fill="none" stroke="rgb(52,101,164)" stroke-width="106" stroke-linejoin="round" d="M 10625,2905 L 10625,8302"/>
+   </g>
+  </g>
+  <g class="com.sun.star.drawing.LineShape">
+   <g id="id5">
+    <rect class="BoundingBox" stroke="none" fill="none" x="5492" y="5592" width="10267" height="107"/>
+    <path fill="none" stroke="rgb(241,13,12)" stroke-width="106" stroke-linejoin="round" d="M 5545,5645 L 15705,5645"/>
+   </g>
+  </g>
+  <g class="TextShape">
+   <g id="id6">
+    <rect class="BoundingBox" stroke="none" fill="none" x="9999" y="3139" width="527" height="2224"/>
+    <text class="SVGTextShape" transform="rotate(-90 10348 5072)"><tspan class="TextParagraph" font-family="Liberation Sans, sans-serif" font-size="247px" font-weight="400"><tspan class="TextPosition" x="10348" y="5072"><tspan fill="rgb(0,0,0)" stroke="none" style="white-space: pre">Prime Meridian</tspan></tspan></tspan></text>
+   </g>
+  </g>
+  <g class="TextShape">
+   <g id="id7">
+    <rect class="BoundingBox" stroke="none" fill="none" x="7163" y="5127" width="1588" height="527"/>
+    <text class="SVGTextShape"><tspan class="TextParagraph" font-family="Liberation Sans, sans-serif" font-size="247px" font-weight="400"><tspan class="TextPosition" x="7529" y="5476"><tspan fill="rgb(0,0,0)" stroke="none" style="white-space: pre">Equator</tspan></tspan></tspan></text>
+   </g>
+  </g>
+  <g class="TextShape">
+   <g id="id8">
+    <rect class="BoundingBox" stroke="none" fill="none" x="7015" y="2927" width="2037" height="527"/>
+    <text class="SVGTextShape"><tspan class="TextParagraph" font-family="Liberation Sans, sans-serif" font-size="247px" font-weight="400"><tspan class="TextPosition" x="7455" y="3276"><tspan fill="rgb(0,0,0)" stroke="none" style="white-space: pre">North Pole</tspan></tspan></tspan></text>
+   </g>
+  </g>
+  <g class="TextShape">
+   <g id="id9">
+    <rect class="BoundingBox" stroke="none" fill="none" x="6915" y="7727" width="2037" height="527"/>
+    <text class="SVGTextShape"><tspan class="TextParagraph" font-family="Liberation Sans, sans-serif" font-size="247px" font-weight="400"><tspan class="TextPosition" x="7334" y="8076"><tspan fill="rgb(0,0,0)" stroke="none" style="white-space: pre">South Pole</tspan></tspan></tspan></text>
+   </g>
+  </g>
+  <g class="TextShape">
+   <g id="id10">
+    <rect class="BoundingBox" stroke="none" fill="none" x="12028" y="2927" width="2037" height="527"/>
+    <text class="SVGTextShape"><tspan class="TextParagraph" font-family="Liberation Sans, sans-serif" font-size="247px" font-weight="400"><tspan class="TextPosition" x="12468" y="3276"><tspan fill="rgb(0,0,0)" stroke="none" style="white-space: pre">North Pole</tspan></tspan></tspan></text>
+   </g>
+  </g>
+  <g class="TextShape">
+   <g id="id11">
+    <rect class="BoundingBox" stroke="none" fill="none" x="12015" y="7727" width="2037" height="527"/>
+    <text class="SVGTextShape"><tspan class="TextParagraph" font-family="Liberation Sans, sans-serif" font-size="247px" font-weight="400"><tspan class="TextPosition" x="12434" y="8076"><tspan fill="rgb(0,0,0)" stroke="none" style="white-space: pre">South Pole</tspan></tspan></tspan></text>
+   </g>
+  </g>
+ </g>
+</svg>

docs/intermediate/tutorial13-hdr/readme.md

@@ -27,10 +27,18 @@ greater than 1.0 meaning you can have a dynamic range of brighter objects.
 ## Switching to HDR
 
 As of writing, wgpu doesn't allow us to use a floating point format such as
-`TextureFormat::Rgba16Float` (not all monitors support that anyways), so we
-will have to render our scene in an HDR format, then convert the values to a
-supported format such as `TextureFormat::Bgra8UnormSrgb` using a technique
-called tonemapping.
+`TextureFormat::Rgba16Float` as the surface texture format (not all
+monitors support that anyways), so we will have to render our scene in
+an HDR format, then convert the values to a supported format such as
+`TextureFormat::Bgra8UnormSrgb` using a technique called tonemapping.
+
+<div class="note">
+
+There are some talks about implementing HDR surface texture support in
+wgpu. Here is a github issues if you want to contribute to that
+effort: https://github.com/gfx-rs/wgpu/issues/2920
+
+</div>
 
 Before we do that though we need to switch to using an HDR texture for rendering.
 
@@ -386,18 +394,966 @@ Here's what it looks like after implementing HDR:
 
 ## Loading HDR textures
 
-## Equirectangular textures and cube maps
+Now that we have an HDR render buffer, we can start leveraging
+HDR textures to their fullest. One of the main uses for HDR
+textures is to store lighting information in the form of an
+environment map.
+
+This map can be used to light objects, display reflections and
+also to make a skybox. We're going to create a skybox using HDR
+texture, but first we need to talk about how environment maps are
+stored.
+
+## Equirectangular textures
+
+An equirectangluar texture is a texture where a sphere is stretched
+across a rectangular surface using what's known as an equirectangular
+projection. This map of the Earth is an example of this projection.
+
+![map of the earth](https://upload.wikimedia.org/wikipedia/commons/thumb/8/83/Equirectangular_projection_SW.jpg/1024px-Equirectangular_projection_SW.jpg)
+
+This projection maps the latitude values of the sphere to the
+horizontal coordinates of the texture. The longitude values get
+mapped to the vertical coordinates. This means that the vertical
+middle of the texture is the equator (0° longitude) of the sphere,
+the horizontal middle is the prime meridian (0° latitude) of the
+sphere, the left and right edges of the texture are the anti-meridian
+(+180°/-180° latitude) the top and bottom edges of the texture are
+the north pole (90° longitude) and south pole (-90° longitude)
+respectively.
+
+![equirectangular diagram](./equirectangular.svg)
+
+This simple projection is easy to use, leading it to be one of the
+most popular projections for storing spherical textures. You can
+see the particular environment map we are going to use below.
 
 ![equirectangular skybox](./kloofendal_43d_clear_puresky.jpg)
 
+## Cube Maps
+
+While we technically can use an equirectangular map directly as long
+as we do some math to figure out the correct coordinates, it is a lot
+more convenient to convert our environment map into a cube map.
+
+<div class="info">
+
+A cube map is special kind of texture that has 6 layers. Each layer
+corresponds to a different face of an imaginary cube that is aligned
+to the X, Y and Z axes. The layers are stored in the following order:
++X, -X, +Y, -Y, +Z, -Z. 
+
+</div>
+
+To prepare to store the cube texture, we are going to create
+a new struct called `CubeTexture` in `texture.rs`.
+
+```rust
+pub struct CubeTexture {
+    texture: wgpu::Texture,
+    sampler: wgpu::Sampler,
+    view: wgpu::TextureView,
+}
+
+impl CubeTexture {
+    pub fn create_2d(
+        device: &wgpu::Device,
+        width: u32,
+        height: u32,
+        format: wgpu::TextureFormat,
+        mip_level_count: u32,
+        usage: wgpu::TextureUsages,
+        mag_filter: wgpu::FilterMode,
+        label: Option<&str>,
+    ) -> Self {
+        let texture = device.create_texture(&wgpu::TextureDescriptor {
+            label,
+            size: wgpu::Extent3d {
+                width,
+                height,
+                // A cube has 6 sides, so we need 6 layers
+                depth_or_array_layers: 6,
+            },
+            mip_level_count,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format,
+            usage,
+            view_formats: &[],
+        });
+
+        let view = texture.create_view(&wgpu::TextureViewDescriptor {
+            label,
+            dimension: Some(wgpu::TextureViewDimension::Cube),
+            array_layer_count: Some(6),
+            ..Default::default()
+        });
+
+        let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
+            label,
+            address_mode_u: wgpu::AddressMode::ClampToEdge,
+            address_mode_v: wgpu::AddressMode::ClampToEdge,
+            address_mode_w: wgpu::AddressMode::ClampToEdge,
+            mag_filter,
+            min_filter: wgpu::FilterMode::Nearest,
+            mipmap_filter: wgpu::FilterMode::Nearest,
+            ..Default::default()
+        });
+
+        Self {
+            texture,
+            sampler,
+            view,
+        }
+    }
+
+    pub fn texture(&self) -> &wgpu::Texture { &self.texture }
+    
+    pub fn view(&self) -> &wgpu::TextureView { &self.view }
+
+    pub fn sampler(&self) -> &wgpu::Sampler { &self.sampler }
+
+}
+```
+
+With this we can now write the code to load the HDR into
+a cube texture.
+
 ## Compute shaders
 
+Up to this point we've been exclusively using render
+pipelines, but I felt this was a good time to introduce
+compute pipelines and by extension compute shaders. Compute
+pipelines are a lot easier to setup. All you need is to tell
+the pipeline what resources you want to use, what code you
+want to run, and how many threads you'd like the GPU to use
+when running your code. We're going to use a compute shader
+to give each pixel in our cube textue a color from the
+HDR image.
+
+Before we can use compute shaders, we need to enable them
+in wgpu. We can do that just need to change the line where
+we specify what features we want to use. In `lib.rs`, change
+the code where we request a device:
+
+```rust
+let (device, queue) = adapter
+    .request_device(
+        &wgpu::DeviceDescriptor {
+            label: None,
+            // UPDATED!
+            features: wgpu::Features::all_webgpu_mask(),
+            // UPDATED!
+            limits: wgpu::Limits::downlevel_defaults(),
+        },
+        None, // Trace path
+    )
+    .await
+    .unwrap();
+```
+
+<div class="warn">
+
+You may have noted that we have switched from
+`downlevel_webgl2_defaults()` to `downlevel_defaults()`.
+This means that we are dropping support for WebGL2. The
+reason for this is that WebGL2 doesn't support compute
+shaders. WebGPU was built with compute shaders in mind. As
+of writing the only browser that supports WebGPU is Chrome,
+and some experimental browsers such as Firefox Nightly.
+
+Consequently we are going to remove the webgl feature from
+`Cargo.toml`. This line in particular:
+
+```toml
+wgpu = { version = "0.17", features = ["webgl"]}
+```
+
+</div>
+
+Now that we've told wgpu that we want to use compute
+shaders, let's create a struct in `resource.rs` that we'll
+use to load the HDR image into our cube map.
+
+```rust
+pub struct HdrLoader {
+    texture_format: wgpu::TextureFormat,
+    equirect_layout: wgpu::BindGroupLayout,
+    equirect_to_cubemap: wgpu::ComputePipeline,
+}
+
+impl HdrLoader {
+    pub fn new(device: &wgpu::Device) -> Self {
+        let module = device.create_shader_module(wgpu::include_wgsl!("equirectangular.wgsl"));
+        let texture_format = wgpu::TextureFormat::Rgba32Float;
+        let equirect_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("HdrLoader::equirect_layout"),
+            entries: &[
+                wgpu::BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Texture {
+                        sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                        multisampled: false,
+                    },
+                    count: None,
+                },
+                wgpu::BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::StorageTexture {
+                        access: wgpu::StorageTextureAccess::WriteOnly,
+                        format: texture_format,
+                        view_dimension: wgpu::TextureViewDimension::D2Array,
+                    },
+                    count: None,
+                },
+            ],
+        });
+
+        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: None,
+            bind_group_layouts: &[&equirect_layout],
+            push_constant_ranges: &[],
+        });
+
+        let equirect_to_cubemap =
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("equirect_to_cubemap"),
+                layout: Some(&pipeline_layout),
+                module: &module,
+                entry_point: "compute_equirect_to_cubemap",
+            });
+
+        Self {
+            equirect_to_cubemap,
+            texture_format,
+            equirect_layout,
+        }
+    }
+
+    pub fn from_equirectangular_bytes(
+        &self,
+        device: &wgpu::Device,
+        queue: &wgpu::Queue,
+        data: &[u8],
+        dst_size: u32,
+        label: Option<&str>,
+    ) -> anyhow::Result<texture::CubeTexture> {
+        let hdr_decoder = HdrDecoder::new(Cursor::new(data))?;
+        let meta = hdr_decoder.metadata();
+        let mut pixels = vec![[0.0, 0.0, 0.0, 0.0]; meta.width as usize * meta.height as usize];
+        hdr_decoder.read_image_transform(
+            |pix| {
+                // There's no Rgb32Float format, so we need
+                // an extra float
+                let rgb = pix.to_hdr();
+                [rgb.0[0], rgb.0[1], rgb.0[2], 1.0f32]
+            },
+            &mut pixels[..],
+        )?;
+
+        let src = texture::Texture::create_2d_texture(
+            device,
+            meta.width,
+            meta.height,
+            self.texture_format,
+            wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST,
+            wgpu::FilterMode::Linear,
+            None,
+        );
+
+        queue.write_texture(
+            wgpu::ImageCopyTexture {
+                texture: &src.texture,
+                mip_level: 0,
+                origin: wgpu::Origin3d::ZERO,
+                aspect: wgpu::TextureAspect::All,
+            },
+            &bytemuck::cast_slice(&pixels),
+            wgpu::ImageDataLayout {
+                offset: 0,
+                bytes_per_row: Some(src.size.width * std::mem::size_of::<[f32; 4]>() as u32),
+                rows_per_image: Some(src.size.height),
+            },
+            src.size,
+        );
+
+        let dst = texture::CubeTexture::create_2d(
+            device,
+            dst_size,
+            dst_size,
+            self.texture_format,
+            1,
+            // We are going to write to `dst` texture so we
+            // need to use a `STORAGE_BINDING`.
+            wgpu::TextureUsages::STORAGE_BINDING
+                | wgpu::TextureUsages::TEXTURE_BINDING,
+            wgpu::FilterMode::Nearest,
+            label,
+        );
+
+        let dst_view = dst.texture().create_view(&wgpu::TextureViewDescriptor {
+            label,
+            // Normally you'd use `TextureViewDimension::Cube`
+            // for a cube texture, but we can't use that
+            // view dimension with a `STORAGE_BINDING`.
+            // We need to access the cube texure layers
+            // directly.
+            dimension: Some(wgpu::TextureViewDimension::D2Array),
+            ..Default::default()
+        });
+
+        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label,
+            layout: &self.equirect_layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(&src.view),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::TextureView(&dst_view),
+                },
+            ],
+        });
+
+        let mut encoder = device.create_command_encoder(&Default::default());
+        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label });
+
+        let num_workgroups = (dst_size + 15) / 16;
+        pass.set_pipeline(&self.equirect_to_cubemap);
+        pass.set_bind_group(0, &bind_group, &[]);
+        pass.dispatch_workgroups(num_workgroups, num_workgroups, 6);
+
+        drop(pass);
+
+        queue.submit([encoder.finish()]);
+
+        Ok(dst)
+    }
+}
+```
+
+The `dispatch_workgroups` call tells the gpu to run our
+code in batchs called workgroups. Each workgroup has a
+number of worker threads called invocations that run the
+code in parallel. Workgroups are organized as a 3d grid
+with the dimensions we pass to `dispatch_workgroups`.
+
+In this example we have a workgroup grid divided into 16x16
+chunks and storing the layer in z dimension.
+
+## The compute shader
+
+Now let's write a compute shader that will convert
+our equirectangular texture to a cube texture. Create a file
+called `equirectangular.wgsl`. We're going to break it down
+chunk by chunk.
+
+```wgsl
+const PI: f32 = 3.1415926535897932384626433832795;
+
+struct Face {
+    forward: vec3<f32>,
+    up: vec3<f32>,
+    right: vec3<f32>,
+}
+```
+
+Two things here:
+
+1. wgsl doesn't have a builtin for PI so we need to specify
+    it ourselves.
+2. each face of the cube map has an orientation to it, so we
+    need to store that.
+
+```wgsl
+@group(0)
+@binding(0)
+var src: texture_2d<f32>;
+
+@group(0)
+@binding(1)
+var dst: texture_storage_2d_array<rgba32float, write>;
+```
+
+Here we have the only two bindings we need. The equirectangular
+`src` texture and our `dst` cube texture. Some things to note:
+about `dst`:
+
+1. While `dst` is a cube texture, it's stored as a array of
+    2d textures.
+2. The type of binding we're using here is a storage texture.
+    An array storage texture to be precise. This is a unique
+    binding only available to compute shaders. It allows us
+    to directly write to the texture.
+3. When using a storage texture binding we need to specify the
+    format of the texture. If you try to bind a texture with
+    a different format, wgpu will panic.
+
+```wgsl
+@compute
+@workgroup_size(16, 16, 1)
+fn compute_equirect_to_cubemap(
+    @builtin(global_invocation_id)
+    gid: vec3<u32>,
+) {
+    // If texture size is not divisible by 32 we
+    // need to make sure we don't try to write to
+    // pixels that don't exist.
+    if gid.x >= u32(textureDimensions(dst).x) {
+        return;
+    }
+
+    var FACES: array<Face, 6> = array(
+        // FACES +X
+        Face(
+            vec3(1.0, 0.0, 0.0),  // forward
+            vec3(0.0, 1.0, 0.0),  // up
+            vec3(0.0, 0.0, -1.0), // right
+        ),
+        // FACES -X
+        Face (
+            vec3(-1.0, 0.0, 0.0),
+            vec3(0.0, 1.0, 0.0),
+            vec3(0.0, 0.0, 1.0),
+        ),
+        // FACES +Y
+        Face (
+            vec3(0.0, -1.0, 0.0),
+            vec3(0.0, 0.0, 1.0),
+            vec3(1.0, 0.0, 0.0),
+        ),
+        // FACES -Y
+        Face (
+            vec3(0.0, 1.0, 0.0),
+            vec3(0.0, 0.0, -1.0),
+            vec3(1.0, 0.0, 0.0),
+        ),
+        // FACES +Z
+        Face (
+            vec3(0.0, 0.0, 1.0),
+            vec3(0.0, 1.0, 0.0),
+            vec3(1.0, 0.0, 0.0),
+        ),
+        // FACES -Z
+        Face (
+            vec3(0.0, 0.0, -1.0),
+            vec3(0.0, 1.0, 0.0),
+            vec3(-1.0, 0.0, 0.0),
+        ),
+    );
+
+    // Get texture coords relative to cubemap face
+    let dst_dimensions = vec2<f32>(textureDimensions(dst));
+    let cube_uv = vec2<f32>(gid.xy) / dst_dimensions * 2.0 - 1.0;
+
+    // Get spherical coordinate from cube_uv
+    let face = FACES[gid.z];
+    let spherical = normalize(face.forward + face.right * cube_uv.x + face.up * cube_uv.y);
+
+    // Get coordinate on the equirectangular texture
+    let inv_atan = vec2(0.1591, 0.3183);
+    let eq_uv = vec2(atan2(spherical.z, spherical.x), asin(spherical.y)) * inv_atan + 0.5;
+    let eq_pixel = vec2<i32>(eq_uv * vec2<f32>(textureDimensions(src)));
+
+    // We use textureLoad() as textureSample() is not allowed in compute shaders
+    var sample = textureLoad(src, eq_pixel, 0);
+
+    textureStore(dst, gid.xy, gid.z, sample);
+}
+```
+
+While I commented some the previous code, there are some
+things I want to go over that wouldn't fit well in a
+comment.
+
+The `workgroup_size` decorator tells the dimensions of the
+workgroup's local grid of invocations. Because we are
+dispatching one workgroup for every pixel in the texture,
+we have each workgroup be a 16x16x1 grid. This means that each workgroup can have 256 threads to work with.
+
+<div class="warn">
+
+For Webgpu each workgroup can only have a max of 256 threads (also
+called invocations).
+
+</div>
+
+With this we can load the environment map in the `new()` function:
+
+```rust
+let hdr_loader = resources::HdrLoader::new(&device);
+let sky_bytes = resources::load_binary("pure-sky.hdr").await?;
+let sky_texture = hdr_loader.from_equirectangular_bytes(
+    &device,
+    &queue,
+    &sky_bytes,
+    1080,
+    Some("Sky Texture"),
+)?;
+```
+
 ## Skybox
 
-![debugging skybox](./debugging-skybox.png)
+No that we have an environment map to render. Let's use
+it to make our skybox. There are different ways to render
+a skybox. A standard way is to render a cube and map the
+environment map on it. While that method works, it can
+have some artifacts in the corners and edges where the
+cubes faces meet.
+
+Instead we are going to render to the entire screen and
+compute the view direction from each pixel, and use that
+to sample the texture. First though we need to create a
+bindgroup for the environment map so that we can use it
+for rendering. Add the following to `new()`:
+
+```rust
+let environment_layout =
+    device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+        label: Some("environment_layout"),
+        entries: &[
+            wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::FRAGMENT,
+                ty: wgpu::BindingType::Texture {
+                    sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                    view_dimension: wgpu::TextureViewDimension::Cube,
+                    multisampled: false,
+                },
+                count: None,
+            },
+            wgpu::BindGroupLayoutEntry {
+                binding: 1,
+                visibility: wgpu::ShaderStages::FRAGMENT,
+                ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::NonFiltering),
+                count: None,
+            },
+        ],
+    });
+
+let environment_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+    label: Some("environment_bind_group"),
+    layout: &environment_layout,
+    entries: &[
+        wgpu::BindGroupEntry {
+            binding: 0,
+            resource: wgpu::BindingResource::TextureView(&sky_texture.view()),
+        },
+        wgpu::BindGroupEntry {
+            binding: 1,
+            resource: wgpu::BindingResource::Sampler(sky_texture.sampler()),
+        },
+    ],
+});
+```
+
+Now that we have the bindgroup, we need a render pipeline
+to render the skybox.
+
+```rust
+// NEW!
+let sky_pipeline = {
+    let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+        label: Some("Sky Pipeline Layout"),
+        bind_group_layouts: &[&camera_bind_group_layout, &environment_layout],
+        push_constant_ranges: &[],
+    });
+    let shader = wgpu::include_wgsl!("sky.wgsl");
+    create_render_pipeline(
+        &device,
+        &layout,
+        hdr.format(),
+        Some(texture::Texture::DEPTH_FORMAT),
+        &[],
+        wgpu::PrimitiveTopology::TriangleList,
+        shader,
+    )
+};
+```
+
+One thing to not here. We added the primitive format to
+`create_render_pipeline()`. Also we changed the depth compare
+function to `CompareFunction::LessEqual` (we'll discuss why when
+we go over the sky shader). Here's the changes to that:
+
+```rust
+
+fn create_render_pipeline(
+    device: &wgpu::Device,
+    layout: &wgpu::PipelineLayout,
+    color_format: wgpu::TextureFormat,
+    depth_format: Option<wgpu::TextureFormat>,
+    vertex_layouts: &[wgpu::VertexBufferLayout],
+    topology: wgpu::PrimitiveTopology, // NEW!
+    shader: wgpu::ShaderModuleDescriptor,
+) -> wgpu::RenderPipeline {
+    let shader = device.create_shader_module(shader);
+
+    device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+        // ...
+        primitive: wgpu::PrimitiveState {
+            topology, // NEW!
+            // ...
+        },
+        depth_stencil: depth_format.map(|format| wgpu::DepthStencilState {
+            format,
+            depth_write_enabled: true,
+            depth_compare: wgpu::CompareFunction::LessEqual, // UDPATED!
+            stencil: wgpu::StencilState::default(),
+            bias: wgpu::DepthBiasState::default(),
+        }),
+        // ...
+    })
+}
+```
+
+Don't forget to add the new bindgroup and pipeline to the
+to `State`.
+
+```rust
+struct State {
+    // ...
+    // NEW!
+    hdr: hdr::HdrPipeline,
+    environment_bind_group: wgpu::BindGroup,
+    sky_pipeline: wgpu::RenderPipeline,
+}
+
+// ...
+impl State {
+    async fn new(window: Window) -> anyhow::Result<Self> {
+        // ...
+        Ok(Self {
+            // ...
+            // NEW!
+            hdr,
+            environment_bind_group,
+            sky_pipeline,
+            debug,
+        })
+    }
+}
+```
+
+Now let's cover `sky.wgsl`.
+
+```wgsl
+struct Camera {
+    view_pos: vec4<f32>,
+    view: mat4x4<f32>,
+    view_proj: mat4x4<f32>,
+    inv_proj: mat4x4<f32>,
+    inv_view: mat4x4<f32>,
+}
+@group(0) @binding(0)
+var<uniform> camera: Camera;
+
+@group(1)
+@binding(0)
+var env_map: texture_cube<f32>;
+@group(1)
+@binding(1)
+var env_sampler: sampler;
+
+struct VertexOutput {
+    @builtin(position) frag_position: vec4<f32>,
+    @location(0) clip_position: vec4<f32>,
+}
+
+@vertex
+fn vs_main(
+    @builtin(vertex_index) id: u32,
+) -> VertexOutput {
+    let uv = vec2<f32>(vec2<u32>(
+        id & 1u,
+        (id >> 1u) & 1u,
+    ));
+    var out: VertexOutput;
+    // out.clip_position = vec4(uv * vec2(4.0, -4.0) + vec2(-1.0, 1.0), 0.0, 1.0);
+    out.clip_position = vec4(uv * 4.0 - 1.0, 1.0, 1.0);
+    out.frag_position = vec4(uv * 4.0 - 1.0, 1.0, 1.0);
+    return out;
+}
+
+@fragment
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
+    let view_pos_homogeneous = camera.inv_proj * in.clip_position;
+    let view_ray_direction = view_pos_homogeneous.xyz / view_pos_homogeneous.w;
+    var ray_direction = normalize((camera.inv_view * vec4(view_ray_direction, 0.0)).xyz);
+
+    let sample = textureSample(env_map, env_sampler, ray_direction);
+    return sample;
+}
+```
+
+Let's break this down:
+
+1. We create a triangle twice the size of the screen.
+2. In the fragment shader we get the view direction from
+    the clip position. We use the inverse projection
+    matrix to get convert the clip coordinates to view
+    direction. Then we use the inverse view matrix to
+    get the direction into world space as that's what we
+    need for to sample the sky box correctly.
+3. We then sample the sky texture with the view direction.
+
+<!-- ![debugging skybox](./debugging-skybox.png) -->
+
+In order for this to work we need to change our camera
+uniforms a bit. We need to add the inverse view matrix,
+and inverse projection matrix to `CameraUniform` struct.
+
+```rust
+#[repr(C)]
+#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+struct CameraUniform {
+    view_position: [f32; 4],
+    view: [[f32; 4]; 4], // NEW!
+    view_proj: [[f32; 4]; 4],
+    inv_proj: [[f32; 4]; 4], // NEW!
+    inv_view: [[f32; 4]; 4], // NEW!
+}
+
+impl CameraUniform {
+    fn new() -> Self {
+        Self {
+            view_position: [0.0; 4],
+            view: cgmath::Matrix4::identity().into(),
+            view_proj: cgmath::Matrix4::identity().into(),
+            inv_proj: cgmath::Matrix4::identity().into(), // NEW!
+            inv_view: cgmath::Matrix4::identity().into(), // NEW!
+        }
+    }
+
+    // UPDATED!
+    fn update_view_proj(&mut self, camera: &camera::Camera, projection: &camera::Projection) {
+        self.view_position = camera.position.to_homogeneous().into();
+        let proj = projection.calc_matrix();
+        let view = camera.calc_matrix();
+        let view_proj = proj * view;
+        self.view = view.into();
+        self.view_proj = view_proj.into();
+        self.inv_proj = proj.invert().unwrap().into();
+        self.inv_view = view.transpose().into();
+    }
+}
+```
+
+Make sure to change the `Camera` definition in
+`shader.wgsl`, and `light.wgsl`. Just as a reminder
+it looks like this:
+
+```wgsl
+struct Camera {
+    view_pos: vec4<f32>,
+    view: mat4x4<f32>,
+    view_proj: mat4x4<f32>,
+    inv_proj: mat4x4<f32>,
+    inv_view: mat4x4<f32>,
+}
+var<uniform> camera: Camera;
+```
+
+<div class="info">
+
+You may have noticed that we removed the `OPENGL_TO_WGPU_MATRIX`. The reason for this is
+that it was messing with the projection of the
+skybox.
 
 ![projection error](./project-error.png)
 
+It wasn't technically needed, so I felt fine
+removing it.
+
+</div>
+
 ## Reflections
 
-![with-reflections](./with-reflections.png)
+Now that we have a sky, we can mess around with
+using it for lighting. This won't be physically
+accurate (we'll look into that later). That being
+said, we have the environment map, we might as
+well use it.
+
+In order to do that though we need to change our
+shader to do lighting in world space instead of
+tangent space because our environment map is in
+world space. Because there are a lot of changes
+I'll post the whole shader here:
+
+```wgsl
+// Vertex shader
+
+struct Camera {
+    view_pos: vec4<f32>,
+    view: mat4x4<f32>,
+    view_proj: mat4x4<f32>,
+    inv_proj: mat4x4<f32>,
+    inv_view: mat4x4<f32>,
+}
+@group(0) @binding(0)
+var<uniform> camera: Camera;
+
+struct Light {
+    position: vec3<f32>,
+    color: vec3<f32>,
+}
+@group(2) @binding(0)
+var<uniform> light: Light;
+
+struct VertexInput {
+    @location(0) position: vec3<f32>,
+    @location(1) tex_coords: vec2<f32>,
+    @location(2) normal: vec3<f32>,
+    @location(3) tangent: vec3<f32>,
+    @location(4) bitangent: vec3<f32>,
+}
+struct InstanceInput {
+    @location(5) model_matrix_0: vec4<f32>,
+    @location(6) model_matrix_1: vec4<f32>,
+    @location(7) model_matrix_2: vec4<f32>,
+    @location(8) model_matrix_3: vec4<f32>,
+    @location(9) normal_matrix_0: vec3<f32>,
+    @location(10) normal_matrix_1: vec3<f32>,
+    @location(11) normal_matrix_2: vec3<f32>,
+}
+
+struct VertexOutput {
+    @builtin(position) clip_position: vec4<f32>,
+    @location(0) tex_coords: vec2<f32>,
+    // Updated!
+    @location(1) world_position: vec3<f32>,
+    @location(2) world_view_position: vec3<f32>,
+    @location(3) world_light_position: vec3<f32>,
+    @location(4) world_normal: vec3<f32>,
+    @location(5) world_tangent: vec3<f32>,
+    @location(6) world_bitangent: vec3<f32>,
+}
+
+@vertex
+fn vs_main(
+    model: VertexInput,
+    instance: InstanceInput,
+) -> VertexOutput {
+    let model_matrix = mat4x4<f32>(
+        instance.model_matrix_0,
+        instance.model_matrix_1,
+        instance.model_matrix_2,
+        instance.model_matrix_3,
+    );
+    let normal_matrix = mat3x3<f32>(
+        instance.normal_matrix_0,
+        instance.normal_matrix_1,
+        instance.normal_matrix_2,
+    );
+
+    // UPDATED!
+    let world_position = model_matrix * vec4<f32>(model.position, 1.0);
+
+    var out: VertexOutput;
+    out.clip_position = camera.view_proj * world_position;
+    out.tex_coords = model.tex_coords;
+    out.world_normal = normalize(normal_matrix * model.normal);
+    out.world_tangent = normalize(normal_matrix * model.tangent);
+    out.world_bitangent = normalize(normal_matrix * model.bitangent);
+    out.world_position = world_position.xyz;
+    out.world_view_position = camera.view_pos.xyz;
+    return out;
+}
+
+// Fragment shader
+
+@group(0) @binding(0)
+var t_diffuse: texture_2d<f32>;
+@group(0)@binding(1)
+var s_diffuse: sampler;
+@group(0)@binding(2)
+var t_normal: texture_2d<f32>;
+@group(0) @binding(3)
+var s_normal: sampler;
+
+@group(3)
+@binding(0)
+var env_map: texture_cube<f32>;
+@group(3)
+@binding(1)
+var env_sampler: sampler;
+
+@fragment
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
+    let object_color: vec4<f32> = textureSample(t_diffuse, s_diffuse, in.tex_coords);
+    let object_normal: vec4<f32> = textureSample(t_normal, s_normal, in.tex_coords);
+
+    // NEW!
+    // Adjust the tangent and bitangent using the Gramm-Schmidt process
+    // This makes sure that they are perpedicular to each other and the
+    // normal of the surface.
+    let world_tangent = normalize(in.world_tangent - dot(in.world_tangent, in.world_normal) * in.world_normal);
+    let world_bitangent = cross(world_tangent, in.world_normal);
+
+    // Convert the normal sample to world space
+    let TBN = mat3x3(
+        world_tangent,
+        world_bitangent,
+        in.world_normal,
+    );
+    let tangent_normal = object_normal.xyz * 2.0 - 1.0;
+    let world_normal = TBN * tangent_normal;
+
+    // Create the lighting vectors
+    let light_dir = normalize(light.position - in.world_position);
+    let view_dir = normalize(in.world_view_position - in.world_position);
+    let half_dir = normalize(view_dir + light_dir);
+
+    let diffuse_strength = max(dot(world_normal, light_dir), 0.0);
+    let diffuse_color = light.color * diffuse_strength;
+
+    let specular_strength = pow(max(dot(world_normal, half_dir), 0.0), 32.0);
+    let specular_color = specular_strength * light.color;
+
+    // NEW!
+    // Calculate reflections
+    let world_reflect = reflect(-view_dir, world_normal);
+    let reflection = textureSample(env_map, env_sampler, world_reflect).rgb;
+    let shininess = 0.1;
+
+    let result = (diffuse_color + specular_color) * object_color.xyz + reflection * shininess;
+
+    return vec4<f32>(result, object_color.a);
+}
+```
+
+A little note on the reflection math. The `view_dir`
+gives us the direction to the camera from the surface.
+The reflection math needs the direction from the
+camera to the surface so we negate `view_dir`. We
+then use `wgsl`'s builtin `reflect` function to
+reflect the inverted `view_dir` about the `world_normal`.
+This gives us a direction that we can use sample the
+environment map to get the color of the sky in that
+direction. Just looking at the reflection component
+gives us the following:
+
+![just-reflections](./just-reflections.png)
+
+Here's the finished scene:
+
+![with-reflections](./with-reflections.png)
+
+## Demo
+
+<div class="warn">
+
+If your browser doesn't support WebGPU, this example
+won't work for you.
+
+</div>
+
+<WasmExample example="tutorial13_hdr"></WasmExample>
+
+<AutoGithubLink/>

wasm-targets.json

@@ -47,6 +47,10 @@
         "package": "tutorial12-camera",
         "out": "docs/.vuepress/components/wasm/tutorial12_camera"
     },
+    {
+        "package": "tutorial13-hdr",
+        "out": "docs/.vuepress/components/wasm/tutorial13_hdr"
+    },
     {
         "package": "pong",
         "out": "docs/.vuepress/components/wasm/pong"