Fabio
c36510c21e
Some checks failed
Linux Build / Linux cpu (push) Has been cancelled
Linux Build / Linux cpu-clang (push) Has been cancelled
Linux Build / Linux vulkan (push) Has been cancelled
mac-build / macOS CPU (push) Has been cancelled
Nix Build / Nix (cpu) (push) Has been cancelled
Nix Build / Nix (vulkan) (push) Has been cancelled
windows-build / Windows CPU (push) Has been cancelled
569 lines
33 KiB
C++
569 lines
33 KiB
C++
#include "engine/framework/core/backend.h"
|
|
#include "engine/framework/codecs/wan_video_vae_runtime.h"
|
|
#include "engine/framework/debug/profiler.h"
|
|
#include "engine/framework/modules/conv_modules.h"
|
|
|
|
#include <ggml-alloc.h>
|
|
#include <algorithm>
|
|
#include <limits>
|
|
#include <ggml-backend.h>
|
|
#include <ggml.h>
|
|
|
|
#include <cmath>
|
|
#include <cstring>
|
|
#include <iostream>
|
|
#include <memory>
|
|
#include <optional>
|
|
#include <stdexcept>
|
|
#include <string>
|
|
#include <utility>
|
|
#include <vector>
|
|
|
|
namespace {
|
|
void check(bool condition, const char * message) {
|
|
if (!condition) throw std::runtime_error(message);
|
|
}
|
|
|
|
void run_case(ggml_backend_t backend, engine::modules::Conv3dConfig config,
|
|
int batch, int depth, int height, int width, bool framework, bool f32_input,
|
|
bool expect_unavailable, bool check_cast = false, bool packed_weights = false,
|
|
bool sample_scalar_reference = false) {
|
|
namespace core = engine::core;
|
|
std::unique_ptr<ggml_context, decltype(&ggml_free)> ctx(
|
|
ggml_init({4 * 1024 * 1024, nullptr, true}), ggml_free);
|
|
check(bool(ctx), "context allocation failed");
|
|
core::ModuleBuildContext build{ctx.get(), "test.cudnn_conv3d", core::BackendType::Cuda};
|
|
auto x = core::make_tensor(build, f32_input ? GGML_TYPE_F32 : GGML_TYPE_F16,
|
|
core::TensorShape::from_dims({batch * config.in_channels, depth, height, width}));
|
|
auto w = core::make_tensor(build, GGML_TYPE_F16, core::TensorShape::from_dims(
|
|
{config.out_channels * config.in_channels, config.kernel_depth, config.kernel_height, config.kernel_width}));
|
|
auto bias = core::make_tensor(build, GGML_TYPE_F32, core::TensorShape::from_dims({config.out_channels}));
|
|
auto reference_w = packed_weights ? core::make_tensor(build, GGML_TYPE_F16, w.shape) : w;
|
|
if (packed_weights) ggml_set_input(reference_w.tensor);
|
|
ggml_set_input(x.tensor);
|
|
ggml_set_input(w.tensor);
|
|
ggml_set_input(bias.tensor);
|
|
const ggml_type output_type = config.cuda_f16_output ? GGML_TYPE_F16 : GGML_TYPE_F32;
|
|
ggml_tensor * output;
|
|
if (framework) {
|
|
config.cuda_cudnn = true;
|
|
config.cuda_prepacked_weights = packed_weights;
|
|
output = engine::modules::Conv3dModule(config).build(build, x,
|
|
{w, config.use_bias ? std::optional<core::TensorValue>(bias) : std::nullopt}).tensor;
|
|
} else {
|
|
output = (config.cuda_f16_output ? ggml_conv_3d_direct_f16 : ggml_conv_3d_direct)(ctx.get(), w.tensor, x.tensor,
|
|
config.stride_width, config.stride_height, config.stride_depth,
|
|
config.padding_width, config.padding_height, config.padding_depth,
|
|
config.dilation_width, config.dilation_height, config.dilation_depth,
|
|
int(config.in_channels), batch, int(config.out_channels));
|
|
check(output->type == output_type && output->op_params[12] == int(config.cuda_f16_output),
|
|
"direct Conv3D output type/default marker changed");
|
|
const auto marker = output->op_params[12];
|
|
output->op_params[12] = 0;
|
|
check(!ggml_backend_supports_op(backend, output), "unmarked direct Conv3D advertised on CUDA");
|
|
output->op_params[12] = marker;
|
|
ggml_conv_3d_direct_set_cudnn(output, config.use_bias ? bias.tensor : nullptr, config.cuda_bf16_round, packed_weights, config.cuda_replicate_depth_left);
|
|
}
|
|
check(output->type == output_type && ggml_is_contiguous(output) &&
|
|
output->nb[0] == ggml_type_size(output_type) &&
|
|
ggml_nbytes(output) == size_t(ggml_nelements(output)) * ggml_type_size(output_type),
|
|
"cuDNN Conv3D output dtype/byte size mismatch");
|
|
check(output->op == GGML_OP_CONV_3D, "native output epilogue added a separate graph operation");
|
|
// Independent GPU oracle: identical convolution with standalone F32 bias
|
|
// addition and ROUND_BF16. Keep its output live alongside the fused result.
|
|
ggml_tensor * reference = nullptr;
|
|
if (config.use_bias || config.cuda_bf16_round || config.cuda_f16_output || packed_weights) {
|
|
reference = ggml_conv_3d_direct(ctx.get(), reference_w.tensor, x.tensor,
|
|
config.stride_width, config.stride_height, config.stride_depth,
|
|
config.padding_width, config.padding_height, config.padding_depth,
|
|
config.dilation_width, config.dilation_height, config.dilation_depth,
|
|
int(config.in_channels), batch, int(config.out_channels));
|
|
ggml_conv_3d_direct_set_cudnn(reference, nullptr, false, false, config.cuda_replicate_depth_left);
|
|
if (config.use_bias) {
|
|
reference = ggml_add(ctx.get(), reference,
|
|
ggml_reshape_4d(ctx.get(), bias.tensor, 1, 1, 1, config.out_channels));
|
|
}
|
|
if (config.cuda_bf16_round) reference = ggml_round_bf16(ctx.get(), reference);
|
|
if (config.cuda_f16_output) reference = ggml_cast(ctx.get(), reference, GGML_TYPE_F16);
|
|
ggml_set_output(reference);
|
|
}
|
|
ggml_set_output(output);
|
|
auto * graph = ggml_new_graph_custom(ctx.get(), 128, false);
|
|
ggml_build_forward_expand(graph, output);
|
|
if (reference) ggml_build_forward_expand(graph, reference);
|
|
ggml_tensor * conv = nullptr;
|
|
for (int i = 0; i < ggml_graph_n_nodes(graph); ++i) {
|
|
auto * node = ggml_graph_node(graph, i);
|
|
if (node->op == GGML_OP_CONV_3D && node == output) conv = node;
|
|
}
|
|
check(conv && conv->op_params[12] == 1, "framework did not select marked direct Conv3D");
|
|
check(conv->src[1] == x.tensor, "marked Conv3D introduced a separate activation cast/copy");
|
|
check(conv->src[2] == (config.use_bias ? bias.tensor : nullptr), "incorrect fused bias dependency");
|
|
check(conv->op_params[13] == int(config.cuda_bf16_round), "incorrect fused BF16 round flag");
|
|
check(conv->op_params[14] == int(packed_weights), "incorrect explicit packed-weight flag");
|
|
conv->op_params[14] = 2;
|
|
check(!ggml_backend_supports_op(backend, conv), "unknown packed-weight flag accepted");
|
|
conv->op_params[14] = int(packed_weights);
|
|
check(conv->op_params[15] == config.cuda_replicate_depth_left + 1, "incorrect replicated depth padding");
|
|
conv->op_params[15] = 2 * config.padding_depth + 2;
|
|
check(!ggml_backend_supports_op(backend, conv), "excess replicated depth padding accepted");
|
|
conv->op_params[15] = config.cuda_replicate_depth_left + 1;
|
|
check(ggml_backend_supports_op(backend, conv) != expect_unavailable,
|
|
"cuDNN Conv3D support does not match the requested build configuration");
|
|
// Reject CPU/Metal scheduling without initializing or executing either backend.
|
|
for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
|
|
const auto dev = ggml_backend_dev_get(i);
|
|
const std::string name = ggml_backend_reg_name(ggml_backend_dev_backend_reg(dev));
|
|
if (name == "CPU" || name == "Metal") {
|
|
check(!ggml_backend_dev_supports_op(dev, conv), "CUDA marker accepted by CPU/Metal");
|
|
}
|
|
}
|
|
conv->op_params[12] = 2;
|
|
check(!ggml_backend_supports_op(backend, conv), "unknown lowering marker accepted");
|
|
conv->op_params[12] = 1;
|
|
conv->op_params[13] = 2;
|
|
check(!ggml_backend_supports_op(backend, conv), "unknown BF16 round flag accepted");
|
|
conv->op_params[13] = int(config.cuda_bf16_round);
|
|
if (config.cuda_f16_output) {
|
|
conv->op_params[13] = 1;
|
|
check(!ggml_backend_supports_op(backend, conv), "F16 output with BF16 rounding advertised");
|
|
conv->op_params[13] = 0;
|
|
}
|
|
auto * saved_bias = conv->src[2];
|
|
conv->src[2] = bias.tensor;
|
|
bias.tensor->type = GGML_TYPE_F16;
|
|
check(!ggml_backend_supports_op(backend, conv), "non-F32 bias advertised");
|
|
bias.tensor->type = GGML_TYPE_F32;
|
|
++bias.tensor->ne[0];
|
|
check(!ggml_backend_supports_op(backend, conv), "incorrect bias channel count advertised");
|
|
--bias.tensor->ne[0];
|
|
if (config.out_channels > 1) {
|
|
bias.tensor->nb[0] += sizeof(float);
|
|
check(!ggml_backend_supports_op(backend, conv), "noncontiguous bias advertised");
|
|
bias.tensor->nb[0] -= sizeof(float);
|
|
}
|
|
conv->src[2] = saved_bias;
|
|
conv->src[3] = bias.tensor;
|
|
check(!ggml_backend_supports_op(backend, conv), "unexpected extra source advertised");
|
|
conv->src[3] = nullptr;
|
|
const auto saved_type = conv->src[1]->type;
|
|
conv->src[1]->type = GGML_TYPE_BF16;
|
|
check(!ggml_backend_supports_op(backend, conv), "unsupported BF16 input advertised");
|
|
conv->src[1]->type = saved_type;
|
|
for (int axis = 0; axis < 4; ++axis) {
|
|
if (conv->src[0]->ne[axis] == 1) continue;
|
|
const auto saved_stride = conv->src[0]->nb[axis];
|
|
conv->src[0]->nb[axis] += sizeof(ggml_fp16_t);
|
|
check(!ggml_backend_supports_op(backend, conv), "noncontiguous weights advertised");
|
|
conv->src[0]->nb[axis] = saved_stride;
|
|
break;
|
|
}
|
|
if (expect_unavailable) return;
|
|
core::validate_backend_graph_supported(backend, graph, "test.cudnn_conv3d");
|
|
std::unique_ptr<ggml_gallocr, decltype(&ggml_gallocr_free)> arena(
|
|
ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend)), ggml_gallocr_free);
|
|
check(bool(arena) && ggml_gallocr_alloc_graph(arena.get(), graph), "graph allocation failed");
|
|
// Release captured backend resources before arena/context destruction, also on failure.
|
|
const auto release = [backend](ggml_cgraph * g) { core::release_backend_graph_resources(backend, g, true); };
|
|
std::unique_ptr<ggml_cgraph, decltype(release)> graph_resources(graph, release);
|
|
|
|
std::vector<float> xf(size_t(ggml_nelements(x.tensor))), wf(size_t(ggml_nelements(w.tensor)));
|
|
std::vector<float> bf(size_t(config.out_channels));
|
|
std::vector<ggml_fp16_t> xh(xf.size()), wh(wf.size());
|
|
// Include non-F16 values and rounding ties. The scalar reference uses exactly
|
|
// the values of the old explicit F32 -> F16 activation cast.
|
|
const float ties[] = {1.00048828125f, 1.00146484375f, -1.00048828125f, -1.00146484375f};
|
|
for (size_t i = 0; i < xf.size(); ++i) {
|
|
xf[i] = float(int(i % 17) - 8) / 16 + float(int(i % 5) - 2) * 0.00013f;
|
|
if (i < 4) xf[i] = ties[i];
|
|
xh[i] = ggml_fp32_to_fp16(xf[i]);
|
|
}
|
|
for (size_t i = 0; i < wf.size(); ++i) {
|
|
wf[i] = float(int(i % 13) - 6) / 32;
|
|
wh[i] = ggml_fp32_to_fp16(wf[i]);
|
|
}
|
|
if (check_cast) {
|
|
check(config.in_channels == 1 && config.out_channels == 1 &&
|
|
config.kernel_depth == 1 && config.kernel_height == 1 && config.kernel_width == 1 &&
|
|
config.stride_depth == 1 && config.stride_height == 1 && config.stride_width == 1 &&
|
|
config.padding_depth == 0 && config.padding_height == 0 && config.padding_width == 0,
|
|
"cast-value check requires identity convolution");
|
|
wf[0] = 1.0f;
|
|
wh[0] = ggml_fp32_to_fp16(1.0f);
|
|
}
|
|
for (size_t i = 0; i < bf.size(); ++i) {
|
|
bf[i] = float(i + 1) / 128;
|
|
if (config.cuda_bf16_round || config.cuda_f16_output || check_cast) bf[i] += float(i + 1) * 0.00013f;
|
|
}
|
|
const int od = int(output->ne[2]), oh = int(output->ne[1]), ow = int(output->ne[0]);
|
|
check(od == (depth + 2 * config.padding_depth - config.dilation_depth * (config.kernel_depth - 1) - 1) / config.stride_depth + 1 &&
|
|
oh == (height + 2 * config.padding_height - config.dilation_height * (config.kernel_height - 1) - 1) / config.stride_height + 1 &&
|
|
ow == (width + 2 * config.padding_width - config.dilation_width * (config.kernel_width - 1) - 1) / config.stride_width + 1 &&
|
|
output->ne[3] == batch * config.out_channels, "incorrect Conv3D output shape");
|
|
std::vector<float> expected(size_t(ggml_nelements(output))), previous;
|
|
if (sample_scalar_reference)
|
|
std::fill(expected.begin(), expected.end(), std::numeric_limits<float>::quiet_NaN());
|
|
std::vector<ggml_fp16_t> previous_half;
|
|
for (int n = 0; n < batch; ++n)
|
|
for (int oc = 0; oc < config.out_channels; ++oc)
|
|
for (int z = 0; z < od; ++z)
|
|
for (int y = 0; y < oh; ++y)
|
|
for (int col = 0; col < ow; ++col) {
|
|
// Keep the large workspace-budget case cheap: cover both batches,
|
|
// channel ends, spatial borders, and the partial final depth tile.
|
|
if (sample_scalar_reference &&
|
|
((oc != 0 && oc != config.out_channels / 2 && oc != config.out_channels - 1) ||
|
|
(z != 0 && z != od / 2 && z != od - 1) ||
|
|
(y != 0 && y != oh / 2 && y != oh - 1) ||
|
|
(col != 0 && col != ow / 2 && col != ow - 1))) continue;
|
|
float sum = 0;
|
|
for (int ic = 0; ic < config.in_channels; ++ic)
|
|
for (int kz = 0; kz < config.kernel_depth; ++kz)
|
|
for (int ky = 0; ky < config.kernel_height; ++ky)
|
|
for (int kx = 0; kx < config.kernel_width; ++kx) {
|
|
const int depth_left = config.cuda_replicate_depth_left >= 0
|
|
? config.cuda_replicate_depth_left : config.padding_depth;
|
|
int iz = z * config.stride_depth - depth_left + kz * config.dilation_depth;
|
|
if (config.cuda_replicate_depth_left >= 0) iz = std::clamp(iz, 0, depth - 1);
|
|
const int iy = y * config.stride_height - config.padding_height + ky * config.dilation_height;
|
|
const int ix = col * config.stride_width - config.padding_width + kx * config.dilation_width;
|
|
if (iz < 0 || iz >= depth || iy < 0 || iy >= height || ix < 0 || ix >= width) continue;
|
|
const size_t xi = (((n * config.in_channels + ic) * depth + iz) * height + iy) * width + ix;
|
|
const size_t wi = (((oc * config.in_channels + ic) * config.kernel_depth + kz) * config.kernel_height + ky) * config.kernel_width + kx;
|
|
sum += ggml_fp16_to_fp32(xh[xi]) * wf[wi];
|
|
}
|
|
const size_t oi = (((n * config.out_channels + oc) * od + z) * oh + y) * ow + col;
|
|
expected[oi] = ggml_fp16_to_fp32(ggml_fp32_to_fp16(sum)) + (config.use_bias ? bf[oc] : 0);
|
|
if (config.cuda_bf16_round) {
|
|
expected[oi] = ggml_bf16_to_fp32(ggml_fp32_to_bf16(expected[oi]));
|
|
}
|
|
if (config.cuda_f16_output) {
|
|
expected[oi] = ggml_fp16_to_fp32(ggml_fp32_to_fp16(expected[oi]));
|
|
}
|
|
}
|
|
std::vector<ggml_fp16_t> packed(wh.size());
|
|
if (packed_weights) {
|
|
const size_t kernel = size_t(config.kernel_depth) * config.kernel_height * config.kernel_width;
|
|
for (size_t oc = 0; oc < size_t(config.out_channels); ++oc)
|
|
for (size_t k = 0; k < kernel; ++k)
|
|
for (size_t ic = 0; ic < size_t(config.in_channels); ++ic)
|
|
packed[(oc * kernel + k) * config.in_channels + ic] =
|
|
wh[(oc * config.in_channels + ic) * kernel + k];
|
|
}
|
|
for (int repeat = 0; repeat < 3; ++repeat) {
|
|
ggml_backend_tensor_set(x.tensor, f32_input ? static_cast<const void *>(xf.data()) : xh.data(), 0,
|
|
f32_input ? xf.size() * sizeof(float) : xh.size() * sizeof(ggml_fp16_t));
|
|
ggml_backend_tensor_set(w.tensor, packed_weights ? packed.data() : wh.data(), 0, wh.size() * sizeof(ggml_fp16_t));
|
|
if (packed_weights) ggml_backend_tensor_set(reference_w.tensor, wh.data(), 0, wh.size() * sizeof(ggml_fp16_t));
|
|
if (config.use_bias) ggml_backend_tensor_set(bias.tensor, bf.data(), 0, bf.size() * sizeof(float));
|
|
check(ggml_backend_graph_compute(backend, graph) == GGML_STATUS_SUCCESS, "CUDA execution failed");
|
|
ggml_backend_synchronize(backend);
|
|
std::vector<float> actual(expected.size());
|
|
std::vector<ggml_fp16_t> actual_half;
|
|
if (config.cuda_f16_output) {
|
|
actual_half.resize(actual.size());
|
|
ggml_backend_tensor_get(output, actual_half.data(), 0, actual_half.size() * sizeof(ggml_fp16_t));
|
|
for (size_t i = 0; i < actual.size(); ++i) actual[i] = ggml_fp16_to_fp32(actual_half[i]);
|
|
} else {
|
|
ggml_backend_tensor_get(output, actual.data(), 0, actual.size() * sizeof(float));
|
|
}
|
|
if (reference) {
|
|
if (config.cuda_f16_output) {
|
|
std::vector<ggml_fp16_t> separate(actual_half.size());
|
|
ggml_backend_tensor_get(reference, separate.data(), 0, separate.size() * sizeof(ggml_fp16_t));
|
|
check(std::memcmp(actual_half.data(), separate.data(), actual_half.size() * sizeof(ggml_fp16_t)) == 0,
|
|
"F16 output differs bitwise from F32 convolution + F32 bias -> F16 GPU oracle");
|
|
} else {
|
|
std::vector<float> separate(actual.size());
|
|
ggml_backend_tensor_get(reference, separate.data(), 0, separate.size() * sizeof(float));
|
|
check(std::memcmp(actual.data(), separate.data(), actual.size() * sizeof(float)) == 0,
|
|
"fused epilogue differs bitwise from standalone GPU bias/ROUND_BF16");
|
|
}
|
|
}
|
|
for (size_t i = 0; i < actual.size(); ++i) {
|
|
check(std::isfinite(actual[i]), "nonfinite Conv3D output");
|
|
if (sample_scalar_reference && !std::isfinite(expected[i])) continue;
|
|
if (check_cast) check(actual[i] == expected[i], "identity convolution changed cast/bias/round values");
|
|
const float relative_tolerance = config.cuda_bf16_round ? 0.01f : 0.001f;
|
|
check(std::isfinite(actual[i]) && std::abs(actual[i] - expected[i]) <= 0.002f + relative_tolerance * std::abs(expected[i]),
|
|
"cuDNN Conv3D differs from scalar cross-correlation");
|
|
}
|
|
if (repeat != 0) check(actual == previous, "repeated/captured Conv3D changed output");
|
|
if (repeat != 0 && config.cuda_f16_output) {
|
|
check(actual_half == previous_half, "repeated/captured F16 output changed bits");
|
|
}
|
|
previous_half = std::move(actual_half);
|
|
previous = std::move(actual);
|
|
}
|
|
}
|
|
|
|
// Storage oracle: the existing F32 fused kernel followed by an independent GPU
|
|
// cast. Exercise both reduction kernels and their channel/spatial tail handling.
|
|
void run_norm_case(ggml_backend_t backend, int channels, bool f32_input,
|
|
bool coalesced, bool conv_available) {
|
|
namespace core = engine::core;
|
|
namespace codecs = engine::codecs;
|
|
std::unique_ptr<ggml_context, decltype(&ggml_free)> ctx(
|
|
ggml_init({4 * 1024 * 1024, nullptr, true}), ggml_free);
|
|
check(bool(ctx), "norm context allocation failed");
|
|
core::ModuleBuildContext build{ctx.get(), "test.channel_rms_silu", core::BackendType::Cuda};
|
|
auto x = core::make_tensor(build, f32_input ? GGML_TYPE_F32 : GGML_TYPE_F16,
|
|
core::TensorShape::from_dims({channels, 2, 3, 17}));
|
|
auto gamma = core::make_tensor(build, GGML_TYPE_F32, core::TensorShape::from_dims({channels}));
|
|
ggml_set_input(x.tensor);
|
|
ggml_set_input(gamma.tensor);
|
|
codecs::WanVideoRMSNorm3dConfig config;
|
|
config.channels = channels;
|
|
config.eps = 1.e-8f;
|
|
config.cuda_fast_lowering = true;
|
|
config.fuse_silu = true;
|
|
config.cuda_f16_output = true;
|
|
auto fused = codecs::WanVideoRMSNorm3dModule(config).build(build, x, {gamma});
|
|
config.cuda_f16_output = false;
|
|
auto f32 = codecs::WanVideoRMSNorm3dModule(config).build(build, x, {gamma});
|
|
check(fused.type == GGML_TYPE_F16 && f32.type == GGML_TYPE_F32 &&
|
|
fused.tensor->op == GGML_OP_RMS_NORM_CHANNELS_SILU &&
|
|
fused.tensor->src[0] == x.tensor && f32.tensor->src[0] == x.tensor,
|
|
"fused module inserted an input cast or selected the wrong operation/type");
|
|
check(fused.tensor->op_params[1] == GGML_RMS_NORM_CHANNELS_LOWERING_CUDA_COALESCED &&
|
|
f32.tensor->op_params[1] == GGML_RMS_NORM_CHANNELS_LOWERING_CUDA_COALESCED,
|
|
"fused module did not select the established CUDA coalesced lowering");
|
|
if (!coalesced) {
|
|
// Factory default (zero) selects the original block reduction.
|
|
fused.tensor->op_params[1] = 0;
|
|
f32.tensor->op_params[1] = 0;
|
|
}
|
|
auto * oracle = ggml_cast(ctx.get(), f32.tensor, GGML_TYPE_F16);
|
|
const size_t count = size_t(ggml_nelements(x.tensor));
|
|
check(ggml_is_contiguous(fused.tensor) && ggml_nbytes(fused.tensor) == count * sizeof(ggml_fp16_t),
|
|
"F16 norm output byte layout mismatch");
|
|
check(ggml_backend_supports_op(backend, fused.tensor), "CUDA rejected F16 RMSNorm/SiLU");
|
|
for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
|
|
auto dev = ggml_backend_dev_get(i);
|
|
const std::string name = ggml_backend_reg_name(ggml_backend_dev_backend_reg(dev));
|
|
if (name == "CPU" || name == "Metal" || name == "Vulkan") {
|
|
check(!ggml_backend_dev_supports_op(dev, fused.tensor), "unsupported backend advertised F16 RMSNorm/SiLU");
|
|
}
|
|
}
|
|
// Explicit F16 requests must fail rather than becoming an implicit cast path.
|
|
for (int invalid = 0; invalid < 3; ++invalid) {
|
|
auto bad = config;
|
|
bad.cuda_f16_output = true;
|
|
if (invalid == 0) bad.cuda_fast_lowering = false;
|
|
if (invalid == 1) bad.fuse_silu = false;
|
|
auto non_cuda = build;
|
|
non_cuda.backend_type = core::BackendType::Cpu;
|
|
bool rejected = false;
|
|
try {
|
|
codecs::WanVideoRMSNorm3dModule module(bad);
|
|
if (invalid == 2) module.build(non_cuda, x, {gamma});
|
|
} catch (const std::runtime_error &) {
|
|
rejected = true;
|
|
}
|
|
check(rejected, "invalid F16 norm configuration silently accepted");
|
|
}
|
|
auto non_cuda = build;
|
|
non_cuda.backend_type = core::BackendType::Cpu;
|
|
auto fallback = codecs::WanVideoRMSNorm3dModule(config).build(non_cuda, x, {gamma});
|
|
check(fallback.type == GGML_TYPE_F32 && fallback.tensor->op == GGML_OP_UNARY &&
|
|
ggml_get_unary_op(fallback.tensor) == GGML_UNARY_OP_SILU,
|
|
"generic fused-SiLU request did not use the existing module composition");
|
|
|
|
auto * graph = ggml_new_graph_custom(ctx.get(), 128, false);
|
|
ggml_set_output(fused.tensor);
|
|
ggml_set_output(oracle);
|
|
ggml_build_forward_expand(graph, fused.tensor);
|
|
ggml_build_forward_expand(graph, oracle);
|
|
ggml_tensor * conv = nullptr;
|
|
ggml_tensor * conv_oracle = nullptr;
|
|
ggml_tensor * weights = nullptr;
|
|
if (conv_available) {
|
|
weights = ggml_new_tensor_4d(ctx.get(), GGML_TYPE_F16, 1, 1, 1, channels * 3);
|
|
ggml_set_input(weights);
|
|
conv = ggml_conv_3d_direct_f16(ctx.get(), weights, fused.tensor,
|
|
1, 1, 1, 0, 0, 0, 1, 1, 1, channels, 1, 3);
|
|
conv_oracle = ggml_conv_3d_direct_f16(ctx.get(), weights, oracle,
|
|
1, 1, 1, 0, 0, 0, 1, 1, 1, channels, 1, 3);
|
|
ggml_conv_3d_direct_set_cudnn(conv);
|
|
ggml_conv_3d_direct_set_cudnn(conv_oracle);
|
|
check(conv->src[1] == fused.tensor, "norm-to-conv inserted a conversion");
|
|
ggml_set_output(conv);
|
|
ggml_set_output(conv_oracle);
|
|
ggml_build_forward_expand(graph, conv);
|
|
ggml_build_forward_expand(graph, conv_oracle);
|
|
}
|
|
core::validate_backend_graph_supported(backend, graph, "test.channel_rms_silu");
|
|
std::unique_ptr<ggml_gallocr, decltype(&ggml_gallocr_free)> arena(
|
|
ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend)), ggml_gallocr_free);
|
|
check(bool(arena) && ggml_gallocr_alloc_graph(arena.get(), graph), "norm graph allocation failed");
|
|
const auto release = [backend](ggml_cgraph * g) { core::release_backend_graph_resources(backend, g, true); };
|
|
std::unique_ptr<ggml_cgraph, decltype(release)> graph_resources(graph, release);
|
|
std::vector<float> xf(count), gf(channels);
|
|
std::vector<ggml_fp16_t> xh(count), actual(count), expected(count), previous;
|
|
std::vector<ggml_fp16_t> wh(size_t(channels) * 3, ggml_fp32_to_fp16(0.f));
|
|
for (int oc = 0; oc < 3; ++oc) wh[size_t(oc) * channels + oc % channels] = ggml_fp32_to_fp16(1.f);
|
|
if (weights) ggml_backend_tensor_set(weights, wh.data(), 0, wh.size() * sizeof(ggml_fp16_t));
|
|
for (int repeat = 0; repeat < 6; ++repeat) {
|
|
// Change data after capture, then verify repeated replay is stable too.
|
|
const int revision = repeat / 3;
|
|
for (size_t i = 0; i < count; ++i) {
|
|
xf[i] = float(int((i * 7 + revision) % 41) - 20) * 0.07137f;
|
|
if (i % 102 == 0) xf[i] = 0.f;
|
|
xh[i] = ggml_fp32_to_fp16(xf[i]);
|
|
}
|
|
for (int c = 0; c < channels; ++c) gf[c] = float(c % 11 - 5) * 0.1937f + revision * 0.017f;
|
|
ggml_backend_tensor_set(x.tensor, f32_input ? static_cast<const void *>(xf.data()) : xh.data(),
|
|
0, count * (f32_input ? sizeof(float) : sizeof(ggml_fp16_t)));
|
|
ggml_backend_tensor_set(gamma.tensor, gf.data(), 0, gf.size() * sizeof(float));
|
|
check(ggml_backend_graph_compute(backend, graph) == GGML_STATUS_SUCCESS, "norm CUDA execution failed");
|
|
ggml_backend_synchronize(backend);
|
|
ggml_backend_tensor_get(fused.tensor, actual.data(), 0, actual.size() * sizeof(ggml_fp16_t));
|
|
ggml_backend_tensor_get(oracle, expected.data(), 0, expected.size() * sizeof(ggml_fp16_t));
|
|
check(actual == expected, "F16 RMSNorm/SiLU differs bitwise from F32 fused kernel + GPU cast");
|
|
for (auto value : actual) check(std::isfinite(ggml_fp16_to_fp32(value)), "nonfinite norm result");
|
|
if (repeat % 3 != 0) check(actual == previous, "repeated/captured F16 norm changed bits");
|
|
if (repeat == 3) check(actual != previous, "captured norm ignored updated input/gamma");
|
|
previous = actual;
|
|
if (conv) {
|
|
std::vector<ggml_fp16_t> got(size_t(ggml_nelements(conv))), ref(got.size());
|
|
ggml_backend_tensor_get(conv, got.data(), 0, got.size() * sizeof(ggml_fp16_t));
|
|
ggml_backend_tensor_get(conv_oracle, ref.data(), 0, ref.size() * sizeof(ggml_fp16_t));
|
|
check(got == ref, "F16 norm-to-conv differs from independent cast-to-conv oracle");
|
|
}
|
|
}
|
|
}
|
|
} // namespace
|
|
|
|
int main(int argc, char ** argv) {
|
|
if ((argc != 3 && argc != 4) || std::string(argv[1]) != "--log" ||
|
|
(argc == 4 && std::string(argv[3]) != "--expect-unavailable")) {
|
|
std::cerr << "Usage: cuda_cudnn_conv3d_test --log PATH [--expect-unavailable]\n";
|
|
return 1;
|
|
}
|
|
engine::debug::configure_logging({true, argv[2]});
|
|
const auto started = std::chrono::steady_clock::now();
|
|
int result = 0;
|
|
try {
|
|
engine::core::ensure_backends_loaded();
|
|
auto reg = ggml_backend_reg_by_name("CUDA");
|
|
if (!reg || ggml_backend_reg_dev_count(reg) == 0) {
|
|
std::cerr << "SKIP: CUDA device unavailable\n";
|
|
result = 77;
|
|
} else {
|
|
std::unique_ptr<ggml_backend, decltype(&ggml_backend_free)> backend(
|
|
ggml_backend_dev_init(ggml_backend_reg_dev_get(reg, 0), nullptr), ggml_backend_free);
|
|
check(bool(backend), "CUDA backend initialization failed");
|
|
engine::modules::Conv3dConfig c;
|
|
c.in_channels = 3; c.out_channels = 5;
|
|
c.kernel_depth = 2; c.kernel_height = 3; c.kernel_width = 2;
|
|
c.stride_depth = 2; c.stride_height = 1; c.stride_width = 2;
|
|
c.padding_depth = 1; c.padding_height = 2; c.padding_width = 1;
|
|
c.dilation_depth = 2; c.dilation_height = 1; c.dilation_width = 2;
|
|
c.use_bias = false;
|
|
const bool unavailable = argc == 4;
|
|
for (int channels : {3, 17, 128, 512, 513}) {
|
|
for (bool f32_input : {false, true}) {
|
|
for (bool coalesced : {false, true}) {
|
|
run_norm_case(backend.get(), channels, f32_input, coalesced, !unavailable);
|
|
}
|
|
}
|
|
}
|
|
run_case(backend.get(), c, 2, 5, 6, 7, false, false, unavailable);
|
|
run_case(backend.get(), c, 2, 5, 6, 7, false, true, unavailable);
|
|
c.use_bias = true;
|
|
run_case(backend.get(), c, 2, 5, 6, 7, true, true, unavailable);
|
|
for (int left : {0, 1, 2}) {
|
|
auto padded = c;
|
|
padded.cuda_replicate_depth_left = left;
|
|
for (bool f32_input : {false, true}) {
|
|
padded.cuda_f16_output = false;
|
|
run_case(backend.get(), padded, 2, 5, 6, 7, true, f32_input, unavailable);
|
|
padded.cuda_f16_output = true;
|
|
run_case(backend.get(), padded, 2, 5, 6, 7, true, f32_input, unavailable, false, true);
|
|
}
|
|
}
|
|
if (!unavailable) {
|
|
auto tiled = c;
|
|
tiled.in_channels = 256; tiled.out_channels = 2048;
|
|
tiled.kernel_depth = tiled.kernel_height = tiled.kernel_width = 3;
|
|
tiled.stride_depth = tiled.stride_height = tiled.stride_width = 1;
|
|
tiled.dilation_depth = tiled.dilation_height = tiled.dilation_width = 1;
|
|
tiled.padding_depth = tiled.padding_height = tiled.padding_width = 1;
|
|
tiled.cuda_replicate_depth_left = 2;
|
|
tiled.cuda_f16_output = true;
|
|
run_case(backend.get(), tiled, 2, 9, 32, 32, true, true, false, false, true, true);
|
|
}
|
|
for (bool f16_output : {false, true}) {
|
|
auto packed = c;
|
|
packed.cuda_f16_output = f16_output;
|
|
for (bool f32_input : {false, true}) {
|
|
run_case(backend.get(), packed, 2, 5, 6, 7, false, f32_input, unavailable, false, true);
|
|
}
|
|
}
|
|
auto typed = c;
|
|
typed.cuda_cudnn = true;
|
|
typed.cuda_f16_output = true;
|
|
for (bool cudnn : {false, true}) {
|
|
auto invalid_typed = typed;
|
|
invalid_typed.cuda_cudnn = cudnn;
|
|
invalid_typed.cuda_bf16_round = cudnn;
|
|
bool rejected_typed = false;
|
|
try {
|
|
engine::modules::Conv3dModule module(invalid_typed);
|
|
} catch (const std::runtime_error &) {
|
|
rejected_typed = true;
|
|
}
|
|
check(rejected_typed, "invalid F16 output configuration accepted");
|
|
}
|
|
typed.use_bias = false;
|
|
run_case(backend.get(), typed, 2, 5, 6, 7, false, false, unavailable);
|
|
typed.use_bias = true;
|
|
run_case(backend.get(), typed, 2, 5, 6, 7, false, true, unavailable);
|
|
run_case(backend.get(), typed, 2, 5, 6, 7, true, false, unavailable);
|
|
run_case(backend.get(), typed, 2, 5, 6, 7, true, true, unavailable);
|
|
auto invalid = c;
|
|
invalid.cuda_bf16_round = true;
|
|
bool rejected = false;
|
|
try {
|
|
engine::modules::Conv3dModule module(invalid);
|
|
} catch (const std::runtime_error &) {
|
|
rejected = true;
|
|
}
|
|
check(rejected, "BF16 rounding accepted without cuDNN selection");
|
|
c.cuda_bf16_round = true;
|
|
run_case(backend.get(), c, 2, 5, 6, 7, true, true, unavailable);
|
|
c.use_bias = false;
|
|
run_case(backend.get(), c, 2, 5, 6, 7, false, false, unavailable);
|
|
c.cuda_bf16_round = false;
|
|
c.use_bias = true;
|
|
c.kernel_depth = c.kernel_height = c.kernel_width = 1;
|
|
c.stride_depth = c.stride_height = c.stride_width = 1;
|
|
c.padding_depth = c.padding_height = c.padding_width = 0;
|
|
c.dilation_depth = c.dilation_height = c.dilation_width = 1;
|
|
run_case(backend.get(), c, 1, 3, 4, 5, true, false, unavailable);
|
|
c.in_channels = c.out_channels = 1;
|
|
c.use_bias = false;
|
|
run_case(backend.get(), c, 2, 3, 4, 5, true, true, unavailable, true);
|
|
c.use_bias = true;
|
|
run_case(backend.get(), c, 2, 3, 4, 5, true, true, unavailable, true);
|
|
c.cuda_bf16_round = true;
|
|
run_case(backend.get(), c, 2, 3, 4, 5, true, true, unavailable, true);
|
|
run_case(backend.get(), c, 2, 3, 4, 5, false, false, unavailable, true);
|
|
c.use_bias = false;
|
|
run_case(backend.get(), c, 2, 3, 4, 5, true, true, unavailable, true);
|
|
typed = c;
|
|
typed.cuda_cudnn = true;
|
|
typed.cuda_bf16_round = false;
|
|
typed.cuda_f16_output = true;
|
|
run_case(backend.get(), typed, 2, 3, 4, 5, true, true, unavailable, true);
|
|
typed.use_bias = true;
|
|
run_case(backend.get(), typed, 2, 3, 4, 5, true, false, unavailable, true);
|
|
std::cout << "PASS: cuDNN Conv3D contracts" << (unavailable ? " (disabled build)" : " and scalar parity") << '\n';
|
|
}
|
|
} catch (const std::exception & error) {
|
|
std::cerr << error.what() << '\n';
|
|
result = 1;
|
|
}
|
|
engine::debug::timing_log_scalar("session.wall_ms", engine::debug::elapsed_ms(started));
|
|
return result;
|
|
}
|