diff --git a/core/inc/SOFIE/RModel.hxx b/core/inc/SOFIE/RModel.hxx index 8153408..8fbcf20 100644 --- a/core/inc/SOFIE/RModel.hxx +++ b/core/inc/SOFIE/RModel.hxx @@ -58,7 +58,7 @@ private: std::vector opIndices; ///< consecutive op indices forming this group std::string inputTensor; ///< input tensor name of the first op std::string outputTensor; ///< output tensor name of the last op - size_t numElements = 0; + std::string lengthExpr; ///< element count: literal for static tensors, runtime expression for dynamic bool isFused() const { return opIndices.size() > 1; } std::string suffix() const { std::string s; diff --git a/core/inc/SOFIE/ROperator_BasicBinary.hxx b/core/inc/SOFIE/ROperator_BasicBinary.hxx index 9a1a963..0f3c233 100644 --- a/core/inc/SOFIE/ROperator_BasicBinary.hxx +++ b/core/inc/SOFIE/ROperator_BasicBinary.hxx @@ -451,97 +451,109 @@ public: return out.str(); } + // Broadcast layout shared by the GPU kernel generation and the launch call, + // computed once so the kernel signature and the call site cannot drift apart. + // Shapes are the Dim members, which are filled in every non-constant path of + // Initialize (for static inputs via ConvertShapeToDim), so one code path + // covers static and dynamic models. + struct GPUBroadcastInfo { + std::vector dimA; + std::vector dimB; + std::vector bcastA; + std::vector bcastB; + bool isAScalar = true; + bool isBScalar = true; + bool isAContiguous = true; + bool isBContiguous = true; + bool needCoords = false; + std::vector dynParams; + }; + + GPUBroadcastInfo GetGPUBroadcastInfo() const { + GPUBroadcastInfo info; + const std::size_t D = fDimShapeY.size(); + // right-aligned rank padding (ONNX broadcast); no-op when Initialize + // already padded the shapes in place + info.dimA.assign(D, Dim{1}); + info.dimB.assign(D, Dim{1}); + for (std::size_t i = 0; i < fDimShapeA.size(); i++) + info.dimA[D - fDimShapeA.size() + i] = fDimShapeA[i]; + for (std::size_t i = 0; i < fDimShapeB.size(); i++) + info.dimB[D - fDimShapeB.size() + i] = fDimShapeB[i]; + + info.bcastA.resize(D); + info.bcastB.resize(D); + for (std::size_t d = 0; d < D; d++) { + // only a static 1 marks a broadcast dimension, a parametric dim is never 1 + info.bcastA[d] = !info.dimA[d].isParam && info.dimA[d].dim == 1; + info.bcastB[d] = !info.dimB[d].isParam && info.dimB[d].dim == 1; + if (!info.bcastA[d]) info.isAScalar = false; + if (!info.bcastB[d]) info.isBScalar = false; + if (info.dimA[d].GetVal() != fDimShapeY[d].GetVal()) info.isAContiguous = false; + if (info.dimB[d].GetVal() != fDimShapeY[d].GetVal()) info.isBContiguous = false; + } + bool generalA = !info.isAScalar && !info.isAContiguous; + bool generalB = !info.isBScalar && !info.isBContiguous; + info.needCoords = generalA || generalB; + if (!info.needCoords) + return info; + + // symbolic names in the emitted stride expressions become size_t kernel args + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fDimShapeY), info.dynParams); + if (generalA) UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(info.dimA), info.dynParams); + if (generalB) UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(info.dimB), info.dynParams); + return info; + } + std::string Generate_GPU_Kernel_ALPAKA(std::string opName) { if (fIsOutputConstant) return ""; + if (fDimShapeY.empty()) + throw std::runtime_error("SOFIE Operator Basic Binary called to Generate without being initialized first"); + + auto info = GetGPUBroadcastInfo(); + const std::size_t D = fDimShapeY.size(); + auto stridesY = UTILITY::ComputeStrideFromShape(fDimShapeY); + auto stridesA = UTILITY::ComputeStrideFromShape(info.dimA); + auto stridesB = UTILITY::ComputeStrideFromShape(info.dimB); + bool generalA = !info.isAScalar && !info.isAContiguous; + bool generalB = !info.isBScalar && !info.isBContiguous; std::string op; op = "\n//------ "+opName+"_"+BinaryOperatorTrait::Name()+"_KERNEL_ALPAKA\n"; op += SP + "struct Binary"+opName+BinaryOperatorTrait::Name()+"Kernel {\n"; op += SP + SP + "template\n"; - op += SP + SP + "ALPAKA_FN_ACC void operator()(TAcc const & acc, T const * A, T const * B, T * C) const {\n"; + op += SP + SP + "ALPAKA_FN_ACC void operator()(TAcc const & acc, T const * A, T const * B, T * C"; + for (auto &p : info.dynParams) + op += ", std::size_t const " + p; + op += ", std::size_t const totalElements) const {\n"; op += SP + SP + SP + "auto idx = alpaka::getIdx(acc)[0];\n"; - op += SP + SP + SP + "if (idx < " + std::to_string(ConvertShapeToLength(fShapeY)) + ") {\n"; - auto stridesA = UTILITY::ComputeStrideFromShape(fShapeA); - auto stridesB = UTILITY::ComputeStrideFromShape(fShapeB); - - for(size_t id_s = 0; id_s < stridesA.size(); ++id_s){ - if(fShapeA[id_s] == 1) - stridesA[id_s] = 0; - } - - for(size_t id_s = 0; id_s < stridesB.size(); ++id_s){ - if(fShapeB[id_s] == 1) - stridesB[id_s] = 0; - } - - auto stridesY = UTILITY::ComputeStrideFromShape(fShapeY); - - // --- Fast-path index simplifications --- - // Check whether A is broadcast (all strides zero → single element) - bool isAScalar = true; - for (const auto& s : stridesA) { if (s != 0) { isAScalar = false; break; } } - // Check whether B is broadcast (all strides zero → single element) - bool isBScalar = true; - for (const auto& s : stridesB) { if (s != 0) { isBScalar = false; break; } } - // Check whether A has the same contiguous layout as Y (no broadcasting) - bool isAContiguous = (fShapeA.size() == fShapeY.size()); - if (isAContiguous) { - for (size_t i = 0; i < fShapeA.size(); ++i) - if (fShapeA[i] != fShapeY[i]) { isAContiguous = false; break; } - } - // Check whether B has the same contiguous layout as Y (no broadcasting) - bool isBContiguous = (fShapeB.size() == fShapeY.size()); - if (isBContiguous) { - for (size_t i = 0; i < fShapeB.size(); ++i) - if (fShapeB[i] != fShapeY[i]) { isBContiguous = false; break; } - } - - std::string flattened_index_A = ""; - std::string flattened_index_B = ""; - - if (isAScalar) { - // A is a single broadcast value - flattened_index_A = "0"; - } else if (isAContiguous) { - // A and Y have identical shapes → direct index - flattened_index_A = "idx"; - } else { - // General broadcast case: decompose idx into per-dim coords - std::string temp = "idx"; - for (size_t id_s = 0; id_s < fShapeA.size(); ++id_s) { - auto strideY = stridesY[id_s]; - auto strideA = stridesA[id_s]; - std::string coord = "(int)(" + temp + " / " + std::to_string(strideY) + ")"; - flattened_index_A += coord + " * " + std::to_string(strideA) + " + "; - temp = temp + " - (" + coord + " * " + std::to_string(strideY) + ")"; - } - if (!flattened_index_A.empty()) - flattened_index_A.erase(flattened_index_A.size() - 3); - } - - if (isBScalar) { - // B is a single broadcast value - flattened_index_B = "0"; - } else if (isBContiguous) { - // B and Y have identical shapes → direct index - flattened_index_B = "idx"; - } else { - // General broadcast case - std::string temp = "idx"; - for (size_t id_s = 0; id_s < fShapeB.size(); ++id_s) { - auto strideY = stridesY[id_s]; - auto strideB = stridesB[id_s]; - std::string coord = "(int)(" + temp + " / " + std::to_string(strideY) + ")"; - flattened_index_B += coord + " * " + std::to_string(strideB) + " + "; - temp = temp + " - (" + coord + " * " + std::to_string(strideY) + ")"; + op += SP + SP + SP + "if (idx < totalElements) {\n"; + + // general broadcast case: decompose idx into per-dim output coords and + // accumulate the input index, skipping broadcast dims (stride 0) + if (info.needCoords) { + op += SP + SP + SP + SP + "std::size_t remaining = idx;\n"; + op += SP + SP + SP + SP + "std::size_t coord;\n"; + if (generalA) op += SP + SP + SP + SP + "std::size_t idxA = 0;\n"; + if (generalB) op += SP + SP + SP + SP + "std::size_t idxB = 0;\n"; + for (std::size_t d = 0; d < D; d++) { + std::string sY = "(" + stridesY[d].GetVal() + ")"; + op += SP + SP + SP + SP + "coord = remaining / " + sY + ";\n"; + if (d + 1 < D) + op += SP + SP + SP + SP + "remaining -= coord * " + sY + ";\n"; + if (generalA && !info.bcastA[d]) + op += SP + SP + SP + SP + "idxA += coord * (" + stridesA[d].GetVal() + ");\n"; + if (generalB && !info.bcastB[d]) + op += SP + SP + SP + SP + "idxB += coord * (" + stridesB[d].GetVal() + ");\n"; } - if (!flattened_index_B.empty()) - flattened_index_B.erase(flattened_index_B.size() - 3); } - - op += "C[idx] = " + BinaryOperatorTrait::Op("A["+flattened_index_A+"]", "B["+flattened_index_B+"]") + ";\n"; - op += "}\n}\n};\n"; + std::string indexA = info.isAScalar ? "0" : (info.isAContiguous ? "idx" : "idxA"); + std::string indexB = info.isBScalar ? "0" : (info.isBContiguous ? "idx" : "idxB"); + op += SP + SP + SP + SP + "C[idx] = " + BinaryOperatorTrait::Op("A["+indexA+"]", "B["+indexB+"]") + ";\n"; + op += SP + SP + SP + "}\n"; + op += SP + SP + "}\n"; + op += SP + "};\n"; return op; } @@ -561,13 +573,18 @@ public: } std::stringstream out; auto length = ConvertDimShapeToLength(fDimShapeY); + auto info = GetGPUBroadcastInfo(); out << "\n//------ "+OpName+"_ALPAKA\n"; out << SP << "auto const elementsPerThread_"<(1));\n"; out << SP << "auto const elementsPerGrid_"<(workDiv_" << fNY << ", binary" << OpName << "Kernel, alpaka::getPtrNative(deviceBuf_" << fNA - << "), alpaka::getPtrNative(deviceBuf_" << fNB << "), alpaka::getPtrNative(deviceBuf_" << fNY << "));\n"; + << "), alpaka::getPtrNative(deviceBuf_" << fNB << "), alpaka::getPtrNative(deviceBuf_" << fNY << ")"; + // dynamic shape params, same order as the kernel signature + for (auto &p : info.dynParams) + out << ", static_cast(" << p << ")"; + out << ", static_cast(" << length << "));\n"; out << SP << "alpaka::enqueue(queue, task_" << OpName << ");\n"; return out.str(); } diff --git a/core/inc/SOFIE/ROperator_BatchNormalization.hxx b/core/inc/SOFIE/ROperator_BatchNormalization.hxx index 8bc3b3d..02cace4 100644 --- a/core/inc/SOFIE/ROperator_BatchNormalization.hxx +++ b/core/inc/SOFIE/ROperator_BatchNormalization.hxx @@ -28,14 +28,11 @@ private: std::string fNMean; std::string fNVar; std::string fNY; + std::string fNFusedScale; // scale/sqrt(var+eps) fused over channels, shape [C] EActivationType fActivation; - std::vector fShapeX; - std::vector fShapeScale; - std::vector fShapeB; - std::vector fShapeMean; - std::vector fShapeVar; - std::vector fShapeY; + std::vector fShapeX; + std::vector fShapeY; std::string fType; @@ -54,6 +51,9 @@ public: fInputTensorNames = { fNX }; fOutputTensorNames = { fNY }; + // fused per-channel scale tensor (created in Initialize) + fNFusedScale = fNScale + "_fused_inv_std_dev"; + if(std::is_same::value){ fType = "float"; } @@ -74,156 +74,102 @@ public: throw std::runtime_error("SOFIE BatchNormalization Op Shape inference need 5 input tensors"); } - for(size_t i = 0; i < input.size(); i++) { - if (input[i].size() != 4) { - throw - std::runtime_error("SOFIE BatchNormalization Op Shape inference only accept tensor with 4 dimensions"); - } - } - auto ret = input; return ret; } void Initialize(RModel& model) override { if (!model.CheckIfTensorAlreadyExist(fNX)) { - throw - std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNX + " fnx is not found in model"); + throw std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNX + " is not found in model"); } if (!model.CheckIfTensorAlreadyExist(fNScale)) { - throw - std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNScale + " fns is not found in model"); + throw std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNScale + " is not found in model"); } - if (!model.CheckIfTensorAlreadyExist(fNB)) { - throw - std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNB + " fnb is not found in model"); + if (!model.CheckIfTensorAlreadyExist(fNB)) { + throw std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNB + " is not found in model"); } if (!model.CheckIfTensorAlreadyExist(fNMean)) { - throw - std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNMean + " fnm is not found in model"); + throw std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNMean + " is not found in model"); } if (!model.CheckIfTensorAlreadyExist(fNVar)) { - throw - std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNVar + " fnv is not found in model"); + throw std::runtime_error("SOFIE BatchNormalization op Input Tensor " + fNVar + " is not found in model"); } - fShapeX = model.GetTensorShape(fNX); - - if (fShapeX.size() < 2 || fShapeX.size() > 4) { - throw - std::runtime_error("SOFIE BatchNormalization Op input tensor " + fNX + " fnx has wrong shape : " + ConvertShapeToString(fShapeX)); + fShapeX = model.GetDimTensorShape(fNX); + // Rank is kept at 2-4, matching ROOT SOFIE's BatchNorm constraint + if (fShapeX.size() < 2 || fShapeX.size() > 4) { + throw std::runtime_error("SOFIE BatchNormalization Op input tensor " + fNX + + " has wrong shape : " + ConvertDimShapeToString(fShapeX)); } - fShapeScale = model.GetTensorShape(fNScale); - fShapeB = model.GetTensorShape(fNB); - fShapeMean = model.GetTensorShape(fNMean); - fShapeVar = model.GetTensorShape(fNVar); fShapeY = fShapeX; model.AddIntermediateTensor(fNY, model.GetTensorType(fNX), fShapeY); - if (fShapeB.size() == 1) { - // Broadcast scale, bias, input_mean and input_var to shape_X - auto original_B = model.GetInitializedTensorData(fNB); - auto original_S = model.GetInitializedTensorData(fNScale); - auto original_M = model.GetInitializedTensorData(fNMean); - auto original_V = model.GetInitializedTensorData(fNVar); - size_t batchSize = fShapeX[0]; - size_t channels = fShapeX[1]; - size_t height = (fShapeX.size() > 2) ? fShapeX[2] : 1; - size_t width = (fShapeX.size() > 3) ? fShapeX[3] : 1; - size_t n = batchSize * channels * height * width; - if (fType == "float") { - float *original_bias = static_cast(original_B.get()); - float *original_scale = static_cast(original_S.get()); - float *original_mean = static_cast(original_M.get()); - float *original_var = static_cast(original_V.get()); - float *new_bias = new float[n]; - float *new_scale = new float[n]; - float *new_mean = new float[n]; - float *new_var = new float[n]; - size_t bs = 0, ch = 0, h = 0, w = 0; - for (ch = 0; ch < channels; ch++) { - for (h = 0; h < height; h++) { - for (w = 0; w < width; w++) { - new_bias[bs * channels * height * width + ch * height * width + h * width + w] = original_bias[ch]; - new_scale[bs * channels * height * width + ch * height * width + h * width + w] = - original_scale[ch]; - new_mean[bs * channels * height * width + ch * height * width + h * width + w] = original_mean[ch]; - new_var[bs * channels * height * width + ch * height * width + h * width + w] = original_var[ch]; - } - } - } - size_t Batchoffset = channels * height * width; - for (bs = 1; bs < batchSize; bs++) { - std::copy(new_bias, new_bias + Batchoffset, new_bias + (bs * Batchoffset)); - std::copy(new_scale, new_scale + Batchoffset, new_scale + (bs * Batchoffset)); - std::copy(new_mean, new_mean + Batchoffset, new_mean + (bs * Batchoffset)); - std::copy(new_var, new_var + Batchoffset, new_var + (bs * Batchoffset)); - } - //// new_var =1. / sqrt(input_var + fepsilon) - for (size_t i = 0; i < n; i++) { - new_var[i] = 1. / sqrt(new_var[i] + fepsilon); - new_scale[i] *= new_var[i]; // include var in new scale - } - std::vector new_bias_shape = {batchSize, channels, height, width}; - std::shared_ptr new_bias_ptr(new_bias, std::default_delete()); - std::shared_ptr new_scale_ptr(new_scale, std::default_delete()); - std::shared_ptr new_mean_ptr(new_mean, std::default_delete()); - std::shared_ptr new_var_ptr(new_var, std::default_delete()); - model.UpdateInitializedTensor(fNB, model.GetTensorType(fNB), new_bias_shape, new_bias_ptr); - model.UpdateInitializedTensor(fNScale, model.GetTensorType(fNScale), new_bias_shape, new_scale_ptr); - model.UpdateInitializedTensor(fNMean, model.GetTensorType(fNMean), new_bias_shape, new_mean_ptr); - model.UpdateInitializedTensor(fNVar, model.GetTensorType(fNVar), new_bias_shape, new_var_ptr); - fShapeB = model.GetTensorShape(fNB); - fShapeScale = model.GetTensorShape(fNScale); - fShapeMean = model.GetTensorShape(fNMean); - fShapeVar = model.GetTensorShape(fNVar); + // Fuse scale with the variance over channels only -> a [C] array, instead of materializing + // the weights to the full [N,C,...] tensor (which would bake the batch size and block any dynamic shape). + //The batch and spatial dims are handled by the runtime index math in the + // kernel / Generate. + auto original_S = model.GetInitializedTensorData(fNScale); + auto original_V = model.GetInitializedTensorData(fNVar); + auto shape_S = model.GetTensorShape(fNScale); + if (shape_S.size() != 1) { + throw std::runtime_error("SOFIE BatchNormalization 'scale' tensor must be 1D (per-channel)."); + } + size_t channels = shape_S[0]; + + //TODO: only float is fused here (mirrors ROOT); add a double branch if needed + if (fType == "float") { + float *original_scale_ptr = static_cast(original_S.get()); + float *original_var_ptr = static_cast(original_V.get()); + float *fused_scale_data = new float[channels]; + for (size_t i = 0; i < channels; i++) { + // scale * (1 / sqrt(variance + epsilon)) + fused_scale_data[i] = original_scale_ptr[i] / std::sqrt(original_var_ptr[i] + fepsilon); } + std::shared_ptr fused_scale_ptr(fused_scale_data, std::default_delete()); + model.AddInitializedTensor(fNFusedScale, model.GetTensorType(fNScale), {channels}, fused_scale_ptr); } } - std::string Generate(std::string OpName) override { - OpName = "op_" + OpName; + std::string Generate(std::string opName) override { + opName = "op_" + opName; if (fShapeX.empty()){ throw std::runtime_error("SOFIE Batch Normalization called to Generate without being initialized first"); } std::stringstream out; - //// Batch Norm op - size_t batchSize = fShapeX[0]; - size_t channels = fShapeX[1]; - size_t height = (fShapeX.size() > 2) ? fShapeX[2] : 1; - size_t width = (fShapeX.size() > 3) ? fShapeX[3] : 1; - size_t n = batchSize * channels * height * width; - - //// copy X into Y - out << "\n\n//---- BatchNorm\n"; - out << SP << "constexpr int " << OpName << "_N =" << batchSize * channels * height * width << ";\n"; - out << SP << "constexpr int "< 2) { + auto spatialShape = fShapeX; + spatialShape.erase(spatialShape.begin(), spatialShape.begin() + 2); + spatial_dim = ConvertDimShapeToLength(spatialShape); + } - if(fActivation == EActivationType::RELU){ - out << SP << "for (int id = 0; id < " << ConvertShapeToLength(fShapeY) << " ; id++){\n"; - out << SP << SP << "tensor_" << fNY << "[id] = ((tensor_" << fNY << "[id] > 0 )? tensor_" << fNY << "[id] : 0);\n"; - out << SP << "}\n"; + // Per-channel affine over a [N, C, spatial] tensor. Weights stay [C] and are indexed by the + // channel c, so batch and spatial can be any (runtime) size. + out << "\n\n//---- BatchNorm" << (fActivation == EActivationType::RELU ? " + ReLU " : " ") << opName << "\n"; + out << SP << "{\n"; + out << SP << " size_t i = 0;\n"; + out << SP << " for (size_t n = 0; n < " << batchSize << "; ++n) {\n"; + out << SP << " for (size_t c = 0; c < " << channels << "; ++c) {\n"; + out << SP << " const float mean_val = tensor_" << fNMean << "[c];\n"; + out << SP << " const float fused_scale_val = tensor_" << fNFusedScale << "[c];\n"; + out << SP << " const float bias_val = tensor_" << fNB << "[c];\n"; + out << SP << " for (size_t sp = 0; sp < " << spatial_dim << "; ++sp) {\n"; + out << SP << " float val = (tensor_" << fNX << "[i] - mean_val) * fused_scale_val + bias_val;\n"; + if (fActivation == EActivationType::RELU) { + out << SP << " tensor_" << fNY << "[i] = (val > 0.0f) ? val : 0.0f;\n"; + } else { + out << SP << " tensor_" << fNY << "[i] = val;\n"; } + out << SP << " i++;\n"; + out << SP << " }\n"; + out << SP << " }\n"; + out << SP << " }\n"; + out << SP << "}\n"; + return out.str(); } @@ -232,7 +178,24 @@ public: if (fShapeX.empty()) throw std::runtime_error("SOFIE BatchNormalization called to Generate without being initialized first"); - std::size_t totalElements = ConvertShapeToLength(fShapeY); + std::string totalElements = ConvertDimShapeToLength(fShapeY); + std::string channels = fShapeX[1].GetVal(); // per-channel weight count (static) + + // spatial_dim = product of dims after [N, C]; "1" for a rank-2 input. + std::string spatial_dim = "1"; + std::vector spatialShape; + if (fShapeX.size() > 2) { + spatialShape.assign(fShapeX.begin() + 2, fShapeX.end()); + spatial_dim = ConvertDimShapeToLength(spatialShape); + } + // symbolic dims inside spatial_dim -> passed as size_t kernel args (kernel-arg convention) + std::vector dynParams; + for (auto &d : spatialShape) + if (d.isParam) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } std::string kname = "BatchNormKernel_" + opName; std::string op; @@ -242,10 +205,12 @@ public: op += SP + SP + "ALPAKA_FN_ACC void operator()(\n"; op += SP + SP + SP + "TAcc const& acc,\n"; op += SP + SP + SP + "T const* __restrict__ X,\n"; - op += SP + SP + SP + "T const* __restrict__ scale,\n"; + op += SP + SP + SP + "T const* __restrict__ fused_scale,\n"; op += SP + SP + SP + "T const* __restrict__ bias,\n"; op += SP + SP + SP + "T const* __restrict__ mean,\n"; op += SP + SP + SP + "T* __restrict__ Y,\n"; + for (auto &p : dynParams) + op += SP + SP + SP + "std::size_t const " + p + ",\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -253,14 +218,13 @@ public: op += SP + SP + SP + "auto const grid_thread_extent = alpaka::getWorkDiv(acc)[0];\n\n"; op += SP + SP + SP + "for (std::size_t i = global_thread_idx; i < totalElements; i += grid_thread_extent) {\n"; - - op += SP + SP + SP + SP + "T val = (X[i] - mean[i]) * scale[i] + bias[i];\n"; - + // weights are per-channel [C]; derive the channel from the flat index (layout [N, C, spatial]). + op += SP + SP + SP + SP + "std::size_t const c = (i / (" + spatial_dim + ")) % (" + channels + ");\n"; + op += SP + SP + SP + SP + "T val = (X[i] - mean[c]) * fused_scale[c] + bias[c];\n"; if (fActivation == EActivationType::RELU) op += SP + SP + SP + SP + "Y[i] = val > static_cast(0) ? val : static_cast(0);\n"; else op += SP + SP + SP + SP + "Y[i] = val;\n"; - op += SP + SP + SP + "}\n"; op += SP + SP + "}\n"; op += SP + "};\n"; @@ -279,29 +243,43 @@ public: if (fShapeX.empty()) throw std::runtime_error("SOFIE BatchNormalization called to Generate without being initialized first"); - std::size_t totalElements = ConvertShapeToLength(fShapeY); + std::string totalElements = ConvertDimShapeToLength(fShapeY); std::string kname = "batchNormKernel_" + opName; + // symbolic spatial dims -> passed after the buffers, matching the kernel signature order. + std::vector dynParams; + if (fShapeX.size() > 2) { + for (auto it = fShapeX.begin() + 2; it != fShapeX.end(); ++it) + if (it->isParam) { + bool seen = false; + for (auto &q : dynParams) if (q == it->param) seen = true; + if (!seen) dynParams.push_back(it->param); + } + } + std::string dynArgs; + for (auto &p : dynParams) dynArgs += ", static_cast(" + p + ")"; + std::stringstream out; out << "\n//------ BATCHNORM_GPU_ALPAKA\n"; out << SP << "auto const elementsPerThread_" << fNY << " = Vec::all(static_cast(1));\n"; out << SP << "auto const elementsPerGrid_" << fNY << " = Vec::all(Idx{" << totalElements << "});\n"; out << SP << "auto const workDiv_" << fNY << " = sofie_workdiv(elementsPerGrid_" << fNY << ");\n"; - + out << SP << "auto task_" << fNY << " = alpaka::createTaskKernel(workDiv_" << fNY << ", " << kname - << ", alpaka::getPtrNative(deviceBuf_" << fNX << ")" - << ", alpaka::getPtrNative(deviceBuf_" << fNScale << ")" - << ", alpaka::getPtrNative(deviceBuf_" << fNB << ")" - << ", alpaka::getPtrNative(deviceBuf_" << fNMean << ")" - << ", alpaka::getPtrNative(deviceBuf_" << fNY << ")" + << ", alpaka::getPtrNative(deviceBuf_" << fNX << ")" + << ", alpaka::getPtrNative(deviceBuf_" << fNFusedScale << ")" + << ", alpaka::getPtrNative(deviceBuf_" << fNB << ")" + << ", alpaka::getPtrNative(deviceBuf_" << fNMean << ")" + << ", alpaka::getPtrNative(deviceBuf_" << fNY << ")" + << dynArgs << ", static_cast(" << totalElements << "));\n"; out << SP <<"alpaka::enqueue(queue, task_" << fNY << ");\n"; - + return out.str(); } - std::vector GetBlasRoutines() override { return { std::string("Copy"), std::string("Axpy") }; } + std::vector GetBlasRoutines() override { return {}; } }; }//SOFIE diff --git a/core/inc/SOFIE/ROperator_Comparision.hxx b/core/inc/SOFIE/ROperator_Comparision.hxx index 1e02d53..6b5e582 100644 --- a/core/inc/SOFIE/ROperator_Comparision.hxx +++ b/core/inc/SOFIE/ROperator_Comparision.hxx @@ -61,6 +61,10 @@ private: std::vector fShapeX1; std::vector fShapeX2; std::vector fShapeY; + std::vector fDimShapeX1; + std::vector fDimShapeX2; + std::vector fDimShapeY; + bool fIsDynamic = false; std::string fNBroadcastedX1; std::string fNBroadcastedX2; ETensorType fTensorType1 = ETensorType::UNDEFINED; @@ -93,10 +97,32 @@ public: if (!model.CheckIfTensorAlreadyExist(fNX2)) { throw std::runtime_error(std::string("SOFIE Comparision Op Input Tensor ") + fNX2 + "is not found in model"); } - fShapeX1 = model.GetTensorShape(fNX1); - fShapeX2 = model.GetTensorShape(fNX2); fTensorType1 = model.GetTensorType(fNX1); fTensorType2 = model.GetTensorType(fNX2); + + fIsDynamic = model.IsDynamicTensor(fNX1) || model.IsDimInputTensor(fNX1) + || model.IsDynamicTensor(fNX2) || model.IsDimInputTensor(fNX2); + if (fIsDynamic) { + // one input may be dynamic while the other is static, so resolve each + // shape on its own; GetDynamicTensorShape throws on a static tensor + fDimShapeX1 = (model.IsDynamicTensor(fNX1) || model.IsDimInputTensor(fNX1)) + ? model.GetDynamicTensorShape(fNX1) + : ConvertShapeToDim(model.GetTensorShape(fNX1)); + fDimShapeX2 = (model.IsDynamicTensor(fNX2) || model.IsDimInputTensor(fNX2)) + ? model.GetDynamicTensorShape(fNX2) + : ConvertShapeToDim(model.GetTensorShape(fNX2)); + if (UTILITY::AreSameShape(fDimShapeX1, fDimShapeX2)) + fDimShapeY = fDimShapeX1; + else + fDimShapeY = UTILITY::MultidirectionalBroadcastShape(fDimShapeX1, fDimShapeX2).second; + model.AddIntermediateTensor(fNY, ETensorType::BOOL, fDimShapeY); + const auto & outNames = model.GetOutputTensorNames(); + fIsModelOutput = (std::find(outNames.begin(), outNames.end(), fNY) != outNames.end()); + return; + } + + fShapeX1 = model.GetTensorShape(fNX1); + fShapeX2 = model.GetTensorShape(fNX2); bool broadcast = !UTILITY::AreSameShape(fShapeX1, fShapeX2); if (broadcast) { // ONNX comparison ops support multidirectional broadcasting (numpy semantics): @@ -196,26 +222,49 @@ public: std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { if (fIsOutputConstant) return ""; opName = "op_" + opName; - if (fShapeY.empty()) + if ((fIsDynamic ? fDimShapeY.empty() : fShapeY.empty())) throw std::runtime_error("SOFIE Comparision Op called to Generate without being initialized first"); - const std::size_t D = fShapeY.size(); - std::size_t totalElements = ConvertShapeToLength(fShapeY); + auto dimX1 = fIsDynamic ? fDimShapeX1 : ConvertShapeToDim(fShapeX1); + auto dimX2 = fIsDynamic ? fDimShapeX2 : ConvertShapeToDim(fShapeX2); + auto dimY = fIsDynamic ? fDimShapeY : ConvertShapeToDim(fShapeY); + + const std::size_t D = dimY.size(); + std::string totalElements = ConvertDimShapeToLength(dimY); - std::vector shapeX1_padded(D, 1); - std::vector shapeX2_padded(D, 1); + std::vector shapeX1_padded(D, Dim(1)); + std::vector shapeX2_padded(D, Dim(1)); { - size_t off1 = D - fShapeX1.size(); - for (size_t i = 0; i < fShapeX1.size(); ++i) - shapeX1_padded[off1 + i] = fShapeX1[i]; - size_t off2 = D - fShapeX2.size(); - for (size_t i = 0; i < fShapeX2.size(); ++i) - shapeX2_padded[off2 + i] = fShapeX2[i]; + size_t off1 = D - dimX1.size(); + for (size_t i = 0; i < dimX1.size(); ++i) + shapeX1_padded[off1 + i] = dimX1[i]; + size_t off2 = D - dimX2.size(); + for (size_t i = 0; i < dimX2.size(); ++i) + shapeX2_padded[off2 + i] = dimX2[i]; } auto stridesX1 = UTILITY::ComputeStrideFromShape(shapeX1_padded); auto stridesX2 = UTILITY::ComputeStrideFromShape(shapeX2_padded); - auto stridesY = UTILITY::ComputeStrideFromShape(fShapeY); + auto stridesY = UTILITY::ComputeStrideFromShape(dimY); + + auto isIdent = [](const std::string &s){ + if (s.empty() || (s[0] >= '0' && s[0] <= '9')) return false; + for (char c : s) + if (!((c>='a'&&c<='z')||(c>='A'&&c<='Z')||(c>='0'&&c<='9')||c=='_')) return false; + return true; + }; + std::vector dynParams; + auto collect = [&](const std::vector& v){ + for (auto &d : v) + if (d.isParam && isIdent(d.param)) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } + }; + collect(dimX1); collect(dimX2); + + auto isOne = [](const Dim &d){ return !d.isParam && d.dim == 1; }; std::string type1 = ConvertTypeToString(fTensorType1); std::string type2 = ConvertTypeToString(fTensorType2); @@ -231,6 +280,8 @@ public: op += SP + SP + SP + type1 + " const* __restrict__ x1,\n"; op += SP + SP + SP + type2 + " const* __restrict__ x2,\n"; op += SP + SP + SP + "uint8_t* __restrict__ output,\n"; + for (auto &p : dynParams) + op += SP + SP + SP + "std::size_t const " + p + ",\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -241,30 +292,30 @@ public: for (std::size_t d = 0; d < D; ++d) { op += SP + SP + SP + SP + "std::size_t const out_" + std::to_string(d) - + " = (elem_idx / " + std::to_string(stridesY[d]) + "u) % " - + std::to_string(fShapeY[d]) + "u;\n"; + + " = (elem_idx / (" + stridesY[d].GetVal() + ")) % (" + + dimY[d].GetVal() + ");\n"; } op += "\n"; op += SP + SP + SP + SP + "std::size_t const x1_idx =\n"; for (std::size_t d = 0; d < D; ++d) { - if (shapeX1_padded[d] == 1) + if (isOne(shapeX1_padded[d])) op += SP + SP + SP + SP + SP + "0u"; else op += SP + SP + SP + SP + SP + "out_" + std::to_string(d) - + " * " + std::to_string(stridesX1[d]) + "u"; + + " * (" + stridesX1[d].GetVal() + ")"; op += (d + 1 < D) ? " +\n" : ";\n\n"; } op += SP + SP + SP + SP + "std::size_t const x2_idx =\n"; for (std::size_t d = 0; d < D; ++d) { - if (shapeX2_padded[d] == 1) + if (isOne(shapeX2_padded[d])) op += SP + SP + SP + SP + SP + "0u"; else op += SP + SP + SP + SP + SP + "out_" + std::to_string(d) - + " * " + std::to_string(stridesX2[d]) + "u"; + + " * (" + stridesX2[d].GetVal() + ")"; op += (d + 1 < D) ? " +\n" : ";\n\n"; } @@ -286,12 +337,34 @@ public: std::string Generate_GPU_ALPAKA(std::string opName) override { if (fIsOutputConstant) return ""; opName = "op_" + opName; - if (fShapeY.empty()) + if ((fIsDynamic ? fDimShapeY.empty() : fShapeY.empty())) throw std::runtime_error("SOFIE Comparision Op called to Generate without being initialized first"); - std::size_t totalElements = ConvertShapeToLength(fShapeY); + auto dimX1 = fIsDynamic ? fDimShapeX1 : ConvertShapeToDim(fShapeX1); + auto dimX2 = fIsDynamic ? fDimShapeX2 : ConvertShapeToDim(fShapeX2); + auto dimY = fIsDynamic ? fDimShapeY : ConvertShapeToDim(fShapeY); + std::string totalElements = ConvertDimShapeToLength(dimY); std::string kname = "comparisonKernel_" + opName; + auto isIdent = [](const std::string &s){ + if (s.empty() || (s[0] >= '0' && s[0] <= '9')) return false; + for (char c : s) + if (!((c>='a'&&c<='z')||(c>='A'&&c<='Z')||(c>='0'&&c<='9')||c=='_')) return false; + return true; + }; + std::vector dynParams; + auto collect = [&](const std::vector& v){ + for (auto &d : v) + if (d.isParam && isIdent(d.param)) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } + }; + collect(dimX1); collect(dimX2); + std::string dynArgs; + for (auto &p : dynParams) dynArgs += ", static_cast(" + p + ")"; + std::stringstream out; out << "\n//------ " << ComparisionTrait::Name() << "_GPU_ALPAKA\n"; out << SP << "auto const elementsPerThread_" << opName << " = Vec::all(static_cast(1));\n"; @@ -302,6 +375,7 @@ public: << ", alpaka::getPtrNative(deviceBuf_" << fNX1 << ")" << ", alpaka::getPtrNative(deviceBuf_" << fNX2 << ")" << ", alpaka::getPtrNative(deviceBuf_" << fNY << ")" + << dynArgs << ", static_cast(" << totalElements << "));\n"; out << SP << "alpaka::enqueue(queue, task_" << opName << ");\n"; diff --git a/core/inc/SOFIE/ROperator_Concat.hxx b/core/inc/SOFIE/ROperator_Concat.hxx index 36ede27..2a4f9ed 100644 --- a/core/inc/SOFIE/ROperator_Concat.hxx +++ b/core/inc/SOFIE/ROperator_Concat.hxx @@ -375,6 +375,18 @@ return out.str(); } + // symbolic names used by the kernel index expressions, shared by the + // kernel signature and the launch call + std::vector GetGPUDynParams() const { + std::vector params; + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fOutputShape), params); + for (std::size_t k = 0; k < fInputShapes.size(); ++k) { + UTILITY::CollectDimParams({fInputShapes[k][fAxis]}, params); + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fInputShapes[k]), params); + } + return params; + } + std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { if (fIsOutputConstant || fIsOutputParamShape) return ""; opName = "op_" + opName; @@ -386,10 +398,12 @@ auto outStrides = UTILITY::ComputeStrideFromShape(fOutputShape); - std::vector prefix(Nin); - prefix[0] = 0; + // cumulative offsets along the concat axis, kept as expression strings + // since an axis dim can be symbolic + std::vector prefix(Nin); + prefix[0] = "0"; for (std::size_t k = 1; k < Nin; ++k) - prefix[k] = prefix[k - 1] + std::stoul(fInputShapes[k - 1][fAxis].GetVal()); + prefix[k] = prefix[k - 1] + " + (" + fInputShapes[k - 1][fAxis].GetVal() + ")"; std::vector> inStrides(Nin); for (std::size_t k = 0; k < Nin; ++k) @@ -403,6 +417,8 @@ op += SP + SP + SP + "TAcc const& acc,\n"; op += SP + SP + SP + "std::array inputs,\n"; op += SP + SP + SP + "T* output,\n"; + for (auto &p : GetGPUDynParams()) + op += SP + SP + SP + "std::size_t const " + p + ",\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -414,26 +430,26 @@ op += SP + SP + SP + SP + "remaining = elem_idx;\n"; for (std::size_t d = 0; d < D; ++d) { - std::string stride_val = outStrides[d].GetVal(); + std::string stride_val = "(" + outStrides[d].GetVal() + ")"; op += SP + SP + SP + SP + "std::size_t const out_" + std::to_string(d) - + " = remaining / " + stride_val + "u;\n"; + + " = remaining / " + stride_val + ";\n"; op += SP + SP + SP + SP + "remaining -= out_" + std::to_string(d) - + " * " + stride_val + "u;\n"; + + " * " + stride_val + ";\n"; } op += "\n"; op += SP + SP + SP + SP + "std::size_t chosen = 0;\n"; for (std::size_t k = 0; k < Nin; ++k) { - std::size_t end_k = prefix[k] + std::stoul(fInputShapes[k][fAxis].GetVal()); + std::string end_k = "(" + prefix[k] + " + (" + fInputShapes[k][fAxis].GetVal() + "))"; op += SP + SP + SP + SP + "chosen += static_cast(" - + std::to_string(end_k) + "u <= out_" + std::to_string(fAxis) + ");\n"; + + end_k + " <= out_" + std::to_string(fAxis) + ");\n"; } op += "\n"; op += SP + SP + SP + SP + "std::size_t const output_idx =\n"; for (std::size_t d = 0; d < D; ++d) { op += SP + SP + SP + SP + SP + "out_" + std::to_string(d) - + " * " + outStrides[d].GetVal() + "u"; + + " * (" + outStrides[d].GetVal() + ")"; op += (d + 1 < D) ? " +\n" : ";\n\n"; } @@ -442,10 +458,10 @@ op += SP + SP + SP + SP + SP + "(chosen == " + std::to_string(k) + "u) * (\n"; for (std::size_t d = 0; d < D; ++d) { std::string coord = (d == static_cast(fAxis)) - ? ("(out_" + std::to_string(d) + " - " + std::to_string(prefix[k]) + "u)") + ? ("(out_" + std::to_string(d) + " - (" + prefix[k] + "))") : ("out_" + std::to_string(d)); op += SP + SP + SP + SP + SP + SP + coord - + " * " + inStrides[k][d].GetVal() + "u"; + + " * (" + inStrides[k][d].GetVal() + ")"; op += (d + 1 < D) ? " +\n" : "\n"; } op += SP + SP + SP + SP + SP + ")"; @@ -493,7 +509,10 @@ out << SP << "auto const elementsPerGrid_"<(workDiv_" << OpName - << ", concatKernel_" << OpName << ", input_ptrs_" << OpName << ", alpaka::getPtrNative(deviceBuf_" << fOutput << "), static_cast(" << length << "));\n"; + << ", concatKernel_" << OpName << ", input_ptrs_" << OpName << ", alpaka::getPtrNative(deviceBuf_" << fOutput << ")"; + for (auto &p : GetGPUDynParams()) + out << ", static_cast(" << p << ")"; + out << ", static_cast(" << length << "));\n"; out << SP << "alpaka::enqueue(queue, task_" << OpName << ");\n"; return out.str(); } diff --git a/core/inc/SOFIE/ROperator_Conv.hxx b/core/inc/SOFIE/ROperator_Conv.hxx index 835a0ff..e893e9f 100644 --- a/core/inc/SOFIE/ROperator_Conv.hxx +++ b/core/inc/SOFIE/ROperator_Conv.hxx @@ -202,7 +202,7 @@ public: } else { // general case (stride not 1) int64_t v = pad - kernel; std::string outStr = "((" + inputDim.param + "+" + std::to_string(v) + ")/" - + std::to_string(stride) + "1)"; + + std::to_string(stride) + "+1)"; return Dim{ outStr, static_cast(-1)}; } } @@ -695,6 +695,8 @@ public: op += SP + SP + SP + "TAcc const& acc,\n"; op += SP + SP + SP + "T const* __restrict__ input,\n"; op += SP + SP + SP + "T* __restrict__ col,\n"; + op += SP + SP + SP + "std::size_t const oDepth, std::size_t const oHeight, std::size_t const oWidth,\n"; + op += SP + SP + SP + "std::size_t const iDepth, std::size_t const iHeight, std::size_t const iWidth,\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -703,8 +705,8 @@ public: op += SP + SP + SP + "for (std::size_t elem_idx = global_thread_idx; elem_idx < totalElements; elem_idx += grid_thread_extent) {\n\n"; - op += SP + SP + SP + SP + "std::size_t const col_row = elem_idx / " + std::to_string(colCols) + "u;\n"; - op += SP + SP + SP + SP + "std::size_t const col_col = elem_idx % " + std::to_string(colCols) + "u;\n\n"; + op += SP + SP + SP + SP + "std::size_t const col_row = elem_idx / (oDepth * oHeight * oWidth);\n"; + op += SP + SP + SP + SP + "std::size_t const col_col = elem_idx % (oDepth * oHeight * oWidth);\n\n"; op += SP + SP + SP + SP + "std::size_t const ic = col_row / " + std::to_string(kernelSize) + "u;\n"; op += SP + SP + SP + SP + "std::size_t const k_rem = col_row % " + std::to_string(kernelSize) + "u;\n"; @@ -723,13 +725,13 @@ public: } if (fDim > 2) { - op += SP + SP + SP + SP + "std::size_t const od = col_col / " + std::to_string(oHeight * oWidth) + "u;\n"; - op += SP + SP + SP + SP + "std::size_t const oh = (col_col / " + std::to_string(oWidth) + "u) % " + std::to_string(oHeight) + "u;\n"; - op += SP + SP + SP + SP + "std::size_t const ow = col_col % " + std::to_string(oWidth) + "u;\n\n"; + op += SP + SP + SP + SP + "std::size_t const od = col_col / (oHeight * oWidth);\n"; + op += SP + SP + SP + SP + "std::size_t const oh = (col_col / oWidth) % oHeight;\n"; + op += SP + SP + SP + SP + "std::size_t const ow = col_col % oWidth;\n\n"; } else if (fDim > 1) { op += SP + SP + SP + SP + "std::size_t const od = 0u;\n"; - op += SP + SP + SP + SP + "std::size_t const oh = col_col / " + std::to_string(oWidth) + "u;\n"; - op += SP + SP + SP + SP + "std::size_t const ow = col_col % " + std::to_string(oWidth) + "u;\n\n"; + op += SP + SP + SP + SP + "std::size_t const oh = col_col / oWidth;\n"; + op += SP + SP + SP + SP + "std::size_t const ow = col_col % oWidth;\n\n"; } else { op += SP + SP + SP + SP + "std::size_t const od = 0u;\n"; op += SP + SP + SP + SP + "std::size_t const oh = 0u;\n"; @@ -763,15 +765,15 @@ public: } op += SP + SP + SP + SP + "bool const in_bounds =\n"; - op += SP + SP + SP + SP + SP + "id_in >= 0 && id_in < " + std::to_string(iDepth) + " &&\n"; - op += SP + SP + SP + SP + SP + "ih_in >= 0 && ih_in < " + std::to_string(iHeight) + " &&\n"; - op += SP + SP + SP + SP + SP + "iw_in >= 0 && iw_in < " + std::to_string(iWidth) + ";\n\n"; + op += SP + SP + SP + SP + SP + "id_in >= 0 && id_in < static_cast(iDepth) &&\n"; + op += SP + SP + SP + SP + SP + "ih_in >= 0 && ih_in < static_cast(iHeight) &&\n"; + op += SP + SP + SP + SP + SP + "iw_in >= 0 && iw_in < static_cast(iWidth) ;\n\n"; op += SP + SP + SP + SP + "if (in_bounds) {\n"; op += SP + SP + SP + SP + SP + "std::size_t const in_idx =\n"; - op += SP + SP + SP + SP + SP + SP + "ic * " + std::to_string(iDepth * iHeight * iWidth) + "u +\n"; - op += SP + SP + SP + SP + SP + SP + "static_cast(id_in) * " + std::to_string(iHeight * iWidth) + "u +\n"; - op += SP + SP + SP + SP + SP + SP + "static_cast(ih_in) * " + std::to_string(iWidth) + "u +\n"; + op += SP + SP + SP + SP + SP + SP + "ic * (iDepth * iHeight * iWidth) +\n"; + op += SP + SP + SP + SP + SP + SP + "static_cast(id_in) * (iHeight * iWidth) +\n"; + op += SP + SP + SP + SP + SP + SP + "static_cast(ih_in) * iWidth +\n"; op += SP + SP + SP + SP + SP + SP + "static_cast(iw_in);\n"; op += SP + SP + SP + SP + SP + "col[elem_idx] = input[in_idx];\n"; op += SP + SP + SP + SP + "} else {\n"; @@ -791,6 +793,7 @@ public: op += SP + SP + SP + "TAcc const& acc,\n"; op += SP + SP + SP + "T const* __restrict__ bias,\n"; op += SP + SP + SP + "T* __restrict__ output,\n"; + op += SP + SP + SP + "std::size_t const spatialSize,\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -798,7 +801,7 @@ public: op += SP + SP + SP + "auto const grid_thread_extent = alpaka::getWorkDiv(acc)[0];\n\n"; op += SP + SP + SP + "for (std::size_t elem_idx = global_thread_idx; elem_idx < totalElements; elem_idx += grid_thread_extent) {\n"; - op += SP + SP + SP + SP + "std::size_t const channel = elem_idx / " + std::to_string(spatialSize) + "u;\n"; + op += SP + SP + SP + SP + "std::size_t const channel = elem_idx / spatialSize;\n"; op += SP + SP + SP + SP + "output[elem_idx] = bias[channel];\n"; op += SP + SP + SP + "}\n"; op += SP + SP + "}\n"; @@ -823,20 +826,20 @@ public: if (fShapeX.empty() || fShapeW.empty() || fShapeY.empty()) throw std::runtime_error("SOFIE Conv Op called to Generate without being initialized first"); - size_t bsize = fShapeX[0].dim; - size_t oDepth = (fDim > 2) ? fShapeY[2].dim : 1; - size_t oHeight = (fDim > 1) ? fShapeY[fDim].dim : 1; - size_t oWidth = fShapeY[fDim + 1].dim; - size_t iDepth = (fDim > 2) ? fShapeX[2].dim : 1; - size_t iHeight = (fDim > 1) ? fShapeX[fDim].dim : 1; - size_t iWidth = fShapeX[fDim + 1].dim; + std::string bsize = fShapeX[0].GetVal(); + std::string oDepth = (fDim > 2) ? fShapeY[2].GetVal() : "1"; + std::string oHeight = (fDim > 1) ? fShapeY[fDim].GetVal() : "1"; + std::string oWidth = fShapeY[fDim + 1].GetVal(); + std::string iDepth = (fDim > 2) ? fShapeX[2].GetVal() : "1"; + std::string iHeight = (fDim > 1) ? fShapeX[fDim].GetVal() : "1"; + std::string iWidth = fShapeX[fDim + 1].GetVal(); size_t outChannels = fShapeW[0]; size_t kernelSize = fAttrKernelShape[0] * fAttrKernelShape[1] * fAttrKernelShape[2]; // gemm dimensions computed from shape members size_t gemm_n = outChannels; // output channels size_t gemm_k = fShapeW[1] * kernelSize; // input channels/group * kernel volume - size_t gemm_m = oDepth * oHeight * oWidth; // output spatial size per channel - size_t colElements = gemm_k * gemm_m; // colRows * colCols + std::string gemm_m = "(" + oDepth + " * " + oHeight + " * " + oWidth + ")"; // output spatial size per channel + std::string colElements = std::to_string(gemm_k) + " * " + gemm_m; // colRows * colCols size_t wTotal = ConvertShapeToLength(fShapeW); // For group conv: per-group output channels and _f offset @@ -865,10 +868,10 @@ public: // Step 2: Batch loop // ----------------------------------------------------------------------- out << SP << "for (std::size_t n = 0; n < " << bsize << "; n++) {\n\n"; - out << SP << SP << "std::size_t const x_offset = n * " - << fShapeX[1].dim * iDepth * iHeight * iWidth << "u;\n"; - out << SP << SP << "std::size_t const out_offset = n * " - << fShapeY[1].dim * gemm_m << "u;\n\n"; + out << SP << SP << "std::size_t const x_offset = n * (" + << std::to_string(fShapeX[1].dim) + " * " + iDepth + " * " + iHeight + " * " + iWidth << ");\n"; + out << SP << SP << "std::size_t const out_offset = n * (" + << std::to_string(fShapeY[1].dim) + " * " + gemm_m << ");\n\n"; // ----------------------------------------------------------------------- // Step 3 + 4: Im2Col then GEMM — structure differs for grouped vs non-grouped @@ -883,12 +886,14 @@ public: out << SP << SP << SP << "alpaka::exec(queue, workDiv_im2col, im2colKernel_" << opName << ", alpaka::getPtrNative(deviceBuf_" << fNX << ") + x_offset" << ", alpaka::getPtrNative(deviceBuf_" << imcol << ")" + << ", static_cast(" << oDepth << "), static_cast(" << oHeight << "), static_cast(" << oWidth << ")" + << ", static_cast(" << iDepth << "), static_cast(" << iHeight << "), static_cast(" << iWidth << ")" << ", static_cast(" << colElements << "));\n"; out << SP << SP << SP << "alpaka::wait(queue);\n"; out << SP << SP << "}\n\n"; if (!fNB.empty()) { - size_t biasElements = gemm_n * gemm_m; + std::string biasElements = std::to_string(gemm_n) + " * " + gemm_m; out << SP << SP << "// Step 4a: broadcast bias into output slice\n"; out << SP << SP << "{\n"; out << SP << SP << SP << "auto const elementsPerThread_bias = Vec::all(static_cast(1));\n"; @@ -897,6 +902,7 @@ public: out << SP << SP << SP << "alpaka::exec(queue, workDiv_bias, biasBroadcastKernel_" << opName << ", alpaka::getPtrNative(deviceBuf_" << fNB << ")" << ", alpaka::getPtrNative(deviceBuf_" << fNY << ") + out_offset" + << ", static_cast(" << gemm_m << ")" << ", static_cast(" << biasElements << "));\n"; out << SP << SP << SP << "alpaka::wait(queue);\n"; out << SP << SP << "}\n\n"; @@ -921,10 +927,10 @@ public: // Grouped convolution: im2col and GEMM per group with group-adjusted input pointer. // Each group processes fShapeW[1] input channels starting at g * fShapeW[1]. out << SP << SP << "for (std::size_t g = 0; g < " << fAttrGroup << "; g++) {\n\n"; - out << SP << SP << SP << "std::size_t const g_in_offset = x_offset + g * " - << fShapeW[1] * iDepth * iHeight * iWidth << "u;\n"; - out << SP << SP << SP << "std::size_t const g_out_offset = out_offset + g * " - << gemm_n * gemm_m << "u;\n"; + out << SP << SP << SP << "std::size_t const g_in_offset = x_offset + g * (" + << std::to_string(fShapeW[1]) + " * " + iDepth + " * " + iHeight + " * " + iWidth << ");\n"; + out << SP << SP << SP << "std::size_t const g_out_offset = out_offset + g * (" + << std::to_string(gemm_n) + " * " + gemm_m << ");\n"; out << SP << SP << SP << "std::size_t const f_offset = g * " << groupFOffset << "u;\n\n"; out << SP << SP << SP << "// im2col for group g (reads only this group's input channels)\n"; @@ -935,12 +941,14 @@ public: out << SP << SP << SP << SP << "alpaka::exec(queue, workDiv_im2col, im2colKernel_" << opName << ", alpaka::getPtrNative(deviceBuf_" << fNX << ") + g_in_offset" << ", alpaka::getPtrNative(deviceBuf_" << imcol << ")" + << ", static_cast(" << oDepth << "), static_cast(" << oHeight << "), static_cast(" << oWidth << ")" + << ", static_cast(" << iDepth << "), static_cast(" << iHeight << "), static_cast(" << iWidth << ")" << ", static_cast(" << colElements << "));\n"; out << SP << SP << SP << SP << "alpaka::wait(queue);\n"; out << SP << SP << SP << "}\n\n"; if (!fNB.empty()) { - size_t groupBiasElements = gemm_n * gemm_m; + std::string groupBiasElements = std::to_string(gemm_n) + " * " + gemm_m; out << SP << SP << SP << "// Broadcast group bias\n"; out << SP << SP << SP << "{\n"; out << SP << SP << SP << SP << "auto const elementsPerThread_bias = Vec::all(static_cast(1));\n"; @@ -949,6 +957,7 @@ public: out << SP << SP << SP << SP << "alpaka::exec(queue, workDiv_bias, biasBroadcastKernel_" << opName << ", alpaka::getPtrNative(deviceBuf_" << fNB << ") + g * " << gemm_n << ", alpaka::getPtrNative(deviceBuf_" << fNY << ") + g_out_offset" + << ", static_cast(" << gemm_m << ")" << ", static_cast(" << groupBiasElements << "));\n"; out << SP << SP << SP << SP << "alpaka::wait(queue);\n"; out << SP << SP << SP << "}\n\n"; @@ -979,17 +988,18 @@ public: std::string GetBlasConfig(){ - size_t oDepth_ = (fDim > 2) ? fShapeY[2].dim : 1; - size_t oHeight_ = (fDim > 1) ? fShapeY[fDim].dim : 1; - size_t oWidth_ = fShapeY[fDim + 1].dim; - size_t kSize_ = fAttrKernelShape[0] * fAttrKernelShape[1] * fAttrKernelShape[2]; - size_t gemm_n_ = fShapeW[0]; - size_t gemm_k_ = fShapeW[1] * kSize_; - size_t gemm_m_ = oDepth_ * oHeight_ * oWidth_; - auto lda = std::to_string(gemm_m_); // ld for xcol^T (gemm_m×gemm_k col-major) - auto ldb = std::to_string(gemm_k_); // ld for xf^T (gemm_k×gemm_n col-major) - auto ldc = std::to_string(gemm_m_); // ld for y^T (gemm_m×gemm_n col-major) - return std::to_string(gemm_m_) + ", " + std::to_string(gemm_n_) + ", " + std::to_string(gemm_k_) + ", " + lda + ", " + ldb + ", " + ldc + ", 'n', 'n'"; + // gemm_m (output spatial) can be dynamic, so build the layout dims as runtime expressions. + // addLayoutConfig is emitted in the ctor where the shape params are in scope, so the layout + // registers for the actual runtime size and matmul finds it. Static dims resolve to numbers. + std::string oDepth_ = (fDim > 2) ? fShapeY[2].GetVal() : "1"; + std::string oHeight_ = (fDim > 1) ? fShapeY[fDim].GetVal() : "1"; + std::string oWidth_ = fShapeY[fDim + 1].GetVal(); + size_t kSize_ = fAttrKernelShape[0] * fAttrKernelShape[1] * fAttrKernelShape[2]; + std::string gemm_n_ = std::to_string(fShapeW[0]); + std::string gemm_k_ = std::to_string(fShapeW[1] * kSize_); + std::string gemm_m_ = "(" + oDepth_ + " * " + oHeight_ + " * " + oWidth_ + ")"; + // lda = gemm_m, ldb = gemm_k, ldc = gemm_m (all col-major) + return gemm_m_ + ", " + gemm_n_ + ", " + gemm_k_ + ", " + gemm_m_ + ", " + gemm_k_ + ", " + gemm_m_ + ", 'n', 'n'"; } }; diff --git a/core/inc/SOFIE/ROperator_Gather.hxx b/core/inc/SOFIE/ROperator_Gather.hxx index 3c16f18..413ba2c 100644 --- a/core/inc/SOFIE/ROperator_Gather.hxx +++ b/core/inc/SOFIE/ROperator_Gather.hxx @@ -276,6 +276,18 @@ public: return out.str(); } +// symbolic names used by the kernel index expressions, shared by the kernel +// signature and the launch call +std::vector GetGPUDynParams() const { + std::vector params; + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fShapeY), params); + UTILITY::CollectDimParams(fShapeY, params); + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fShapeIndices), params); + UTILITY::CollectDimParams({fShapeX[fAttrAxis]}, params); + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fShapeX), params); + return params; +} + std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { if (fIsOutputConstant) return ""; opName = "op_" + opName; @@ -301,6 +313,8 @@ std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { op += SP + SP + SP + "T const* __restrict__ input,\n"; op += SP + SP + SP + "int64_t const* __restrict__ indices,\n"; op += SP + SP + SP + "T* __restrict__ output,\n"; + for (auto &p : GetGPUDynParams()) + op += SP + SP + SP + "std::size_t const " + p + ",\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -311,8 +325,8 @@ std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { for (std::size_t d = 0; d < D; ++d) { op += SP + SP + SP + SP + "std::size_t const out_" + std::to_string(d) - + " = (elem_idx / " + stridesY[d].GetVal() + "u) % " - + fShapeY[d].GetVal() + "u;\n"; + + " = (elem_idx / (" + stridesY[d].GetVal() + ")) % (" + + fShapeY[d].GetVal() + ");\n"; } op += "\n"; @@ -325,14 +339,14 @@ std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { for (std::size_t i = 0; i < q; ++i) { op += SP + SP + SP + SP + SP + "out_" + std::to_string(fAttrAxis + i) - + " * " + stridesIndices[i].GetVal() + "u"; + + " * (" + stridesIndices[i].GetVal() + ")"; op += (i + 1 < q) ? " +\n" : ";\n"; } } op += "\n"; op += SP + SP + SP + SP + "int64_t k = indices[i_index];\n"; - op += SP + SP + SP + SP + "if (k < 0) k += " + fShapeX[fAttrAxis].GetVal() + ";\n"; + op += SP + SP + SP + SP + "if (k < 0) k += (" + fShapeX[fAttrAxis].GetVal() + ");\n"; op += SP + SP + SP + SP + "if (k < 0) k = 0;\n"; op += SP + SP + SP + SP + "if (k >= static_cast(" + fShapeX[fAttrAxis].GetVal() + ")) " + "k = static_cast(" + fShapeX[fAttrAxis].GetVal() + ") - 1;\n\n"; @@ -342,15 +356,15 @@ std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { // + sum over j in [axis+1, r): out_{j-1+q} * stridesX[j] // (the dims after axis in Y are shifted by q-1 relative to X) op += SP + SP + SP + SP + "std::size_t const input_idx =\n"; - op += SP + SP + SP + SP + SP + "static_cast(k) * " + stridesX[fAttrAxis].GetVal() + "u"; + op += SP + SP + SP + SP + SP + "static_cast(k) * (" + stridesX[fAttrAxis].GetVal() + ")"; for (std::size_t j = 0; j < static_cast(fAttrAxis); ++j) { op += " +\n" + SP + SP + SP + SP + SP - + "out_" + std::to_string(j) + " * " + stridesX[j].GetVal() + "u"; + + "out_" + std::to_string(j) + " * (" + stridesX[j].GetVal() + ")"; } for (std::size_t j = fAttrAxis + 1; j < r; ++j) { // in Y, the coord for X's dim j lives at output dim q + j - 1 op += " +\n" + SP + SP + SP + SP + SP - + "out_" + std::to_string(q + j - 1) + " * " + stridesX[j].GetVal() + "u"; + + "out_" + std::to_string(q + j - 1) + " * (" + stridesX[j].GetVal() + ")"; } op += ";\n\n"; @@ -387,8 +401,10 @@ std::string Generate_GPU_ALPAKA(std::string opName) override { << ", " << kname << ", alpaka::getPtrNative(deviceBuf_" << fNX << ")" << ", alpaka::getPtrNative(deviceBuf_" << fNIndices << ")" - << ", alpaka::getPtrNative(deviceBuf_" << fNY << ")" - << ", static_cast(" << totalElements << "));\n"; + << ", alpaka::getPtrNative(deviceBuf_" << fNY << ")"; + for (auto &p : GetGPUDynParams()) + out << ", static_cast(" << p << ")"; + out << ", static_cast(" << totalElements << "));\n"; out << SP << "alpaka::enqueue(queue, task_" << opName << ");\n"; return out.str(); } diff --git a/core/inc/SOFIE/ROperator_Reduce.hxx b/core/inc/SOFIE/ROperator_Reduce.hxx index f3e7170..8f1ee15 100644 --- a/core/inc/SOFIE/ROperator_Reduce.hxx +++ b/core/inc/SOFIE/ROperator_Reduce.hxx @@ -28,9 +28,9 @@ private: std::string fNX; std::string fNAxes; std::string fNY; - std::vector fShapeX; - std::vector fShapeY; - std::vector fShapeYNotPruned; // needed for fKeepdims=0 + std::vector fShapeX; + std::vector fShapeY; + std::vector fShapeYNotPruned; // needed for fKeepdims=0 public: @@ -82,7 +82,6 @@ public: // set to 1 the reduced dims outputShape[fAttrAxes[j]] = 1; } - fShapeYNotPruned = outputShape; // in case of pruning dimension we need to sort axes attributes if (fkeepdims == 0) { auto ax = fAttrAxes; @@ -98,6 +97,32 @@ public: } return ret; } + + // dynamic (Dim) shape inference, mirrors the size_t version above + std::vector DoShapeInference(const std::vector & input) { + auto ret = input; + auto & outputShape = ret; + for (size_t j = 0; j < fAttrAxes.size(); j++) { + if (fAttrAxes[j] < 0) fAttrAxes[j] += outputShape.size(); + if (fAttrAxes[j] < 0 || (size_t) fAttrAxes[j] >= outputShape.size()) + throw std::runtime_error("SOFIE Reduce Op - invalid axes values " + std::to_string(fAttrAxes[j])); + outputShape[fAttrAxes[j]] = Dim{1}; + } + fShapeYNotPruned = outputShape; + if (fkeepdims == 0) { + auto ax = fAttrAxes; + std::sort(ax.begin(), ax.end()); + for (size_t j = 0; j < ax.size(); j++) { + if (outputShape.size() > 0) { + outputShape.erase(outputShape.begin() + ax[j]); + for (size_t k = j+1; k < ax.size(); k++) + ax[k] -= 1; + } + } + } + return ret; + } + void Initialize(RModel& model) override { fUseSession = model.UseSession(); @@ -106,7 +131,7 @@ public: // input must be a graph input, or already initialized intermediate tensor throw std::runtime_error("SOFIE Reduce Op Input Tensor " + fNX + " is not found in model"); } - fShapeX = model.GetTensorShape(fNX); + fShapeX = model.GetDimTensorShape(fNX); // check if tensor with axes is provided if (!fNAxes.empty()) { auto ax_shptr = model.GetInitializedTensorData(fNAxes); @@ -121,10 +146,10 @@ public: fAttrAxes[i] = i; } // find shape of Y and add it in the list of intermediate tensors - fShapeY = ShapeInference({fShapeX})[0]; + fShapeY = DoShapeInference(fShapeX); model.AddIntermediateTensor(fNY, model.GetTensorType(fNX), fShapeY); if (model.Verbose()){ - std::cout << Name() << " : " << fNX << " -> " << fNY << " shape " << ConvertShapeToString(fShapeY) << std::endl; + std::cout << Name() << " : " << fNX << " -> " << fNY << " shape " << ConvertDimShapeToString(fShapeY) << std::endl; } model.AddNeededStdLib("algorithm"); } @@ -135,8 +160,8 @@ public: throw std::runtime_error("SOFIE Reduce Op called to Generate without being initialized first"); } - size_t inputLength = SOFIE::ConvertShapeToLength(fShapeX); - size_t outputLength = SOFIE::ConvertShapeToLength(fShapeY); + std::string inputLength = SOFIE::ConvertDimShapeToLength(fShapeX); + std::string outputLength = SOFIE::ConvertDimShapeToLength(fShapeY); auto inputStrides = SOFIE::UTILITY::ComputeStrideFromShape(fShapeX); // output stride (or not pruned vector) @@ -174,7 +199,7 @@ public: } } } - size_t reducedLength = inputLength / outputLength; + std::string reducedLength = "((" + inputLength + ") / (" + outputLength + "))"; if (reduceDims == kLast) { //std::cout << "reduction for operator " << opName << " is last" << std::endl; // new faster implementation using a single loop @@ -267,8 +292,8 @@ public: for (size_t k = 0; k < dim; k++) { if (std::find(fAttrAxes.begin(), fAttrAxes.end(), k) == fAttrAxes.end()) { // do for not reducing axes - out << SP << SP << "size_t i_" << k << " = i / " << inputStrides[k] << " % " << fShapeX[k] << ";\n"; - out << SP << SP << "outputIndex += i_" << k << " * " << outputStrides[k] << ";\n"; + out << SP << SP << "size_t i_" << k << " = i / (" << inputStrides[k].GetVal() << ") % (" << fShapeX[k].GetVal() << ");\n"; + out << SP << SP << "outputIndex += i_" << k << " * (" << outputStrides[k].GetVal() << ");\n"; } } // now compute reduction @@ -313,9 +338,6 @@ public: const std::size_t Dx = fShapeX.size(); auto inputStrides = UTILITY::ComputeStrideFromShape(fShapeX); auto outputStrides = UTILITY::ComputeStrideFromShape(fShapeYNotPruned); - std::size_t inputLength = ConvertShapeToLength(fShapeX); - std::size_t outputLength = ConvertShapeToLength(fShapeY); - std::size_t reducedLength = inputLength / outputLength; // Partition axes into keep (non-reduced) and reduce sets. std::vector redAxes, keepAxes; @@ -326,12 +348,21 @@ public: keepAxes.push_back(d); } - // Row-major strides for decomposing the flat reduction index r into - // per-axis coordinates. + // Row-major strides (as runtime expressions) for decomposing the flat + // reduction index r into per-axis coordinates. // redStrides[i] = product of fShapeX[redAxes[j]] for j > i - std::vector redStrides(redAxes.size(), 1); + std::vector redStrides(redAxes.size(), "1"); for (int ri = (int)redAxes.size() - 2; ri >= 0; --ri) - redStrides[ri] = redStrides[ri + 1] * fShapeX[redAxes[ri + 1]]; + redStrides[ri] = "(" + redStrides[ri + 1] + " * " + fShapeX[redAxes[ri + 1]].GetVal() + ")"; + + // dynamic shape params referenced by the baked index math, passed as size_t kernel args + std::vector dynParams; + for (auto &d : fShapeX) + if (d.isParam) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } std::string kname = "ReduceKernel_" + Name() + "_" + fNY; @@ -344,7 +375,10 @@ public: op += SP + SP + SP + "T const* __restrict__ input,\n"; op += SP + SP + SP + "T* __restrict__ output,\n"; op += SP + SP + SP + "std::size_t const reducedLength,\n"; - op += SP + SP + SP + "std::size_t const outputLength) const {\n\n"; + op += SP + SP + SP + "std::size_t const outputLength"; + for (auto &p : dynParams) + op += ",\n" + SP + SP + SP + "std::size_t const " + p; + op += ") const {\n\n"; // ---- shared memory (fixed 256 slots, matches block size) ---- op += SP + SP + SP + "auto& shmem = alpaka::declareSharedVar(acc);\n\n"; @@ -359,8 +393,8 @@ public: for (std::size_t d = 0; d < Dx; ++d) { if (std::find(redAxes.begin(), redAxes.end(), d) == redAxes.end()) { op += SP + SP + SP + "std::size_t const oy_" + std::to_string(d) - + " = (out_idx / " + std::to_string(outputStrides[d]) + "u) % " - + std::to_string(fShapeYNotPruned[d]) + "u;\n"; + + " = (out_idx / (" + outputStrides[d].GetVal() + ")) % (" + + fShapeYNotPruned[d].GetVal() + ");\n"; } } op += "\n"; @@ -377,8 +411,8 @@ public: for (std::size_t ri = 0; ri < redAxes.size(); ++ri) { std::size_t rd = redAxes[ri]; op += SP + SP + SP + SP + "std::size_t const r_" + std::to_string(rd) - + " = (r / " + std::to_string(redStrides[ri]) + "u) % " - + std::to_string(fShapeX[rd]) + "u;\n"; + + " = (r / (" + redStrides[ri] + ")) % (" + + fShapeX[rd].GetVal() + ");\n"; } // Compute flat input index. @@ -386,7 +420,7 @@ public: for (std::size_t d = 0; d < Dx; ++d) { bool isReduced = std::find(redAxes.begin(), redAxes.end(), d) != redAxes.end(); std::string coord = isReduced ? "r_" + std::to_string(d) : "oy_" + std::to_string(d); - op += SP + SP + SP + SP + SP + coord + " * " + std::to_string(inputStrides[d]) + "u"; + op += SP + SP + SP + SP + SP + coord + " * (" + inputStrides[d].GetVal() + ")"; op += (d + 1 < Dx) ? " +\n" : ";\n"; } @@ -423,7 +457,7 @@ public: op += SP + SP + SP + "if (thread_id == 0u) {\n"; op += SP + SP + SP + SP + "T result = shmem[0];\n"; if (Op == ReduceMean) - op += SP + SP + SP + SP + "result /= static_cast(" + std::to_string(reducedLength) + "u);\n"; + op += SP + SP + SP + SP + "result /= static_cast(reducedLength);\n"; else if (Op == ReduceL2) op += SP + SP + SP + SP + "result = std::sqrt(result);\n"; op += SP + SP + SP + SP + "output[out_idx] = result;\n"; @@ -443,25 +477,36 @@ public: if (fShapeX.empty() || fShapeY.empty()) throw std::runtime_error("SOFIE Reduce Op called to Generate without being initialized first"); - std::size_t inputLength = ConvertShapeToLength(fShapeX); - std::size_t outputLength = ConvertShapeToLength(fShapeY); - std::size_t reducedLength = inputLength / outputLength; + std::string inputLength = ConvertDimShapeToLength(fShapeX); + std::string outputLength = ConvertDimShapeToLength(fShapeY); + std::string reducedLength = "(" + inputLength + ") / (" + outputLength + ")"; std::string kname = "reduceKernel_" + Name() + "_" + fNY; + std::vector dynParams; + for (auto &d : fShapeX) + if (d.isParam) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } + std::string dynArgs; + for (auto &p : dynParams) dynArgs += ", static_cast(" + p + ")"; + std::stringstream out; out << "\n//------ " << Name() << "_GPU_ALPAKA\n"; // Grid: one block per output element; Block: 256 threads cooperate to // reduce the corresponding slice. out << SP << "alpaka::WorkDivMembers workDiv_" << fNY << "(\n"; - out << SP << SP << "Vec::all(Idx{" << outputLength << "u}),\n"; + out << SP << SP << "Vec::all(Idx{" << outputLength << "}),\n"; out << SP << SP << "Vec::all(Idx{256u}),\n"; out << SP << SP << "Vec::all(Idx{1u}));\n"; out << SP << "alpaka::exec(queue, workDiv_" << fNY << ", " << kname << ", alpaka::getPtrNative(deviceBuf_" << fNX << ")" << ", alpaka::getPtrNative(deviceBuf_" << fNY << ")" - << ", static_cast(" << reducedLength << "u)" - << ", static_cast(" << outputLength << "u));\n"; + << ", static_cast(" << reducedLength << ")" + << ", static_cast(" << outputLength << ")" + << dynArgs << ");\n"; return out.str(); } diff --git a/core/inc/SOFIE/ROperator_Slice.hxx b/core/inc/SOFIE/ROperator_Slice.hxx index fb738cf..cd20685 100644 --- a/core/inc/SOFIE/ROperator_Slice.hxx +++ b/core/inc/SOFIE/ROperator_Slice.hxx @@ -494,6 +494,18 @@ public: return out.str(); } + // symbolic names used by the kernel index expressions, shared by the + // kernel signature and the launch call + std::vector GetGPUDynParams() const { + std::vector params; + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fShapeOutput), params); + UTILITY::CollectDimParams(fShapeOutput, params); + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(fShapeInput), params); + UTILITY::CollectDimParams(fStart, params); + UTILITY::CollectDimParams(fSteps, params); + return params; + } + std::string Generate_GPU_Kernel_ALPAKA(std::string opName) override { if (fIsOutputConstant) return ""; opName = "op_" + opName; @@ -505,7 +517,6 @@ public: auto inputStrides = UTILITY::ComputeStrideFromShape(fShapeInput); auto outputStrides = UTILITY::ComputeStrideFromShape(fShapeOutput); - std::size_t totalElements = ConvertShapeToLength(fShapeOutput); std::string kname = "SliceKernel_" + opName; std::string op; @@ -516,6 +527,8 @@ public: op += SP + SP + SP + "TAcc const& acc,\n"; op += SP + SP + SP + "T const* __restrict__ input,\n"; op += SP + SP + SP + "T* __restrict__ output,\n"; + for (auto &p : GetGPUDynParams()) + op += SP + SP + SP + "std::size_t const " + p + ",\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -526,8 +539,8 @@ public: for (std::size_t d = 0; d < D; ++d) { op += SP + SP + SP + SP + "std::size_t const out_" + std::to_string(d) - + " = (elem_idx / " + outputStrides[d].GetVal() + "u) % " - + fShapeOutput[d].GetVal() + "u;\n"; + + " = (elem_idx / (" + outputStrides[d].GetVal() + ")) % (" + + fShapeOutput[d].GetVal() + ");\n"; } op += "\n"; @@ -538,12 +551,12 @@ public: op += SP + SP + SP + SP + "std::size_t const input_idx =\n"; for (std::size_t d = 0; d < D; ++d) { // input coordinate for this dim: start + out_d * step - std::string input_coord = "(" + fStart[d].GetVal() - + " + out_" + std::to_string(d) - + " * " + fSteps[d].GetVal() + ")"; + std::string input_coord = "((" + fStart[d].GetVal() + + ") + out_" + std::to_string(d) + + " * (" + fSteps[d].GetVal() + "))"; op += SP + SP + SP + SP + SP + "static_cast(" + input_coord + ")" - + " * " + inputStrides[d].GetVal() + "u"; + + " * (" + inputStrides[d].GetVal() + ")"; op += (d + 1 < D) ? " +\n" : ";\n\n"; } @@ -567,7 +580,7 @@ public: if (fShapeInput.empty() || fShapeOutput.empty()) throw std::runtime_error("SOFIE Slice Op called to Generate without being initialized first"); - std::size_t totalElements = ConvertShapeToLength(fShapeOutput); + auto totalElements = ConvertDimShapeToLength(fShapeOutput); std::string kname = "sliceKernel_" + opName; std::stringstream out; @@ -578,8 +591,10 @@ public: out << SP << "alpaka::exec(queue, workDiv_" << opName << ", " << kname << ", alpaka::getPtrNative(deviceBuf_" << fNData << ")" - << ", alpaka::getPtrNative(deviceBuf_" << fNOutput << ")" - << ", static_cast(" << totalElements << "));\n"; + << ", alpaka::getPtrNative(deviceBuf_" << fNOutput << ")"; + for (auto &p : GetGPUDynParams()) + out << ", static_cast(" << p << ")"; + out << ", static_cast(" << totalElements << "));\n"; return out.str(); } diff --git a/core/inc/SOFIE/ROperator_Tile.hxx b/core/inc/SOFIE/ROperator_Tile.hxx index 5a3921e..fda5de4 100644 --- a/core/inc/SOFIE/ROperator_Tile.hxx +++ b/core/inc/SOFIE/ROperator_Tile.hxx @@ -19,8 +19,8 @@ private: std::string fNRepeats; std::string fNInput; std::string fNY; - std::vector fShapeInput; - std::vector fShapeY; + std::vector fShapeInput; + std::vector fShapeY; std::vector fRepeats; public: @@ -50,7 +50,7 @@ public: if (model.CheckIfTensorAlreadyExist(fNRepeats) == false) throw std::runtime_error("SOFIE Tile Op Repeats Tensor is not found in model"); - fShapeInput = model.GetTensorShape(fNInput); + fShapeInput = model.GetDimTensorShape(fNInput); if (!model.IsInitializedTensor(fNRepeats)) throw std::runtime_error("SOFIE Tile Op: non-initialized repeats input is not supported"); @@ -75,13 +75,19 @@ public: if (fRepeats.size()){ model.RemoveInitializedTensor(fNRepeats); } - fShapeY = ShapeInference({fShapeInput, fRepeats})[0]; + fShapeY = fShapeInput; + for (size_t i = 0; i < fRepeats.size(); i++) { + if (fShapeInput[i].isParam) + fShapeY[i] = Dim{fShapeInput[i].GetVal() + " * " + std::to_string(fRepeats[i]), static_cast(-1)}; + else + fShapeY[i] = Dim{fShapeInput[i].dim * fRepeats[i]}; + } model.AddIntermediateTensor(fNY, model.GetTensorType(fNInput), fShapeY); if (model.Verbose()) - std::cout << "Tile: " << fNInput << " " << ConvertShapeToString(fShapeInput) - << " -> " << fNY << " with shape " << ConvertShapeToString(fShapeY) + std::cout << "Tile: " << fNInput << " " << ConvertDimShapeToString(fShapeInput) + << " -> " << fNY << " with shape " << ConvertDimShapeToString(fShapeY) << " given repeats " << ConvertShapeToString(fRepeats) << std::endl; } @@ -101,11 +107,11 @@ public: out << SP << "const int input_shape[" << fShapeInput.size() << "] = {"; for (size_t i = 0; i < fShapeInput.size(); ++i) { if (i > 0) out << ", "; - out << fShapeInput[i]; + out << fShapeInput[i].GetVal(); } out << "};\n"; - out << SP << "int inputLength = " << ConvertShapeToLength(fShapeInput) << ";\n"; + out << SP << "int inputLength = " << ConvertDimShapeToLength(fShapeInput) << ";\n"; out << SP << "int s = 1;\n"; // Read repeats from the tensor at runtime so the generated code remains @@ -152,13 +158,26 @@ public: auto inputStrides = UTILITY::ComputeStrideFromShape(fShapeInput); auto outputStrides = UTILITY::ComputeStrideFromShape(fShapeY); - std::size_t totalElements = ConvertShapeToLength(fShapeY); // If fRepeats is populated, repeats were known at generation time and // we can bake fShapeInput[d] as literals — no runtime repeats pointer needed. // If fRepeats is empty (future: runtime repeats), pass repeats as a kernel arg. bool repeatsKnown = !fRepeats.empty(); + auto isIdent = [](const std::string &s){ + if (s.empty() || (s[0] >= '0' && s[0] <= '9')) return false; + for (char c : s) + if (!((c>='a'&&c<='z')||(c>='A'&&c<='Z')||(c>='0'&&c<='9')||c=='_')) return false; + return true; + }; + std::vector dynParams; + for (auto &d : fShapeInput) + if (d.isParam && isIdent(d.param)) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } + std::string kname = "TileKernel_" + opName; std::string op; @@ -171,6 +190,8 @@ public: op += SP + SP + SP + "T* __restrict__ output,\n"; if (!repeatsKnown) op += SP + SP + SP + "int64_t const* __restrict__ repeats,\n"; + for (auto &p : dynParams) + op += SP + SP + SP + "std::size_t const " + p + ",\n"; op += SP + SP + SP + "std::size_t const totalElements) const {\n\n"; op += SP + SP + SP + "auto const global_thread_idx = alpaka::getIdx(acc)[0];\n"; @@ -182,8 +203,8 @@ public: // Decompose output linear index — output strides always compile-time for (std::size_t d = 0; d < D; ++d) { op += SP + SP + SP + SP + "std::size_t const out_" + std::to_string(d) - + " = (elem_idx / " + std::to_string(outputStrides[d]) + "u) % " - + std::to_string(fShapeY[d]) + "u;\n"; + + " = (elem_idx / (" + outputStrides[d].GetVal() + ")) % (" + + fShapeY[d].GetVal() + ");\n"; } op += "\n"; @@ -195,8 +216,8 @@ public: op += SP + SP + SP + SP + "std::size_t const input_idx =\n"; for (std::size_t d = 0; d < D; ++d) { op += SP + SP + SP + SP + SP - + "(out_" + std::to_string(d) + " % " + std::to_string(fShapeInput[d]) + "u)" - + " * " + std::to_string(inputStrides[d]) + "u"; + + "(out_" + std::to_string(d) + " % (" + fShapeInput[d].GetVal() + "))" + + " * (" + inputStrides[d].GetVal() + ")"; op += (d + 1 < D) ? " +\n" : ";\n\n"; } @@ -220,16 +241,32 @@ public: throw std::runtime_error("SOFIE Operator Tile called to Generate without being initialized first"); bool repeatsKnown = !fRepeats.empty(); - std::size_t totalElements = ConvertShapeToLength(fShapeY); + std::string totalElements = ConvertDimShapeToLength(fShapeY); std::string kname = "tileKernel_" + opName; + auto isIdent = [](const std::string &s){ + if (s.empty() || (s[0] >= '0' && s[0] <= '9')) return false; + for (char c : s) + if (!((c>='a'&&c<='z')||(c>='A'&&c<='Z')||(c>='0'&&c<='9')||c=='_')) return false; + return true; + }; + std::vector dynParams; + for (auto &d : fShapeInput) + if (d.isParam && isIdent(d.param)) { + bool seen = false; + for (auto &q : dynParams) if (q == d.param) seen = true; + if (!seen) dynParams.push_back(d.param); + } + // Build argument list once, reused for both getValidWorkDiv and exec std::string args = "alpaka::getPtrNative(deviceBuf_" + fNInput + "), " + "alpaka::getPtrNative(deviceBuf_" + fNY + ")"; if (!repeatsKnown) args += ", alpaka::getPtrNative(deviceBuf_" + fNRepeats + ")"; - args += ", static_cast(" + std::to_string(totalElements) + ")"; + for (auto &p : dynParams) + args += ", static_cast(" + p + ")"; + args += ", static_cast(" + totalElements + ")"; std::stringstream out; out << "\n//------ TILE_GPU_ALPAKA\n"; diff --git a/core/inc/SOFIE/ROperator_Transpose.hxx b/core/inc/SOFIE/ROperator_Transpose.hxx index 03dad41..6a6fe1a 100644 --- a/core/inc/SOFIE/ROperator_Transpose.hxx +++ b/core/inc/SOFIE/ROperator_Transpose.hxx @@ -173,6 +173,17 @@ public: return out.str(); } + // symbolic names used by the kernel stride expressions, shared by the kernel signature and the launch call + + std::vector GetGPUDynParams() const { + auto dimShapeData = fDimShapeData.empty() ? ConvertShapeToDim(fShapeData) : fDimShapeData; + auto dimShapeOutput = fDimShapeOutput.empty() ? ConvertShapeToDim(fShapeOutput) : fDimShapeOutput; + std::vector params; + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(dimShapeOutput), params); + UTILITY::CollectDimParams(UTILITY::ComputeStrideFromShape(dimShapeData), params); + return params; + } + std::string Generate_GPU_Kernel_ALPAKA(std::string OpName) { std::string op; OpName = "op_" + OpName; @@ -180,6 +191,8 @@ public: op += SP + "struct TransposeKernel_" + OpName + " {\n"; op += SP + SP + "template\n"; op += SP + SP + "ALPAKA_FN_ACC void operator()(TAcc const& acc, T const* input, T* output,"; + for (auto &p : GetGPUDynParams()) + op += "const std::size_t " + p + ","; op += "const std::size_t totalElements) const {\n"; op += SP + SP + SP + SP + "auto const idx = alpaka::getIdx(acc)[0];\n"; op += SP + SP + SP + SP + "if(idx >= totalElements) return;\n"; @@ -193,12 +206,12 @@ public: auto outputStrides = UTILITY::ComputeStrideFromShape(dimShapeOutput); for (size_t k = 0; k < dimShapeData.size(); k++) { - op += SP + SP + SP + SP + "coord = remaining / " - + outputStrides[k].GetVal() + "u;\n"; - op += SP + SP + SP + SP + "remaining = remaining - coord * " - + outputStrides[k].GetVal() + "u;\n"; - op += SP + SP + SP + SP + "input_idx += coord * " - + inputStrides[fAttrPerm[k]].GetVal() + "u;\n"; + op += SP + SP + SP + SP + "coord = remaining / (" + + outputStrides[k].GetVal() + ");\n"; + op += SP + SP + SP + SP + "remaining = remaining - coord * (" + + outputStrides[k].GetVal() + ");\n"; + op += SP + SP + SP + SP + "input_idx += coord * (" + + inputStrides[fAttrPerm[k]].GetVal() + ");\n"; } op += SP + SP + SP + SP + "output[idx] = input[input_idx];\n"; @@ -226,7 +239,10 @@ public: out << SP << "auto const workDiv_" << fNOutput << " = sofie_workdiv(elementsPerGrid_" << fNOutput << ");\n"; out << SP << "auto task_" << OpName << " = alpaka::createTaskKernel(workDiv_" << fNOutput << ", transposeKernel_" << OpName << ", alpaka::getPtrNative(deviceBuf_" << fNData - << "), alpaka::getPtrNative(deviceBuf_" << fNOutput << "), static_cast(" << length << "));\n"; + << "), alpaka::getPtrNative(deviceBuf_" << fNOutput << ")"; + for (auto &p : GetGPUDynParams()) + out << ", static_cast(" << p << ")"; + out << ", static_cast(" << length << "));\n"; out << SP <<"alpaka::enqueue(queue, task_" << OpName << ");\n"; return out.str(); } diff --git a/core/inc/SOFIE/SOFIE_common.hxx b/core/inc/SOFIE/SOFIE_common.hxx index e36df0a..5a77e4f 100644 --- a/core/inc/SOFIE/SOFIE_common.hxx +++ b/core/inc/SOFIE/SOFIE_common.hxx @@ -547,6 +547,9 @@ void UnidirectionalBroadcast(const T* data, const std::vector& shape, co std::vector ComputeStrideFromShape(const std::vector & shape); std::vector ComputeStrideFromShape(const std::vector & shape); +/// collect the symbolic identifiers used in the dims, deduplicated in first-seen order +void CollectDimParams(const std::vector & dims, std::vector & params); + /// function to check if a >> 0 and a < MAX using a single comparison //// use trick casting to unsigned values so it becomes a single comparison inline bool is_a_ge_zero_and_a_lt_b(int a, int b) { diff --git a/core/src/RModel.cxx b/core/src/RModel.cxx index 377171c..1e569e6 100644 --- a/core/src/RModel.cxx +++ b/core/src/RModel.cxx @@ -648,12 +648,25 @@ void RModel::Initialize(const std::map & inputParams, bool // Build set of initialized tensors consumed by at least one runtime operator (need for later) std::unordered_set runtimeInitializedInputs; + std::vector initFailures; // diagnostic: collect every Initialize failure in one pass for(size_t op_idx = 0; op_idx < fOperators.size(); ++op_idx){ if (verbose) { auto& r = *fOperators[op_idx].get(); std::cout << "Initializing operator " << i << " " << typeid(r).name() << std::endl; } - fOperators[op_idx]->Initialize(*this); + try { + fOperators[op_idx]->Initialize(*this); + } catch (const std::exception &e) { + std::string msg = "op#" + std::to_string(op_idx) + " kind=" + + std::to_string(static_cast(fOperators[op_idx]->GetKind())) + " in["; + for (auto &t : fOperators[op_idx]->GetOpInputTensors()) msg += std::string{t} + " "; + msg += "] out["; + for (auto &t : fOperators[op_idx]->GetOpOutputTensors()) msg += std::string{t} + " "; + msg += "] : " + std::string(e.what()); + std::cerr << "[SOFIE][init-fail] " << msg << std::endl; + initFailures.push_back(msg); + continue; // keep going so one run reports all failures, not just the first + } for(auto &it:fOperators[op_idx]->GetOpOutputTensors()){ std::string name = std::string{it}; // check if tensor is not an initialized or output tensor and it is not already in the list @@ -677,6 +690,16 @@ void RModel::Initialize(const std::map & inputParams, bool i++; } + if (!initFailures.empty()) { + std::cerr << "\n[SOFIE] ===== " << initFailures.size() + << " operator(s) failed to initialize =====\n"; + for (auto &f : initFailures) std::cerr << " " << f << "\n"; + std::cerr << "[SOFIE] note: \"not found\" errors are usually cascades from an earlier failed op; " + "the real gaps are the \"dynamic\"/\"not supported\" ones\n" << std::endl; + throw std::runtime_error("SOFIE: " + std::to_string(initFailures.size()) + + " operator(s) failed to initialize (see list above)"); + } + // loop on initialized tensors and make the integers as constant to be // not written in a weight file and check if the tensors flagged as not writable are really not writable, // i.e. are not used by non constant operators diff --git a/core/src/RModel_ALPAKA.cxx b/core/src/RModel_ALPAKA.cxx index 9e0e84c..d50de5b 100644 --- a/core/src/RModel_ALPAKA.cxx +++ b/core/src/RModel_ALPAKA.cxx @@ -52,9 +52,9 @@ void RModel::ComputeEltwiseFusionGroups() { auto firstInputs = fOperators[i]->GetOpInputTensors(); group.inputTensor = firstInputs.empty() ? "" : std::string(firstInputs[0]); - // Extend chain: only if CURRENT op is elementwise and its single output can be fused + // Extend chain: only if CURRENT op is elementwise with runtime (non constant) output that can be fused size_t current = i; - while (fOperators[current]->IsElementwise()) { + while (fOperators[current]->IsElementwise() && !fOperators[current]->IsOutputConstant()) { auto curOutputs = fOperators[current]->GetOpOutputTensors(); if (curOutputs.size() != 1) break; std::string curOut = std::string(curOutputs[0]); @@ -64,7 +64,7 @@ void RModel::ComputeEltwiseFusionGroups() { // Must be strictly the next op in sequence and itself elementwise with single input if (nextIdx != current + 1) break; if (opAssigned[nextIdx]) break; - if (!fOperators[nextIdx]->IsElementwise()) break; + if (!fOperators[nextIdx]->IsElementwise() || fOperators[nextIdx]->IsOutputConstant()) break; auto nextInputs = fOperators[nextIdx]->GetOpInputTensors(); if (nextInputs.size() != 1) break; @@ -77,11 +77,19 @@ void RModel::ComputeEltwiseFusionGroups() { auto lastOutputs = fOperators[current]->GetOpOutputTensors(); group.outputTensor = lastOutputs.empty() ? "" : std::string(lastOutputs[0]); - // Element count from intermediate tensor info (all op outputs are intermediates) + // Element count: literal for static tensors, runtime expression for dynamic ones if (!group.outputTensor.empty()) { auto it = fIntermediateTensorInfos.find(group.outputTensor); - if (it != fIntermediateTensorInfos.end()) - group.numElements = ConvertShapeToLength(it->second.shape); + if (it != fIntermediateTensorInfos.end()) { + group.lengthExpr = std::to_string(ConvertShapeToLength(it->second.shape)); + } else { + auto itDyn = fDynamicTensorInfos.find(group.outputTensor); + if (itDyn != fDynamicTensorInfos.end()) + group.lengthExpr = ConvertDimShapeToLength(itDyn->second.shape); + else if (group.isFused()) + throw std::runtime_error("SOFIE eltwise fusion: output tensor " + group.outputTensor + + " not found in intermediate or dynamic tensor infos"); + } } size_t gIdx = fEltwiseFusionGroups.size(); @@ -248,17 +256,18 @@ void RModel::GenerateGPU_ALPAKA_Buffers() { // add also the dynamic tensors (only declarations, allocation will be done later) if (!fDynamicTensorInfos.empty()) { fGC += "//--- declare the dynamic tensors\n"; - fGC += "using bufDev_float = alpaka::Buf, size_t>;\n"; - fGC += "using bufDev_double = alpaka::Buf, size_t>;\n"; - fGC += "using bufDev_int64 = alpaka::Buf, size_t>;\n"; - for (auto &i : fDynamicTensorInfos) { + if (fFusionIntermediateTensors.count(i.first)) continue; if (i.second.type == ETensorType::FLOAT) { - fGC += "bufDev_float bufDev_" + i.first + ";\n"; + fGC += "BufF1D deviceBuf_" + i.first + " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{1}));\n"; } else if (i.second.type == ETensorType::DOUBLE) { - fGC += "bufDev_double bufDev_" + i.first + ";\n"; + fGC += "BufD1D deviceBuf_" + i.first + " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{1}));\n"; + } else if (i.second.type == ETensorType::INT32) { + fGC += "BufI321D deviceBuf_" + i.first + " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{1}));\n"; } else if (i.second.type == ETensorType::INT64) { - fGC += "bufDev_int64 bufDev_" + i.first + ";\n"; + fGC += "BufI641D deviceBuf_" + i.first + " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{1}));\n"; + } else if (i.second.type == ETensorType::BOOL) { + fGC += "BufUI81D deviceBuf_" + i.first + " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{1}));\n"; } } } @@ -269,10 +278,20 @@ void RModel::GenerateDynamicTensorInfo_GPU_ALPAKA() { std::stringstream out; for (auto &i : fDynamicTensorInfos) { + if (fFusionIntermediateTensors.count(i.first)) continue; auto length = ConvertDimShapeToLength(i.second.shape); out << SP << "if (" << length << " > 0) {\n"; - out << "auto bufDev_" + i.first + - " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{" << length << "}));\n"; + if (i.second.type == ETensorType::FLOAT) { + out << SP << "deviceBuf_" << i.first << " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{" << length << "}));\n"; + } else if (i.second.type == ETensorType::DOUBLE) { + out << SP << "deviceBuf_" << i.first << " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{" << length << "}));\n"; + } else if (i.second.type == ETensorType::INT32) { + out << SP << "deviceBuf_" << i.first << " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{" << length << "}));\n"; + } else if (i.second.type == ETensorType::INT64) { + out << SP << "deviceBuf_" << i.first << " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{" << length << "}));\n"; + } else if (i.second.type == ETensorType::BOOL) { + out << SP << "deviceBuf_" << i.first << " = alpaka::allocBuf(devAcc, Ext1D::all(Idx{" << length << "}));\n"; + } out << SP << "}\n"; } fGC += out.str(); @@ -432,11 +451,11 @@ void RModel::GenerateOutput_GPU_ALPAKA() { fusedCode += "\n//------ FUSED_ELTWISE_GPU_ALPAKA" + sfx + "\n"; fusedCode += SP + "{\n"; fusedCode += SP + SP + "auto const elementsPerThread_fused" + sfx + " = Vec::all(static_cast(1));\n"; - fusedCode += SP + SP + "auto const elementsPerGrid_fused" + sfx + " = Vec::all(Idx{" + std::to_string(grp.numElements) + "});\n"; + fusedCode += SP + SP + "auto const elementsPerGrid_fused" + sfx + " = Vec::all(Idx{" + grp.lengthExpr + "});\n"; fusedCode += SP + SP + "auto const workDiv_fused" + sfx + " = sofie_workdiv(elementsPerGrid_fused" + sfx + ");\n"; fusedCode += SP + SP + "auto task_fused" + sfx + " = alpaka::createTaskKernel(workDiv_fused" + sfx + ", " + kname + ", alpaka::getPtrNative(deviceBuf_" + grp.inputTensor + "), alpaka::getPtrNative(deviceBuf_" + grp.outputTensor + - "), static_cast(" + std::to_string(grp.numElements) + "));\n"; + "), static_cast(" + grp.lengthExpr + "));\n"; fusedCode += SP + SP + "alpaka::enqueue(queue, task_fused" + sfx + ");\n"; fusedCode += SP + "}\n"; if (fProfile) { @@ -714,10 +733,22 @@ void RModel::GenerateSessionCode_GPU_ALPAKA() { } if (!fShapeParams.empty()) { - for (auto &p : fShapeParams) { - fGC += ",\n"; - fGC += " size_t " + p.first + " = " + p.second; + // emit params in declaration order (like the infer signature), not unordered_map order, + // so ctor and infer agree on arg order for multi-symbol models + std::unordered_map seenParam; + for (auto &name : fInputTensorNames) { + if (IsDimInputTensor(name)) { + for (auto &d : GetDynamicTensorShape(name)) { + if (d.isParam && seenParam.count(d.param) == 0) { + seenParam[d.param] = 1; + fGC += ",\n size_t " + d.param + " = " + fShapeParams[d.param]; + } + } + } } + for (auto &p : fShapeParams) + if (seenParam.count(p.first) == 0) + fGC += ",\n size_t " + p.first + " = " + p.second; } fGC += ") {\n"; diff --git a/core/src/SOFIE_common.cxx b/core/src/SOFIE_common.cxx index a2bafde..af3f5fb 100644 --- a/core/src/SOFIE_common.cxx +++ b/core/src/SOFIE_common.cxx @@ -556,6 +556,31 @@ std::vector UTILITY::ComputeStrideFromShape(const std::vector & shape) return strides; } +void UTILITY::CollectDimParams(const std::vector & dims, std::vector & params) { + auto isAlpha = [](char c) { return (c >= 'a' && c <= 'z') || (c >= 'A' && c <= 'Z') || c == '_'; }; + auto isAlnum = [&](char c) { return isAlpha(c) || (c >= '0' && c <= '9'); }; + for (auto & d : dims) { + if (!d.isParam) continue; + const std::string & e = d.param; + std::size_t i = 0; + while (i < e.size()) { + if (isAlpha(e[i])) { + std::size_t j = i; + while (j < e.size() && isAlnum(e[j])) j++; + std::string tok = e.substr(i, j - i); + // skip function names from expressions like std::max(...) + bool skip = (tok == "std" || tok == "max" || tok == "min"); + for (auto & p : params) + if (p == tok) skip = true; + if (!skip) params.push_back(tok); + i = j; + } else { + i++; + } + } + } +} + struct FreeBlock { std::size_t offset; std::size_t size; diff --git a/test/TestCustomModelsFromONNXForAlpakaCuda.cxx b/test/TestCustomModelsFromONNXForAlpakaCuda.cxx index fccacbe..da13ac8 100644 --- a/test/TestCustomModelsFromONNXForAlpakaCuda.cxx +++ b/test/TestCustomModelsFromONNXForAlpakaCuda.cxx @@ -64,6 +64,17 @@ #include "Tile5D_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/Tile5D.ref.hxx" +#include "DynamicTile_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicEqual_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicAddBroadcast_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicGather_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicTranspose_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicConcat_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicSlice_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicNegRelu_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicLinear_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicConv1DNoBias_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicConv2DNoBias_FromONNX_GPU_ALPAKA.hxx" #include "GatherAxis0_FromONNX_GPU_ALPAKA.hxx" #include "GatherAxis1_FromONNX_GPU_ALPAKA.hxx" @@ -139,6 +150,10 @@ #include "ReduceMax_FromONNX_GPU_ALPAKA.hxx" #include "ReduceMax_axis0_FromONNX_GPU_ALPAKA.hxx" #include "ReduceMax_mid_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicReduceSumLast_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicReduceMeanMid_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicReduceMaxFirst_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicReduceSumMulti_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/ReduceMean.ref.hxx" #include "input_models/references/ReduceProd.ref.hxx" #include "input_models/references/ReduceL2.ref.hxx" @@ -162,10 +177,14 @@ #include "input_models/references/ConvWithStridesNoPadding.ref.hxx" #include "ConvWithAsymmetricPadding_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicConv1D_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/ConvWithAsymmetricPadding.ref.hxx" #include "BatchNorm_FromONNX_GPU_ALPAKA.hxx" #include "BatchNormRelu_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicBatchNorm4D_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicBatchNormDynSpatialRelu_FromONNX_GPU_ALPAKA.hxx" +#include "DynamicBatchNorm2D_FromONNX_GPU_ALPAKA.hxx" #include "LayerNorm_FromONNX_GPU_ALPAKA.hxx" #include "LayerNormScaleBias_FromONNX_GPU_ALPAKA.hxx" @@ -453,6 +472,46 @@ TEST_F(SofieAlpakaTest, Transpose) EXPECT_LE(std::abs(res_ptr[i] - expected[i]), TOLERANCE); } +TEST_F(SofieAlpakaTest, DynamicTranspose) +{ + // X[N,3,n_pf] perm(0,2,1) -> Y[N,n_pf,3], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 3; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t sz = N * C * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < sz; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicTranspose::Session session("", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t p = 0; p < P; ++p) + for (std::size_t c = 0; c < C; ++c) { + float expected = in_ptr[n * C * P + c * P + p]; + float got = res[n * P * C + p * C + c]; + EXPECT_LE(std::abs(got - expected), TOLERANCE) << "n=" << n << " p=" << p << " c=" << c; + } + } +} + TEST_F(SofieAlpakaTest, Concat0D) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -496,6 +555,52 @@ TEST_F(SofieAlpakaTest, Concat0D) } } +TEST_F(SofieAlpakaTest, DynamicConcat) +{ + // A[N,2,n_pf] + B[N,3,n_pf] axis 1 -> Y[N,5,n_pf], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t Ca = 2, Cb = 3, Cy = 5; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t szA = N * Ca * P, szB = N * Cb * P, szY = N * Cy * P; + + auto a_h = alpaka::allocBuf(host, Ext1D::all(Idx{szA})); + auto b_h = alpaka::allocBuf(host, Ext1D::all(Idx{szB})); + float* a_ptr = reinterpret_cast(alpaka::getPtrNative(a_h)); + float* b_ptr = reinterpret_cast(alpaka::getPtrNative(b_h)); + for (Idx i = 0; i < szA; ++i) a_ptr[i] = static_cast(i % 7) - 3.0f; + for (Idx i = 0; i < szB; ++i) b_ptr[i] = static_cast(i % 5) + 10.0f; + + auto a_d = alpaka::allocBuf(device, Ext1D::all(Idx{szA})); + auto b_d = alpaka::allocBuf(device, Ext1D::all(Idx{szB})); + alpaka::memcpy(queue, a_d, a_h); + alpaka::memcpy(queue, b_d, b_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{szY})); + { + SOFIE_DynamicConcat::Session session("", N, P); + auto result = session.infer(N, P, a_d, b_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t j = 0; j < Cy; ++j) + for (std::size_t p = 0; p < P; ++p) { + float expected = (j < Ca) ? a_ptr[n * Ca * P + j * P + p] + : b_ptr[n * Cb * P + (j - Ca) * P + p]; + float got = res[n * Cy * P + j * P + p]; + EXPECT_LE(std::abs(got - expected), TOLERANCE) << "n=" << n << " j=" << j << " p=" << p; + } + } +} + TEST_F(SofieAlpakaTest, ScatterElements) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -722,6 +827,43 @@ TEST_F(SofieAlpakaTest, Tile5D) EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE); } +TEST_F(SofieAlpakaTest, DynamicTile) +{ + // X[N,2] -> Tile([2,3]) -> Y[2N,6], with N dynamic. Run at two batch sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t inCols = 2, outCols = 6; + + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t inRows = N, outRows = 2 * N; + const std::size_t inSize = inRows * inCols, outSize = outRows * outCols; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) + in_ptr[i] = static_cast(i + 1); + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicTile::Session session("", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t r = 0; r < outRows; ++r) + for (std::size_t c = 0; c < outCols; ++c) { + float expected = in_ptr[(r % inRows) * inCols + (c % inCols)]; + EXPECT_LE(std::abs(res[r * outCols + c] - expected), TOLERANCE); + } + } +} + TEST_F(SofieAlpakaTest, GatherAxis0) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -996,6 +1138,47 @@ TEST_F(SofieAlpakaTest, ExpandDiffSize) EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE); } +TEST_F(SofieAlpakaTest, DynamicGather) +{ + // X[N,3,n_pf], indices [2,0], axis 1 -> Y[N,2,n_pf], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 3, K = 2; + const std::size_t idxs[2] = {2, 0}; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t inSize = N * C * P, outSize = N * K * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicGather::Session session("", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t i = 0; i < K; ++i) + for (std::size_t p = 0; p < P; ++p) { + float expected = in_ptr[n * C * P + idxs[i] * P + p]; + float got = res[n * K * P + i * P + p]; + EXPECT_LE(std::abs(got - expected), TOLERANCE) << "n=" << n << " i=" << i << " p=" << p; + } + } +} + TEST_F(SofieAlpakaTest, GatherND_Ex1) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -1259,6 +1442,121 @@ TEST_F(SofieAlpakaTest, Equal) EXPECT_EQ(res_ptr[i], correct[i]) << "i=" << i; } +TEST_F(SofieAlpakaTest, DynamicEqual) +{ + // X1[N,3] == X2[N,3] -> Y[N,3] (bool/uint8), N dynamic. Run at two batch sizes. + const std::size_t cols = 3; + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t sz = N * cols; + + auto x1_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + auto x2_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* x1p = reinterpret_cast(alpaka::getPtrNative(x1_h)); + float* x2p = reinterpret_cast(alpaka::getPtrNative(x2_h)); + for (Idx i = 0; i < sz; ++i) { + x1p[i] = static_cast(i % 3); + x2p[i] = static_cast(i % 2); // mix of equal and unequal + } + + auto x1_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + auto x2_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, x1_d, x1_h); + alpaka::memcpy(queue, x2_d, x2_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicEqual::Session session("", N); + auto result = session.infer(N, x1_d, x2_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + uint8_t* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t i = 0; i < sz; ++i) { + uint8_t expected = (x1p[i] == x2p[i]) ? 1 : 0; + EXPECT_EQ(res[i], expected); + } + } +} + +TEST_F(SofieAlpakaTest, DynamicAddBroadcast) +{ + // X[N,4,n_pf] + bias[4,1] -> Y[N,4,n_pf], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 4; + const float bias[4] = {0.5f, -1.0f, 0.25f, 2.0f}; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t sz = N * C * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < sz; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicAddBroadcast::Session session("DynamicAddBroadcast_FromONNX_GPU_ALPAKA.dat", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t i = 0; i < sz; ++i) { + std::size_t c = (i / P) % C; + float expected = in_ptr[i] + bias[c]; + EXPECT_LE(std::abs(res[i] - expected), TOLERANCE) << "i=" << i << " N=" << N << " P=" << P; + } + } +} + +TEST_F(SofieAlpakaTest, DynamicNegRelu) +{ + // X[N,3,n_pf] -> Neg -> Relu -> Y, fused eltwise chain, N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 3; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t sz = N * C * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < sz; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicNegRelu::Session session("", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t i = 0; i < sz; ++i) { + float expected = std::max(0.0f, -in_ptr[i]); + EXPECT_LE(std::abs(res[i] - expected), TOLERANCE) << "i=" << i << " N=" << N << " P=" << P; + } + } +} + TEST_F(SofieAlpakaTest, LessOrEqual) { std::vector input1 = {1.0f, 2.0f, 3.0f}; @@ -1539,6 +1837,46 @@ TEST_F(SofieAlpakaTest, Slice_Neg) EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE) << "i=" << i; } +TEST_F(SofieAlpakaTest, DynamicSlice) +{ + // X[N,4,n_pf] axis 1 [1:3] -> Y[N,2,n_pf], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 4, K = 2, start = 1; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t inSize = N * C * P, outSize = N * K * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicSlice::Session session("", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t j = 0; j < K; ++j) + for (std::size_t p = 0; p < P; ++p) { + float expected = in_ptr[n * C * P + (start + j) * P + p]; + float got = res[n * K * P + j * P + p]; + EXPECT_LE(std::abs(got - expected), TOLERANCE) << "n=" << n << " j=" << j << " p=" << p; + } + } +} + TEST_F(SofieAlpakaTest, Sin) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -2214,6 +2552,160 @@ TEST_F(SofieAlpakaTest, ReduceMax_mid) EXPECT_NEAR(res_ptr[i], correct[i], TOLERANCE) << " i=" << i; } +// ── Dynamic (symbolic-shape) Reduce tests: each runs at two runtime sizes so a +// baked-in shape would fail one of them. References are computed in-test. + +// DynamicReduceSumLast: X[N,4] -> ReduceSum(axis=-1, keepdims=0) -> Y[N] +// last-axis (kLast), pruned output, negative axis, static reduced-length. +TEST_F(SofieAlpakaTest, DynamicReduceSumLast) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t cols = 4; + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t inSize = N * cols, outSize = N; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i + 1); + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicReduceSumLast::Session session("", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) { + float expected = 0.f; + for (std::size_t c = 0; c < cols; ++c) expected += in_ptr[n * cols + c]; + EXPECT_LE(std::abs(res[n] - expected), TOLERANCE); + } + } +} + +// DynamicReduceMeanMid: X[N,4,M] -> ReduceMean(axis=1, keepdims=1) -> Y[N,1,M] +// middle-axis (kMiddle) with TWO symbolic dims. Reconstruct the session per size so +// the returned Y buffer matches outSize; the output length N*M is a symmetric product +// so the unordered_map ctor-arg order is immaterial, and infer args stay in +// declaration order (N, M). +TEST_F(SofieAlpakaTest, DynamicReduceMeanMid) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t K = 4; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ms[] = {1, 3}; + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], M = Ms[t]; + const std::size_t inSize = N * K * M, outSize = N * M; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i + 1); + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicReduceMeanMid::Session session("", N, M); + auto result = session.infer(N, M, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t m = 0; m < M; ++m) { + float sum = 0.f; + for (std::size_t k = 0; k < K; ++k) + sum += in_ptr[n * K * M + k * M + m]; + float expected = sum / static_cast(K); + EXPECT_LE(std::abs(res[n * M + m] - expected), TOLERANCE); + } + } +} + +// DynamicReduceMaxFirst: X[N,4] -> ReduceMax(axis=0, keepdims=0) -> Y[4] +// reduces the DYNAMIC axis, so reducedLength is symbolic (=N); kFirst path, Max op. +TEST_F(SofieAlpakaTest, DynamicReduceMaxFirst) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t cols = 4; + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t inSize = N * cols, outSize = cols; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast((i * 7 + 3) % 11); + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicReduceMaxFirst::Session session("", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t c = 0; c < cols; ++c) { + float expected = in_ptr[c]; // n == 0 + for (std::size_t n = 1; n < N; ++n) + if (in_ptr[n * cols + c] > expected) expected = in_ptr[n * cols + c]; + EXPECT_LE(std::abs(res[c] - expected), TOLERANCE); + } + } +} + +// DynamicReduceSumMulti: X[N,3,2] -> ReduceSum(axes=[1,2], keepdims=0) -> Y[N] +// multi-axis reduce -> exercises the redStrides product string in the kernel. +TEST_F(SofieAlpakaTest, DynamicReduceSumMulti) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t slice = 3 * 2; + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t inSize = N * slice, outSize = N; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i + 1); + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicReduceSumMulti::Session session("", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) { + float expected = 0.f; + for (std::size_t j = 0; j < slice; ++j) expected += in_ptr[n * slice + j]; + EXPECT_LE(std::abs(res[n] - expected), TOLERANCE); + } + } +} + TEST_F(SofieAlpakaTest, ConvWithPadding) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -2442,6 +2934,185 @@ TEST_F(SofieAlpakaTest, ConvWithAsymmetricPadding) } } +// dynamic 1D Conv (k=3, pad=1) + bias, run at (N,n_pf) = (1,1) and (8,5) +TEST_F(SofieAlpakaTest, DynamicConv1D) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t Cin = 2, Cout = 3, K = 3; + const float bias[3] = {0.5f, -0.5f, 1.0f}; + auto Wv = [&](std::size_t oc, std::size_t ic, std::size_t kk) { + return static_cast(oc * (Cin * K) + ic * K + kk); + }; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t inSize = N * Cin * P, outSize = N * Cout * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i % 10) - 5.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicConv1D::Session session("DynamicConv1D_FromONNX_GPU_ALPAKA.dat", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t oc = 0; oc < Cout; ++oc) + for (std::size_t l = 0; l < P; ++l) { + float acc = bias[oc]; + for (std::size_t ic = 0; ic < Cin; ++ic) + for (std::size_t kk = 0; kk < K; ++kk) { + int64_t li = static_cast(l) + static_cast(kk) - 1; // pad=1 + if (li >= 0 && li < static_cast(P)) + acc += in_ptr[n * Cin * P + ic * P + li] * Wv(oc, ic, kk); + } + std::size_t idx = n * Cout * P + oc * P + l; + EXPECT_LE(std::abs(res[idx] - acc), TOLERANCE) << "n=" << n << " oc=" << oc << " l=" << l; + } + } +} + +TEST_F(SofieAlpakaTest, DynamicLinear) +{ + // X[N,4] -> Gemm(W[4,3],B[3]) -> Relu -> Y[N,3], N dynamic. Gemm+Relu fuses to gemmrelu. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t In = 4, Out = 3; + const float W[4][3] = {{0.5f, -1.0f, 0.25f}, + {0.75f, 0.5f, -0.5f}, + {-0.25f, 1.0f, 0.75f}, + {1.5f, -0.75f, 0.5f}}; + const float B[3] = {0.5f, -0.5f, 0.25f}; + + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t inSize = N * In, outSize = N * Out; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicLinear::Session session("DynamicLinear_FromONNX_GPU_ALPAKA.dat", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t j = 0; j < Out; ++j) { + float acc = B[j]; + for (std::size_t i = 0; i < In; ++i) acc += in_ptr[n * In + i] * W[i][j]; + float expected = acc > 0.0f ? acc : 0.0f; + EXPECT_LE(std::abs(res[n * Out + j] - expected), TOLERANCE) << "n=" << n << " j=" << j << " N=" << N; + } + } +} + +TEST_F(SofieAlpakaTest, DynamicConv1DNoBias) +{ + // X[N,2,n_pf] k=1 no-bias conv -> Y[N,3,n_pf], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t Cin = 2, Cout = 3; + const float W[3][2] = {{0.5f, -1.0f}, {0.25f, 0.75f}, {-0.5f, 1.5f}}; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t inSize = N * Cin * P, outSize = N * Cout * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicConv1DNoBias::Session session("DynamicConv1DNoBias_FromONNX_GPU_ALPAKA.dat", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t oc = 0; oc < Cout; ++oc) + for (std::size_t p = 0; p < P; ++p) { + float expected = 0.0f; + for (std::size_t ic = 0; ic < Cin; ++ic) + expected += W[oc][ic] * in_ptr[n * Cin * P + ic * P + p]; + float got = res[n * Cout * P + oc * P + p]; + EXPECT_LE(std::abs(got - expected), TOLERANCE) << "n=" << n << " oc=" << oc << " p=" << p; + } + } +} + +TEST_F(SofieAlpakaTest, DynamicConv2DNoBias) +{ + // X[N,2,n_pf,4] 1x1 no-bias conv -> Y[N,3,n_pf,4], N and n_pf dynamic. Run at two sizes. + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t Cin = 2, Cout = 3, Q = 4; + const float W[3][2] = {{0.5f, -1.0f}, {0.25f, 0.75f}, {-0.5f, 1.5f}}; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t inSize = N * Cin * P * Q, outSize = N * Cout * P * Q; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{inSize})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < inSize; ++i) in_ptr[i] = static_cast(i % 7) - 3.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{inSize})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{outSize})); + { + SOFIE_DynamicConv2DNoBias::Session session("DynamicConv2DNoBias_FromONNX_GPU_ALPAKA.dat", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t n = 0; n < N; ++n) + for (std::size_t oc = 0; oc < Cout; ++oc) + for (std::size_t p = 0; p < P; ++p) + for (std::size_t q = 0; q < Q; ++q) { + float expected = 0.0f; + for (std::size_t ic = 0; ic < Cin; ++ic) + expected += W[oc][ic] * in_ptr[n * Cin * P * Q + ic * P * Q + p * Q + q]; + float got = res[n * Cout * P * Q + oc * P * Q + p * Q + q]; + EXPECT_LE(std::abs(got - expected), TOLERANCE) << "n=" << n << " oc=" << oc << " p=" << p << " q=" << q; + } + } +} + TEST_F(SofieAlpakaTest, BatchNormalization) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -2518,6 +3189,132 @@ TEST_F(SofieAlpakaTest, BatchNormalizationRelu) } } +// dynamic BatchNorm, run at two batch sizes. weights are per-channel [C=4] (same values in +// each test); fused scale is scale[c]/sqrt(var[c]+eps). the static BN tests above cover the +// shared code path, these cover the dynamic shapes. + +// X[N,4,2,3]: dynamic batch, static spatial (6) +TEST_F(SofieAlpakaTest, DynamicBatchNorm4D) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 4, spatial = 2 * 3; + const float scale[4] = {1.0f, 2.0f, 0.5f, 1.5f}, bias[4] = {0.1f, -0.2f, 0.3f, 0.0f}; + const float mean_[4] = {0.5f, 1.0f, -0.5f, 2.0f}, var_[4] = {1.0f, 4.0f, 0.25f, 2.0f}; + const float eps = 1e-5f; + + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t sz = N * C * spatial; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < sz; ++i) in_ptr[i] = static_cast(i % 10) - 5.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicBatchNorm4D::Session session("DynamicBatchNorm4D_FromONNX_GPU_ALPAKA.dat", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t i = 0; i < sz; ++i) { + std::size_t c = (i / spatial) % C; + float fused = scale[c] / std::sqrt(var_[c] + eps); + float expected = (in_ptr[i] - mean_[c]) * fused + bias[c]; + EXPECT_LE(std::abs(res[i] - expected), TOLERANCE) << "i=" << i; + } + } +} + +// X[N,4,n_pf] + relu: dynamic batch and spatial. rebuilt per size so Y matches its length; +// Y length is a symmetric product so the ctor arg order doesn't matter, infer is (N, n_pf). +TEST_F(SofieAlpakaTest, DynamicBatchNormDynSpatialRelu) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 4; + const float scale[4] = {1.0f, 2.0f, 0.5f, 1.5f}, bias[4] = {0.1f, -0.2f, 0.3f, 0.0f}; + const float mean_[4] = {0.5f, 1.0f, -0.5f, 2.0f}, var_[4] = {1.0f, 4.0f, 0.25f, 2.0f}; + const float eps = 1e-5f; + + const std::size_t Ns[] = {1, 8}; + const std::size_t Ps[] = {1, 5}; // n_pf + for (int t = 0; t < 2; ++t) { + const std::size_t N = Ns[t], P = Ps[t]; + const std::size_t sz = N * C * P; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < sz; ++i) in_ptr[i] = static_cast(i % 10) - 5.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicBatchNormDynSpatialRelu::Session session("DynamicBatchNormDynSpatialRelu_FromONNX_GPU_ALPAKA.dat", N, P); + auto result = session.infer(N, P, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t i = 0; i < sz; ++i) { + std::size_t c = (i / P) % C; + float fused = scale[c] / std::sqrt(var_[c] + eps); + float expected = (in_ptr[i] - mean_[c]) * fused + bias[c]; + expected = expected > 0.0f ? expected : 0.0f; // relu + EXPECT_LE(std::abs(res[i] - expected), TOLERANCE) << "i=" << i; + } + } +} + +// X[N,4]: rank 2, spatial is 1 +TEST_F(SofieAlpakaTest, DynamicBatchNorm2D) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + const std::size_t C = 4; + const float scale[4] = {1.0f, 2.0f, 0.5f, 1.5f}, bias[4] = {0.1f, -0.2f, 0.3f, 0.0f}; + const float mean_[4] = {0.5f, 1.0f, -0.5f, 2.0f}, var_[4] = {1.0f, 4.0f, 0.25f, 2.0f}; + const float eps = 1e-5f; + + for (std::size_t N : {std::size_t(1), std::size_t(8)}) { + const std::size_t sz = N * C; + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + float* in_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < sz; ++i) in_ptr[i] = static_cast(i % 10) - 5.0f; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{sz})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sz})); + { + SOFIE_DynamicBatchNorm2D::Session session("DynamicBatchNorm2D_FromONNX_GPU_ALPAKA.dat", N); + auto result = session.infer(N, input_d); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res = reinterpret_cast(alpaka::getPtrNative(result_h)); + for (std::size_t i = 0; i < sz; ++i) { + std::size_t c = i % C; + float fused = scale[c] / std::sqrt(var_[c] + eps); + float expected = (in_ptr[i] - mean_[c]) * fused + bias[c]; + EXPECT_LE(std::abs(res[i] - expected), TOLERANCE) << "i=" << i; + } + } +} + TEST_F(SofieAlpakaTest, LayerNorm) { constexpr float TOLERANCE = 1e-4f; diff --git a/test/input_models/DynamicAddBroadcast.onnx b/test/input_models/DynamicAddBroadcast.onnx new file mode 100644 index 0000000..12b7cbc Binary files /dev/null and b/test/input_models/DynamicAddBroadcast.onnx differ diff --git a/test/input_models/DynamicBatchNorm2D.onnx b/test/input_models/DynamicBatchNorm2D.onnx new file mode 100644 index 0000000..a8219dd Binary files /dev/null and b/test/input_models/DynamicBatchNorm2D.onnx differ diff --git a/test/input_models/DynamicBatchNorm4D.onnx b/test/input_models/DynamicBatchNorm4D.onnx new file mode 100644 index 0000000..8aecf1a Binary files /dev/null and b/test/input_models/DynamicBatchNorm4D.onnx differ diff --git a/test/input_models/DynamicBatchNormDynSpatialRelu.onnx b/test/input_models/DynamicBatchNormDynSpatialRelu.onnx new file mode 100644 index 0000000..626e3fa Binary files /dev/null and b/test/input_models/DynamicBatchNormDynSpatialRelu.onnx differ diff --git a/test/input_models/DynamicConcat.onnx b/test/input_models/DynamicConcat.onnx new file mode 100644 index 0000000..bef8904 Binary files /dev/null and b/test/input_models/DynamicConcat.onnx differ diff --git a/test/input_models/DynamicConv1D.onnx b/test/input_models/DynamicConv1D.onnx new file mode 100644 index 0000000..b97f8f8 Binary files /dev/null and b/test/input_models/DynamicConv1D.onnx differ diff --git a/test/input_models/DynamicConv1DNoBias.onnx b/test/input_models/DynamicConv1DNoBias.onnx new file mode 100644 index 0000000..387ebd4 Binary files /dev/null and b/test/input_models/DynamicConv1DNoBias.onnx differ diff --git a/test/input_models/DynamicConv2DNoBias.onnx b/test/input_models/DynamicConv2DNoBias.onnx new file mode 100644 index 0000000..86480de Binary files /dev/null and b/test/input_models/DynamicConv2DNoBias.onnx differ diff --git a/test/input_models/DynamicEqual.onnx b/test/input_models/DynamicEqual.onnx new file mode 100644 index 0000000..8627ba1 Binary files /dev/null and b/test/input_models/DynamicEqual.onnx differ diff --git a/test/input_models/DynamicGather.onnx b/test/input_models/DynamicGather.onnx new file mode 100644 index 0000000..b9ee64e Binary files /dev/null and b/test/input_models/DynamicGather.onnx differ diff --git a/test/input_models/DynamicLinear.onnx b/test/input_models/DynamicLinear.onnx new file mode 100644 index 0000000..b452679 Binary files /dev/null and b/test/input_models/DynamicLinear.onnx differ diff --git a/test/input_models/DynamicNegRelu.onnx b/test/input_models/DynamicNegRelu.onnx new file mode 100644 index 0000000..6c07396 Binary files /dev/null and b/test/input_models/DynamicNegRelu.onnx differ diff --git a/test/input_models/DynamicReduceMaxFirst.onnx b/test/input_models/DynamicReduceMaxFirst.onnx new file mode 100644 index 0000000..13fa588 Binary files /dev/null and b/test/input_models/DynamicReduceMaxFirst.onnx differ diff --git a/test/input_models/DynamicReduceMeanMid.onnx b/test/input_models/DynamicReduceMeanMid.onnx new file mode 100644 index 0000000..cb0c0cb Binary files /dev/null and b/test/input_models/DynamicReduceMeanMid.onnx differ diff --git a/test/input_models/DynamicReduceSumLast.onnx b/test/input_models/DynamicReduceSumLast.onnx new file mode 100644 index 0000000..03f7bd2 Binary files /dev/null and b/test/input_models/DynamicReduceSumLast.onnx differ diff --git a/test/input_models/DynamicReduceSumMulti.onnx b/test/input_models/DynamicReduceSumMulti.onnx new file mode 100644 index 0000000..0a4dd9f Binary files /dev/null and b/test/input_models/DynamicReduceSumMulti.onnx differ diff --git a/test/input_models/DynamicSlice.onnx b/test/input_models/DynamicSlice.onnx new file mode 100644 index 0000000..a5451af Binary files /dev/null and b/test/input_models/DynamicSlice.onnx differ diff --git a/test/input_models/DynamicTile.onnx b/test/input_models/DynamicTile.onnx new file mode 100644 index 0000000..b1f557a Binary files /dev/null and b/test/input_models/DynamicTile.onnx differ diff --git a/test/input_models/DynamicTranspose.onnx b/test/input_models/DynamicTranspose.onnx new file mode 100644 index 0000000..3f53452 Binary files /dev/null and b/test/input_models/DynamicTranspose.onnx differ