Stabilize smoke workload targets and runtime guards

AnastaZIuk · AnastaZIuk · commit f8d8679d1862 · 2026-02-17T14:16:36.000+01:00
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
@@ -26,6 +26,11 @@ set(WEBVULKAN_SPIRV_WASM_PACKAGE "lights0123/llvm-spir" CACHE STRING "Wasmer pac
 set(WEBVULKAN_SPIRV_WASM_ENTRYPOINT "clspv" CACHE STRING "Wasmer command used for SPIR-V probe in clang wasm smoke")
 set(WEBVULKAN_RUNTIME_BENCH_ITERATIONS "5" CACHE STRING "Timed dispatch iterations per lavapipe runtime mode smoke")
 set(WEBVULKAN_RUNTIME_WARMUP_ITERATIONS "1" CACHE STRING "Warmup dispatch iterations per lavapipe runtime mode smoke")
+option(
+  WEBVULKAN_ENABLE_EXPERIMENTAL_ATOMIC_WORKLOAD_SMOKE
+  "Enable experimental workload smoke targets that currently require unsupported LLVM interpreter intrinsics"
+  OFF
+)
 
 if(NOT COMMAND CPMAddPackage)
   if(NOT DEFINED WEBVULKAN_CPM_VERSION OR WEBVULKAN_CPM_VERSION STREQUAL "")
@@ -383,54 +388,60 @@ add_dependencies(lavapipe_runtime_smoke_hot_loop
   lavapipe_runtime_smoke_raw_llvm_ir_hot_loop
 )
 
-webvulkan_add_lavapipe_runtime_mode_smoke_target(
-  lavapipe_runtime_smoke_fast_wasm_atomic_single_counter
-  fast_wasm
-  hot_loop_single_dispatch
-  atomic_single_counter
-)
-webvulkan_add_lavapipe_runtime_mode_smoke_target(
-  lavapipe_runtime_smoke_raw_llvm_ir_atomic_single_counter
-  raw_llvm_ir
-  hot_loop_single_dispatch
-  atomic_single_counter
+add_custom_target(lavapipe_runtime_smoke_shader_workloads)
+add_dependencies(lavapipe_runtime_smoke_shader_workloads
+  lavapipe_runtime_smoke_fast_wasm_micro
+  lavapipe_runtime_smoke_raw_llvm_ir_micro
 )
 
-webvulkan_add_lavapipe_runtime_mode_smoke_target(
-  lavapipe_runtime_smoke_fast_wasm_atomic_per_workgroup
-  fast_wasm
-  hot_loop_single_dispatch
-  atomic_per_workgroup
-)
-webvulkan_add_lavapipe_runtime_mode_smoke_target(
-  lavapipe_runtime_smoke_raw_llvm_ir_atomic_per_workgroup
-  raw_llvm_ir
-  hot_loop_single_dispatch
-  atomic_per_workgroup
-)
+if(WEBVULKAN_ENABLE_EXPERIMENTAL_ATOMIC_WORKLOAD_SMOKE)
+  webvulkan_add_lavapipe_runtime_mode_smoke_target(
+    lavapipe_runtime_smoke_fast_wasm_no_race_unique_writes
+    fast_wasm
+    hot_loop_single_dispatch
+    no_race_unique_writes
+  )
+  webvulkan_add_lavapipe_runtime_mode_smoke_target(
+    lavapipe_runtime_smoke_raw_llvm_ir_no_race_unique_writes
+    raw_llvm_ir
+    hot_loop_single_dispatch
+    no_race_unique_writes
+  )
+  webvulkan_add_lavapipe_runtime_mode_smoke_target(
+    lavapipe_runtime_smoke_fast_wasm_atomic_single_counter
+    fast_wasm
+    hot_loop_single_dispatch
+    atomic_single_counter
+  )
+  webvulkan_add_lavapipe_runtime_mode_smoke_target(
+    lavapipe_runtime_smoke_raw_llvm_ir_atomic_single_counter
+    raw_llvm_ir
+    hot_loop_single_dispatch
+    atomic_single_counter
+  )
 
-webvulkan_add_lavapipe_runtime_mode_smoke_target(
-  lavapipe_runtime_smoke_fast_wasm_no_race_unique_writes
-  fast_wasm
-  hot_loop_single_dispatch
-  no_race_unique_writes
-)
-webvulkan_add_lavapipe_runtime_mode_smoke_target(
-  lavapipe_runtime_smoke_raw_llvm_ir_no_race_unique_writes
-  raw_llvm_ir
-  hot_loop_single_dispatch
-  no_race_unique_writes
-)
+  webvulkan_add_lavapipe_runtime_mode_smoke_target(
+    lavapipe_runtime_smoke_fast_wasm_atomic_per_workgroup
+    fast_wasm
+    hot_loop_single_dispatch
+    atomic_per_workgroup
+  )
+  webvulkan_add_lavapipe_runtime_mode_smoke_target(
+    lavapipe_runtime_smoke_raw_llvm_ir_atomic_per_workgroup
+    raw_llvm_ir
+    hot_loop_single_dispatch
+    atomic_per_workgroup
+  )
 
-add_custom_target(lavapipe_runtime_smoke_shader_workloads)
-add_dependencies(lavapipe_runtime_smoke_shader_workloads
-  lavapipe_runtime_smoke_fast_wasm_atomic_single_counter
-  lavapipe_runtime_smoke_raw_llvm_ir_atomic_single_counter
-  lavapipe_runtime_smoke_fast_wasm_atomic_per_workgroup
-  lavapipe_runtime_smoke_raw_llvm_ir_atomic_per_workgroup
-  lavapipe_runtime_smoke_fast_wasm_no_race_unique_writes
-  lavapipe_runtime_smoke_raw_llvm_ir_no_race_unique_writes
-)
+  add_dependencies(lavapipe_runtime_smoke_shader_workloads
+    lavapipe_runtime_smoke_fast_wasm_no_race_unique_writes
+    lavapipe_runtime_smoke_raw_llvm_ir_no_race_unique_writes
+    lavapipe_runtime_smoke_fast_wasm_atomic_single_counter
+    lavapipe_runtime_smoke_raw_llvm_ir_atomic_single_counter
+    lavapipe_runtime_smoke_fast_wasm_atomic_per_workgroup
+    lavapipe_runtime_smoke_raw_llvm_ir_atomic_per_workgroup
+  )
+endif()
 
 add_custom_target(lavapipe_runtime_smoke)
 add_dependencies(lavapipe_runtime_smoke lavapipe_runtime_smoke_fast_wasm lavapipe_runtime_smoke_raw_llvm_ir)
diff --git a/tests/wasm/tools/smoke_runtime.mjs b/tests/wasm/tools/smoke_runtime.mjs
@@ -47,12 +47,82 @@ function runtimeShaderThreadgroupSizeX(workloadName) {
   return workloadName === "write_const" ? 1 : 64;
 }
 
-async function compileRuntimeSpirv(storeValue, workloadName) {
-  const storeConst = `0x${(storeValue >>> 0).toString(16)}u`;
-  const threadgroupSizeX = runtimeShaderThreadgroupSizeX(workloadName);
-  const dxcWasmJs = process.env.WEBVULKAN_DXC_WASM_JS || "";
-  if (dxcWasmJs) {
-    const hlslSource = `
+function runtimeBenchProfileDescriptor(profileName) {
+  switch (profileName) {
+    case "micro":
+      return { dispatchesPerSubmit: 1024, dispatchX: 1, dispatchY: 1, dispatchZ: 1 };
+    case "realistic":
+      return { dispatchesPerSubmit: 64, dispatchX: 4, dispatchY: 1, dispatchZ: 1 };
+    case "hot_loop_single_dispatch":
+      return { dispatchesPerSubmit: 1, dispatchX: 256, dispatchY: 1, dispatchZ: 1 };
+    default:
+      throw new Error(`Unsupported runtime bench profile '${profileName}'`);
+  }
+}
+
+function runtimeShaderEntrypoint(workloadName) {
+  switch (workloadName) {
+    case "atomic_single_counter":
+      return "atomic_single_counter";
+    case "atomic_per_workgroup":
+      return "atomic_per_workgroup";
+    case "no_race_unique_writes":
+      return "no_race_unique_writes";
+    case "write_const":
+      return "write_const";
+    default:
+      throw new Error(`Unsupported runtime shader workload '${workloadName}'`);
+  }
+}
+
+function runtimeShaderHlslSource(
+  storeConst,
+  threadgroupSizeX,
+  workloadName,
+  dispatchInvocationsPerSubmit,
+  dispatchWorkgroupsPerSubmit
+) {
+  const entrypoint = runtimeShaderEntrypoint(workloadName);
+  if (workloadName === "atomic_single_counter") {
+    return `
+RWStructuredBuffer<uint> OutBuf : register(u0);
+
+[numthreads(${threadgroupSizeX}, 1, 1)]
+void ${entrypoint}(uint3 tid : SV_DispatchThreadID) {
+  InterlockedAdd(OutBuf[0], 1u);
+}
+`;
+  }
+
+  if (workloadName === "atomic_per_workgroup") {
+    return `
+RWStructuredBuffer<uint> OutBuf : register(u0);
+
+[numthreads(${threadgroupSizeX}, 1, 1)]
+void ${entrypoint}(uint3 groupThreadId : SV_GroupThreadID) {
+  if (groupThreadId.x == 0u) {
+    InterlockedAdd(OutBuf[0], 1u);
+  }
+}
+`;
+  }
+
+  if (workloadName === "no_race_unique_writes") {
+    return `
+RWStructuredBuffer<uint> OutBuf : register(u0);
+
+[numthreads(${threadgroupSizeX}, 1, 1)]
+void ${entrypoint}(uint3 tid : SV_DispatchThreadID) {
+  if (tid.x == 0u && tid.y == 0u && tid.z == 0u) {
+    OutBuf[0] = ${dispatchInvocationsPerSubmit}u;
+  }
+  uint idx = tid.x;
+  OutBuf[1u + idx] = idx + 1u;
+}
+`;
+  }
+
+  return `
 RWStructuredBuffer<uint> OutBuf : register(u0);
 
 uint webvulkan_mix_const(uint v, uint salt) {
@@ -72,14 +142,36 @@ uint webvulkan_compile_time_chain() {
 }
 
 [numthreads(${threadgroupSizeX}, 1, 1)]
-void write_const(uint3 tid : SV_DispatchThreadID) {
+void ${entrypoint}(uint3 tid : SV_DispatchThreadID) {
   uint folded = webvulkan_compile_time_chain();
   if (folded == 0xdeadbeefu) {
     OutBuf[1] = folded;
   }
   OutBuf[0] = ${storeConst};
 }
 `;
+}
+
+async function compileRuntimeSpirv(storeValue, workloadName) {
+  const storeConst = `0x${(storeValue >>> 0).toString(16)}u`;
+  const threadgroupSizeX = runtimeShaderThreadgroupSizeX(workloadName);
+  const shaderEntrypoint = runtimeShaderEntrypoint(workloadName);
+  const profile = runtimeBenchProfileDescriptor(runtimeBenchProfile);
+  const dispatchWorkgroupsPerSubmit =
+    profile.dispatchesPerSubmit *
+    profile.dispatchX *
+    profile.dispatchY *
+    profile.dispatchZ;
+  const dispatchInvocationsPerSubmit = dispatchWorkgroupsPerSubmit * threadgroupSizeX;
+  const dxcWasmJs = process.env.WEBVULKAN_DXC_WASM_JS || "";
+  if (dxcWasmJs) {
+    const hlslSource = runtimeShaderHlslSource(
+      storeConst,
+      threadgroupSizeX,
+      workloadName,
+      dispatchInvocationsPerSubmit,
+      dispatchWorkgroupsPerSubmit
+    );
 
     const scratchDir = await mkdtemp(join(tmpdir(), "webvulkan-dxc-wasm-"));
     const inputFile = "runtime_smoke.hlsl";
@@ -94,7 +186,7 @@ void write_const(uint3 tid : SV_DispatchThreadID) {
       "-T",
       "cs_6_0",
       "-E",
-      "write_const",
+      shaderEntrypoint,
       "-Fo",
       outputFile,
       inputFile
@@ -117,10 +209,14 @@ void write_const(uint3 tid : SV_DispatchThreadID) {
     return {
       provider: `dxc-wasm:${dxcWasmJs}`,
       bytes,
-      entrypoint: "write_const"
+      entrypoint: shaderEntrypoint
     };
   }
 
+  if (workloadName !== "write_const") {
+    throw new Error(`Workload '${workloadName}' requires WEBVULKAN_DXC_WASM_JS`);
+  }
+
   const wasmerBin = process.env.WEBVULKAN_WASMER_BIN;
   if (!wasmerBin) {
     throw new Error("WEBVULKAN_WASMER_BIN is required when SMOKE_REQUIRE_RUNTIME_SPIRV=1");
@@ -146,7 +242,7 @@ uint webvulkan_compile_time_chain() {
 }
 
 __attribute__((reqd_work_group_size(${threadgroupSizeX}, 1, 1)))
-__kernel void write_const(__global uint* out) {
+__kernel void ${shaderEntrypoint}(__global uint* out) {
   uint folded = webvulkan_compile_time_chain();
   if (folded == (uint)0xdeadbeefu) {
     out[1] = folded;
@@ -200,7 +296,7 @@ __kernel void write_const(__global uint* out) {
       return {
         provider: attempt.provider,
         bytes: result.stdout,
-        entrypoint: "write_const"
+        entrypoint: shaderEntrypoint
       };
     }
 
@@ -519,7 +615,14 @@ async function runFastWasmSmoke(shaderValue) {
   console.log(`  runtime_wasm.bytes=${runtimeWasm.bytes.length}`);
 
   console.log("runtime smoke discover_key");
-  invokeSmokeOnce();
+  if (runtimeShaderWorkload === "write_const") {
+    invokeSmokeOnce();
+  } else {
+    const discoverRc = smokeFn();
+    if (discoverRc !== 0) {
+      console.log(`runtime smoke discover_key observed_nonzero_rc=${discoverRc} before runtime module registration`);
+    }
+  }
 
   const hasCapturedKey = runtime.ccall("webvulkan_runtime_has_captured_shader_key", "number", [], []) !== 0;
   if (!hasCapturedKey) {
@@ -576,6 +679,9 @@ async function runFastWasmSmoke(shaderValue) {
 }
 
 async function runRawLlvmIrSmoke(shaderValue) {
+  if (runtimeShaderWorkload !== "write_const") {
+    throw new Error(`raw_llvm_ir mode currently supports only write_const workload, got '${runtimeShaderWorkload}'`);
+  }
   const spirv = await compileRuntimeSpirv(shaderValue, runtimeShaderWorkload);
   setRuntimeBenchProfile(runtimeBenchProfileValue);
   setRuntimeShaderWorkload(runtimeShaderWorkloadValue);
