17656823807

Committed 11 Sep 2025 08:42PM UTC coverage: 60.447% (+1.1%) from 59.335%

Build # 17656823807

Build Type

Pull #219

github

Committed by

web-flow

Commit Message

Merge d5416236f into 6c1992b40

Pull Request Pull Request #219: stdlib Library Nodes and ConstantNodes

Run Details

460 of 1635 new or added lines in 81 files covered. (28.13%)

93 existing lines in 35 files now uncovered.

9385 of 15526 relevant lines covered (60.45%)

107.21 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

21.88

/src/codegen/language_extensions/cuda_language_extension.cpp

#include "sdfg/codegen/language_extensions/cuda_language_extension.h"

#include "sdfg/codegen/language_extensions/cpp_language_extension.h"
#include "sdfg/codegen/utils.h"
#include "sdfg/data_flow/library_node.h"
#include "sdfg/data_flow/tasklet.h"

namespace sdfg {
namespace codegen {

constexpr const char* code_to_string(data_flow::TaskletCode c) {
    switch (c) {
        case data_flow::TaskletCode::assign:
            return "=";
        case data_flow::TaskletCode::neg:
            return "-";
        case data_flow::TaskletCode::add:
            return "+";
        case data_flow::TaskletCode::sub:
            return "-";
        case data_flow::TaskletCode::mul:
            return "*";
        case data_flow::TaskletCode::div:
            return "/";
        case data_flow::TaskletCode::fma:
            return "__daisy_fma";
        case data_flow::TaskletCode::mod:
            return "%";
        case data_flow::TaskletCode::max:
            return "__daisy_max";
        case data_flow::TaskletCode::min:
            return "__daisy_min";
        case data_flow::TaskletCode::minnum:
            return "minnum";
        case data_flow::TaskletCode::maxnum:
            return "maxnum";
        case data_flow::TaskletCode::minimum:
            return "minimum";
        case data_flow::TaskletCode::maximum:
            return "maximum";
        case data_flow::TaskletCode::trunc:
            return "trunc";
        case data_flow::TaskletCode::logical_and:
            return "&&";
        case data_flow::TaskletCode::logical_or:
            return "||";
        case data_flow::TaskletCode::bitwise_and:
            return "&";
        case data_flow::TaskletCode::bitwise_or:
            return "|";
        case data_flow::TaskletCode::bitwise_xor:
            return "^";
        case data_flow::TaskletCode::bitwise_not:
            return "~";
        case data_flow::TaskletCode::shift_left:
            return "<<";
        case data_flow::TaskletCode::shift_right:
            return ">>";
        case data_flow::TaskletCode::olt:
            return "<";
        case data_flow::TaskletCode::ole:
            return "<=";
        case data_flow::TaskletCode::oeq:
            return "==";
        case data_flow::TaskletCode::one:
            return "!=";
        case data_flow::TaskletCode::oge:
            return ">=";
        case data_flow::TaskletCode::ogt:
            return ">";
        case data_flow::TaskletCode::ord:
            return "==";
        case data_flow::TaskletCode::ult:
            return "<";
        case data_flow::TaskletCode::ule:
            return "<=";
        case data_flow::TaskletCode::ueq:
            return "==";
        case data_flow::TaskletCode::une:
            return "!=";
        case data_flow::TaskletCode::uge:
            return ">=";
        case data_flow::TaskletCode::ugt:
            return ">";
        case data_flow::TaskletCode::uno:
            return "!=";
        case data_flow::TaskletCode::abs:
            return "abs";
        case data_flow::TaskletCode::acos:
            return "acos";
        case data_flow::TaskletCode::acosf:
            return "acosf";
        case data_flow::TaskletCode::acosl:
            return "acosl";
        case data_flow::TaskletCode::acosh:
            return "acosh";
        case data_flow::TaskletCode::acoshf:
            return "acoshf";
        case data_flow::TaskletCode::acoshl:
            return "acoshl";
        case data_flow::TaskletCode::asin:
            return "asin";
        case data_flow::TaskletCode::asinf:
            return "asinf";
        case data_flow::TaskletCode::asinl:
            return "asinl";
        case data_flow::TaskletCode::asinh:
            return "asinh";
        case data_flow::TaskletCode::asinhf:
            return "asinhf";
        case data_flow::TaskletCode::asinhl:
            return "asinhl";
        case data_flow::TaskletCode::atan:
            return "atan";
        case data_flow::TaskletCode::atanf:
            return "atanf";
        case data_flow::TaskletCode::atanl:
            return "atanl";
        case data_flow::TaskletCode::atan2:
            return "atan2";
        case data_flow::TaskletCode::atan2f:
            return "atan2f";
        case data_flow::TaskletCode::atan2l:
            return "atan2l";
        case data_flow::TaskletCode::atanh:
            return "atanh";
        case data_flow::TaskletCode::atanhf:
            return "atanhf";
        case data_flow::TaskletCode::atanhl:
            return "atanhl";
        case data_flow::TaskletCode::cabs:
            return "cabs";
        case data_flow::TaskletCode::cabsf:
            return "cabsf";
        case data_flow::TaskletCode::cabsl:
            return "cabsl";
        case data_flow::TaskletCode::ceil:
            return "ceil";
        case data_flow::TaskletCode::ceilf:
            return "ceilf";
        case data_flow::TaskletCode::ceill:
            return "ceill";
        case data_flow::TaskletCode::copysign:
            return "copysign";
        case data_flow::TaskletCode::copysignf:
            return "copysignf";
        case data_flow::TaskletCode::copysignl:
            return "copysignl";
        case data_flow::TaskletCode::cos:
            return "cos";
        case data_flow::TaskletCode::cosf:
            return "cosf";
        case data_flow::TaskletCode::cosl:
            return "cosl";
        case data_flow::TaskletCode::cosh:
            return "cosh";
        case data_flow::TaskletCode::coshf:
            return "coshf";
        case data_flow::TaskletCode::coshl:
            return "coshl";
        case data_flow::TaskletCode::cbrt:
            return "cbrt";
        case data_flow::TaskletCode::cbrtf:
            return "cbrtf";
        case data_flow::TaskletCode::cbrtl:
            return "cbrtl";
        case data_flow::TaskletCode::erf:
            return "erf";
        case data_flow::TaskletCode::erff:
            return "erff";
        case data_flow::TaskletCode::erfl:
            return "erfl";
        case data_flow::TaskletCode::exp10:
            return "exp10";
        case data_flow::TaskletCode::exp10f:
            return "exp10f";
        case data_flow::TaskletCode::exp10l:
            return "exp10l";
        case data_flow::TaskletCode::exp2:
            return "exp2";
        case data_flow::TaskletCode::exp2f:
            return "exp2f";
        case data_flow::TaskletCode::exp2l:
            return "exp2l";
        case data_flow::TaskletCode::exp:
            return "exp";
        case data_flow::TaskletCode::expf:
            return "expf";
        case data_flow::TaskletCode::expl:
            return "expl";
        case data_flow::TaskletCode::expm1:
            return "expm1";
        case data_flow::TaskletCode::expm1f:
            return "expm1f";
        case data_flow::TaskletCode::expm1l:
            return "expm1l";
        case data_flow::TaskletCode::fabs:
            return "fabs";
        case data_flow::TaskletCode::fabsf:
            return "fabsf";
        case data_flow::TaskletCode::fabsl:
            return "fabsl";
        case data_flow::TaskletCode::floor:
            return "floor";
        case data_flow::TaskletCode::floorf:
            return "floorf";
        case data_flow::TaskletCode::floorl:
            return "floorl";
        case data_flow::TaskletCode::fls:
            return "fls";
        case data_flow::TaskletCode::flsl:
            return "flsl";
        case data_flow::TaskletCode::fmax:
            return "fmax";
        case data_flow::TaskletCode::fmaxf:
            return "fmaxf";
        case data_flow::TaskletCode::fmaxl:
            return "fmaxl";
        case data_flow::TaskletCode::fmin:
            return "fmin";
        case data_flow::TaskletCode::fminf:
            return "fminf";
        case data_flow::TaskletCode::fminl:
            return "fminl";
        case data_flow::TaskletCode::fmod:
            return "fmod";
        case data_flow::TaskletCode::fmodf:
            return "fmodf";
        case data_flow::TaskletCode::fmodl:
            return "fmodl";
        case data_flow::TaskletCode::frexp:
            return "frexp";
        case data_flow::TaskletCode::frexpf:
            return "frexpf";
        case data_flow::TaskletCode::frexpl:
            return "frexpl";
        case data_flow::TaskletCode::labs:
            return "labs";
        case data_flow::TaskletCode::ldexp:
            return "ldexp";
        case data_flow::TaskletCode::ldexpf:
            return "ldexpf";
        case data_flow::TaskletCode::ldexpl:
            return "ldexpl";
        case data_flow::TaskletCode::log10:
            return "log10";
        case data_flow::TaskletCode::log10f:
            return "log10f";
        case data_flow::TaskletCode::log10l:
            return "log10l";
        case data_flow::TaskletCode::log2:
            return "log2";
        case data_flow::TaskletCode::log2f:
            return "log2f";
        case data_flow::TaskletCode::log2l:
            return "log2l";
        case data_flow::TaskletCode::log:
            return "log";
        case data_flow::TaskletCode::logf:
            return "logf";
        case data_flow::TaskletCode::logl:
            return "logl";
        case data_flow::TaskletCode::logb:
            return "logb";
        case data_flow::TaskletCode::logbf:
            return "logbf";
        case data_flow::TaskletCode::logbl:
            return "logbl";
        case data_flow::TaskletCode::log1p:
            return "log1p";
        case data_flow::TaskletCode::log1pf:
            return "log1pf";
        case data_flow::TaskletCode::log1pl:
            return "log1pl";
        case data_flow::TaskletCode::modf:
            return "modf";
        case data_flow::TaskletCode::modff:
            return "modff";
        case data_flow::TaskletCode::modfl:
            return "modfl";
        case data_flow::TaskletCode::nearbyint:
            return "nearbyint";
        case data_flow::TaskletCode::nearbyintf:
            return "nearbyintf";
        case data_flow::TaskletCode::nearbyintl:
            return "nearbyintl";
        case data_flow::TaskletCode::pow:
            return "pow";
        case data_flow::TaskletCode::powf:
            return "powf";
        case data_flow::TaskletCode::powl:
            return "powl";
        case data_flow::TaskletCode::rint:
            return "rint";
        case data_flow::TaskletCode::rintf:
            return "rintf";
        case data_flow::TaskletCode::rintl:
            return "rintl";
        case data_flow::TaskletCode::lrint:
            return "lrint";
        case data_flow::TaskletCode::llrint:
            return "llrint";
        case data_flow::TaskletCode::lround:
            return "lround";
        case data_flow::TaskletCode::llround:
            return "llround";
        case data_flow::TaskletCode::round:
            return "round";
        case data_flow::TaskletCode::roundf:
            return "roundf";
        case data_flow::TaskletCode::roundl:
            return "roundl";
        case data_flow::TaskletCode::roundeven:
            return "roundeven";
        case data_flow::TaskletCode::roundevenf:
            return "roundevenf";
        case data_flow::TaskletCode::roundevenl:
            return "roundevenl";
        case data_flow::TaskletCode::sin:
            return "sin";
        case data_flow::TaskletCode::sinf:
            return "sinf";
        case data_flow::TaskletCode::sinl:
            return "sinl";
        case data_flow::TaskletCode::sinh:
            return "sinh";
        case data_flow::TaskletCode::sinhf:
            return "sinhf";
        case data_flow::TaskletCode::sinhl:
            return "sinhl";
        case data_flow::TaskletCode::sqrt:
            return "sqrt";
        case data_flow::TaskletCode::sqrtf:
            return "sqrtf";
        case data_flow::TaskletCode::sqrtl:
            return "sqrtl";
        case data_flow::TaskletCode::rsqrt:
            return "rsqrt";
        case data_flow::TaskletCode::rsqrtf:
            return "rsqrtf";
        case data_flow::TaskletCode::rsqrtl:
            return "rsqrtl";
        case data_flow::TaskletCode::tan:
            return "tan";
        case data_flow::TaskletCode::tanf:
            return "tanf";
        case data_flow::TaskletCode::tanl:
            return "tanl";
        case data_flow::TaskletCode::tanh:
            return "tanh";
        case data_flow::TaskletCode::tanhf:
            return "tanhf";
        case data_flow::TaskletCode::tanhl:
            return "tanhl";
    };
    throw std::invalid_argument("Invalid tasklet code");
};

std::string CUDALanguageExtension::primitive_type(const types::PrimitiveType prim_type) {
    switch (prim_type) {
        case types::PrimitiveType::Void:
            return "void";
        case types::PrimitiveType::Bool:
            return "bool";
        case types::PrimitiveType::Int8:
            return "signed char";
        case types::PrimitiveType::Int16:
            return "short";
        case types::PrimitiveType::Int32:
            return "int";
        case types::PrimitiveType::Int64:
            return "long long";
        case types::PrimitiveType::Int128:
            return "__int128";
        case types::PrimitiveType::UInt8:
            return "char";
        case types::PrimitiveType::UInt16:
            return "unsigned short";
        case types::PrimitiveType::UInt32:
            return "unsigned int";
        case types::PrimitiveType::UInt64:
            return "unsigned long long";
        case types::PrimitiveType::UInt128:
            return "unsigned __int128";
        case types::PrimitiveType::Half:
            return "__fp16";
        case types::PrimitiveType::BFloat:
            return "__bf16";
        case types::PrimitiveType::Float:
            return "float";
        case types::PrimitiveType::Double:
            return "double";
        case types::PrimitiveType::X86_FP80:
            return "long double";
        case types::PrimitiveType::FP128:
            return "__float128";
        case types::PrimitiveType::PPC_FP128:
            return "__float128";
    }

    throw std::runtime_error("Unknown primitive type");
};

std::string CUDALanguageExtension::
    declaration(const std::string& name, const types::IType& type, bool use_initializer, bool use_alignment) {
    std::stringstream val;

    if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {
        if (scalar_type->storage_type() == types::StorageType_NV_Shared) {
            val << "__shared__ ";
        } else if (scalar_type->storage_type() == types::StorageType_NV_Constant) {
            val << "__constant__ ";
        }
        val << primitive_type(scalar_type->primitive_type());
        val << " ";
        val << name;
    } else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {
        auto& element_type = array_type->element_type();
        val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]", element_type);
    } else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {
        if (pointer_type->has_pointee_type()) {
            const types::IType& pointee = pointer_type->pointee_type();

            const bool pointee_is_function_or_array = dynamic_cast<const types::Function*>(&pointee) ||
                                                      dynamic_cast<const types::Array*>(&pointee);

            // Parenthesise *only* when it is needed to bind tighter than [] or ()
            std::string decorated = pointee_is_function_or_array ? "(*" + name + ")" : "*" + name;

            val << declaration(decorated, pointee);
        } else {
            val << "void*";
            val << " " << name;
        }
    } else if (auto ref_type = dynamic_cast<const Reference*>(&type)) {
        val << declaration("&" + name, ref_type->reference_type());
    } else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {
        if (structure_type->storage_type() == types::StorageType_NV_Shared) {
            val << "__shared__ ";
        } else if (structure_type->storage_type() == types::StorageType_NV_Constant) {
            val << "__constant__ ";
        }
        val << structure_type->name();
        val << " ";
        val << name;
    } else if (auto function_type = dynamic_cast<const types::Function*>(&type)) {
        std::stringstream params;
        for (size_t i = 0; i < function_type->num_params(); ++i) {
            params << declaration("", function_type->param_type(symbolic::integer(i)));
            if (i + 1 < function_type->num_params()) params << ", ";
        }
        if (function_type->is_var_arg()) {
            // ISO C++ forbids empty parameter lists before ...
            if (function_type->num_params() > 0) {
                params << ", ";
                params << "...";
            }
        }

        const std::string fun_name = name + "(" + params.str() + ")";
        val << declaration(fun_name, function_type->return_type());
    } else {
        throw std::runtime_error("Unknown declaration type");
    }

    if (use_alignment && type.alignment() > 0) {
        val << " __attribute__((aligned(" << type.alignment() << ")))";
    }

    if (use_initializer && !type.initializer().empty()) {
        val << " = " << type.initializer();
    }

    return val.str();
};

std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {
    std::stringstream val;

    val << "reinterpret_cast";
    val << "<";
    val << declaration("", type);
    val << ">";
    val << "(" << name << ")";

    return val.str();
};

std::string CUDALanguageExtension::subset(const Function& function, const types::IType& type, const data_flow::Subset& sub) {
    if (sub.empty()) {
        return "";
    }

    if (dynamic_cast<const types::Scalar*>(&type)) {
        return "";
    } else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {
        std::string subset_str = "[" + this->expression(sub.at(0)) + "]";

        if (sub.size() > 1) {
            data_flow::Subset element_subset(sub.begin() + 1, sub.end());
            auto& element_type = array_type->element_type();
            return subset_str + subset(function, element_type, element_subset);
        } else {
            return subset_str;
        }
    } else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {
        std::string subset_str = "[" + this->expression(sub.at(0)) + "]";

        data_flow::Subset element_subset(sub.begin() + 1, sub.end());
        auto& pointee_type = pointer_type->pointee_type();
        return subset_str + subset(function, pointee_type, element_subset);
    } else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {
        auto& definition = function.structure(structure_type->name());

        std::string subset_str = ".member_" + this->expression(sub.at(0));
        if (sub.size() > 1) {
            auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));
            auto& member_type = definition.member_type(member);
            data_flow::Subset element_subset(sub.begin() + 1, sub.end());
            return subset_str + subset(function, member_type, element_subset);
        } else {
            return subset_str;
        }
    }

    throw std::invalid_argument("Invalid subset type");
};

std::string CUDALanguageExtension::expression(const symbolic::Expression& expr) {
    CPPSymbolicPrinter printer;
    return printer.apply(expr);
};

std::string CUDALanguageExtension::access_node(const data_flow::AccessNode& node) {
    if (dynamic_cast<const data_flow::ConstantNode*>(&node)) {
        std::string name = node.data();
        if (symbolic::is_nullptr(symbolic::symbol(name))) {
            return this->expression(symbolic::__nullptr__());
        }
        return name;
    } else {
        std::string name = node.data();
        if (this->external_variables_.find(name) != this->external_variables_.end()) {
            return "(&" + name + ")";
        }
        return name;
    }
};

std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {
    std::string op = code_to_string(tasklet.code());
    std::vector<std::string> arguments;
    for (size_t i = 0; i < tasklet.inputs().size(); ++i) {
        arguments.push_back(tasklet.input(i));
    }

    if (tasklet.code() == data_flow::TaskletCode::assign) {
        return arguments.at(0);
    } else if (data_flow::is_infix(tasklet.code())) {
        switch (data_flow::arity(tasklet.code())) {
            case 1:
                return op + arguments.at(0);
            case 2:
                return arguments.at(0) + " " + op + " " + arguments.at(1);
            default:
                throw std::runtime_error("Unsupported arity");
        }
    } else {
        return op + "(" + helpers::join(arguments, ", ") + ")";
    }
};

std::string CUDALanguageExtension::zero(const types::PrimitiveType prim_type) {
    switch (prim_type) {
        case types::Void:
            throw InvalidSDFGException("No zero for void type possible");
        case types::Bool:
            return "false";
        case types::Int8:
            return "0";
        case types::Int16:
            return "0";
        case types::Int32:
            return "0";
        case types::Int64:
            return "0ll";
        case types::Int128:
            return "0";
        case types::UInt8:
            return "0u";
        case types::UInt16:
            return "0u";
        case types::UInt32:
            return "0u";
        case types::UInt64:
            return "0ull";
        case types::UInt128:
            return "0";
        case types::Half:
            return "CUDART_ZERO_FP16";
        case types::BFloat:
            return "CUDART_ZERO_BF16";
        case types::Float:
            return "0.0f";
        case types::Double:
            return "0.0";
        case types::X86_FP80:
            return "0.0l";
        case types::FP128:
            return "0.0";
        case types::PPC_FP128:
            return "0.0";
    }
}

} // namespace codegen
} // namespace sdfg

1	#include "sdfg/codegen/language_extensions/cuda_language_extension.h"
2
3	#include "sdfg/codegen/language_extensions/cpp_language_extension.h"
4	#include "sdfg/codegen/utils.h"
5	#include "sdfg/data_flow/library_node.h"
6	#include "sdfg/data_flow/tasklet.h"
7
8	namespace sdfg {
9	namespace codegen {
10
11	constexpr const char* code_to_string(data_flow::TaskletCode c) {	×
12	switch (c) {	×
13	case data_flow::TaskletCode::assign:
14	return "=";	×
15	case data_flow::TaskletCode::neg:
16	return "-";	×
17	case data_flow::TaskletCode::add:
18	return "+";	×
19	case data_flow::TaskletCode::sub:
20	return "-";	×
21	case data_flow::TaskletCode::mul:
22	return "*";	×
23	case data_flow::TaskletCode::div:
24	return "/";	×
25	case data_flow::TaskletCode::fma:
26	return "__daisy_fma";	×
27	case data_flow::TaskletCode::mod:
28	return "%";	×
29	case data_flow::TaskletCode::max:
30	return "__daisy_max";	×
31	case data_flow::TaskletCode::min:
32	return "__daisy_min";	×
33	case data_flow::TaskletCode::minnum:
34	return "minnum";	×
35	case data_flow::TaskletCode::maxnum:
36	return "maxnum";	×
37	case data_flow::TaskletCode::minimum:
38	return "minimum";	×
39	case data_flow::TaskletCode::maximum:
40	return "maximum";	×
41	case data_flow::TaskletCode::trunc:
42	return "trunc";	×
43	case data_flow::TaskletCode::logical_and:
44	return "&&";	×
45	case data_flow::TaskletCode::logical_or:
46	return "\|\|";	×
47	case data_flow::TaskletCode::bitwise_and:
48	return "&";	×
49	case data_flow::TaskletCode::bitwise_or:
50	return "\|";	×
51	case data_flow::TaskletCode::bitwise_xor:
52	return "^";	×
53	case data_flow::TaskletCode::bitwise_not:
54	return "~";	×
55	case data_flow::TaskletCode::shift_left:
56	return "<<";	×
57	case data_flow::TaskletCode::shift_right:
58	return ">>";	×
59	case data_flow::TaskletCode::olt:
60	return "<";	×
61	case data_flow::TaskletCode::ole:
62	return "<=";	×
63	case data_flow::TaskletCode::oeq:
64	return "==";	×
65	case data_flow::TaskletCode::one:
66	return "!=";	×
67	case data_flow::TaskletCode::oge:
68	return ">=";	×
69	case data_flow::TaskletCode::ogt:
70	return ">";	×
71	case data_flow::TaskletCode::ord:
72	return "==";	×
73	case data_flow::TaskletCode::ult:
74	return "<";	×
75	case data_flow::TaskletCode::ule:
76	return "<=";	×
77	case data_flow::TaskletCode::ueq:
78	return "==";	×
79	case data_flow::TaskletCode::une:
80	return "!=";	×
81	case data_flow::TaskletCode::uge:
82	return ">=";	×
83	case data_flow::TaskletCode::ugt:
84	return ">";	×
85	case data_flow::TaskletCode::uno:
86	return "!=";	×
87	case data_flow::TaskletCode::abs:
88	return "abs";	×
89	case data_flow::TaskletCode::acos:
90	return "acos";	×
91	case data_flow::TaskletCode::acosf:
92	return "acosf";	×
93	case data_flow::TaskletCode::acosl:
94	return "acosl";	×
95	case data_flow::TaskletCode::acosh:
96	return "acosh";	×
97	case data_flow::TaskletCode::acoshf:
98	return "acoshf";	×
99	case data_flow::TaskletCode::acoshl:
100	return "acoshl";	×
101	case data_flow::TaskletCode::asin:
102	return "asin";	×
103	case data_flow::TaskletCode::asinf:
104	return "asinf";	×
105	case data_flow::TaskletCode::asinl:
106	return "asinl";	×
107	case data_flow::TaskletCode::asinh:
108	return "asinh";	×
109	case data_flow::TaskletCode::asinhf:
110	return "asinhf";	×
111	case data_flow::TaskletCode::asinhl:
112	return "asinhl";	×
113	case data_flow::TaskletCode::atan:
114	return "atan";	×
115	case data_flow::TaskletCode::atanf:
116	return "atanf";	×
117	case data_flow::TaskletCode::atanl:
118	return "atanl";	×
119	case data_flow::TaskletCode::atan2:
120	return "atan2";	×
121	case data_flow::TaskletCode::atan2f:
122	return "atan2f";	×
123	case data_flow::TaskletCode::atan2l:
124	return "atan2l";	×
125	case data_flow::TaskletCode::atanh:
126	return "atanh";	×
127	case data_flow::TaskletCode::atanhf:
128	return "atanhf";	×
129	case data_flow::TaskletCode::atanhl:
130	return "atanhl";	×
131	case data_flow::TaskletCode::cabs:
132	return "cabs";	×
133	case data_flow::TaskletCode::cabsf:
134	return "cabsf";	×
135	case data_flow::TaskletCode::cabsl:
136	return "cabsl";	×
137	case data_flow::TaskletCode::ceil:
138	return "ceil";	×
139	case data_flow::TaskletCode::ceilf:
140	return "ceilf";	×
141	case data_flow::TaskletCode::ceill:
142	return "ceill";	×
143	case data_flow::TaskletCode::copysign:
144	return "copysign";	×
145	case data_flow::TaskletCode::copysignf:
146	return "copysignf";	×
147	case data_flow::TaskletCode::copysignl:
148	return "copysignl";	×
149	case data_flow::TaskletCode::cos:
150	return "cos";	×
151	case data_flow::TaskletCode::cosf:
152	return "cosf";	×
153	case data_flow::TaskletCode::cosl:
154	return "cosl";	×
155	case data_flow::TaskletCode::cosh:
156	return "cosh";	×
157	case data_flow::TaskletCode::coshf:
158	return "coshf";	×
159	case data_flow::TaskletCode::coshl:
160	return "coshl";	×
161	case data_flow::TaskletCode::cbrt:
162	return "cbrt";	×
163	case data_flow::TaskletCode::cbrtf:
164	return "cbrtf";	×
165	case data_flow::TaskletCode::cbrtl:
166	return "cbrtl";	×
167	case data_flow::TaskletCode::erf:
168	return "erf";	×
169	case data_flow::TaskletCode::erff:
170	return "erff";	×
171	case data_flow::TaskletCode::erfl:
172	return "erfl";	×
173	case data_flow::TaskletCode::exp10:
174	return "exp10";	×
175	case data_flow::TaskletCode::exp10f:
176	return "exp10f";	×
177	case data_flow::TaskletCode::exp10l:
178	return "exp10l";	×
179	case data_flow::TaskletCode::exp2:
180	return "exp2";	×
181	case data_flow::TaskletCode::exp2f:
182	return "exp2f";	×
183	case data_flow::TaskletCode::exp2l:
184	return "exp2l";	×
185	case data_flow::TaskletCode::exp:
186	return "exp";	×
187	case data_flow::TaskletCode::expf:
188	return "expf";	×
189	case data_flow::TaskletCode::expl:
190	return "expl";	×
191	case data_flow::TaskletCode::expm1:
192	return "expm1";	×
193	case data_flow::TaskletCode::expm1f:
194	return "expm1f";	×
195	case data_flow::TaskletCode::expm1l:
196	return "expm1l";	×
197	case data_flow::TaskletCode::fabs:
198	return "fabs";	×
199	case data_flow::TaskletCode::fabsf:
200	return "fabsf";	×
201	case data_flow::TaskletCode::fabsl:
202	return "fabsl";	×
203	case data_flow::TaskletCode::floor:
204	return "floor";	×
205	case data_flow::TaskletCode::floorf:
206	return "floorf";	×
207	case data_flow::TaskletCode::floorl:
208	return "floorl";	×
209	case data_flow::TaskletCode::fls:
210	return "fls";	×
211	case data_flow::TaskletCode::flsl:
212	return "flsl";	×
213	case data_flow::TaskletCode::fmax:
214	return "fmax";	×
215	case data_flow::TaskletCode::fmaxf:
216	return "fmaxf";	×
217	case data_flow::TaskletCode::fmaxl:
218	return "fmaxl";	×
219	case data_flow::TaskletCode::fmin:
220	return "fmin";	×
221	case data_flow::TaskletCode::fminf:
222	return "fminf";	×
223	case data_flow::TaskletCode::fminl:
224	return "fminl";	×
225	case data_flow::TaskletCode::fmod:
226	return "fmod";	×
227	case data_flow::TaskletCode::fmodf:
228	return "fmodf";	×
229	case data_flow::TaskletCode::fmodl:
230	return "fmodl";	×
231	case data_flow::TaskletCode::frexp:
232	return "frexp";	×
233	case data_flow::TaskletCode::frexpf:
234	return "frexpf";	×
235	case data_flow::TaskletCode::frexpl:
236	return "frexpl";	×
237	case data_flow::TaskletCode::labs:
238	return "labs";	×
239	case data_flow::TaskletCode::ldexp:
240	return "ldexp";	×
241	case data_flow::TaskletCode::ldexpf:
242	return "ldexpf";	×
243	case data_flow::TaskletCode::ldexpl:
244	return "ldexpl";	×
245	case data_flow::TaskletCode::log10:
246	return "log10";	×
247	case data_flow::TaskletCode::log10f:
248	return "log10f";	×
249	case data_flow::TaskletCode::log10l:
250	return "log10l";	×
251	case data_flow::TaskletCode::log2:
252	return "log2";	×
253	case data_flow::TaskletCode::log2f:
254	return "log2f";	×
255	case data_flow::TaskletCode::log2l:
256	return "log2l";	×
257	case data_flow::TaskletCode::log:
258	return "log";	×
259	case data_flow::TaskletCode::logf:
260	return "logf";	×
261	case data_flow::TaskletCode::logl:
262	return "logl";	×
263	case data_flow::TaskletCode::logb:
264	return "logb";	×
265	case data_flow::TaskletCode::logbf:
266	return "logbf";	×
267	case data_flow::TaskletCode::logbl:
268	return "logbl";	×
269	case data_flow::TaskletCode::log1p:
270	return "log1p";	×
271	case data_flow::TaskletCode::log1pf:
272	return "log1pf";	×
273	case data_flow::TaskletCode::log1pl:
274	return "log1pl";	×
275	case data_flow::TaskletCode::modf:
276	return "modf";	×
277	case data_flow::TaskletCode::modff:
278	return "modff";	×
279	case data_flow::TaskletCode::modfl:
280	return "modfl";	×
281	case data_flow::TaskletCode::nearbyint:
282	return "nearbyint";	×
283	case data_flow::TaskletCode::nearbyintf:
284	return "nearbyintf";	×
285	case data_flow::TaskletCode::nearbyintl:
286	return "nearbyintl";	×
287	case data_flow::TaskletCode::pow:
288	return "pow";	×
289	case data_flow::TaskletCode::powf:
290	return "powf";	×
291	case data_flow::TaskletCode::powl:
292	return "powl";	×
293	case data_flow::TaskletCode::rint:
294	return "rint";	×
295	case data_flow::TaskletCode::rintf:
296	return "rintf";	×
297	case data_flow::TaskletCode::rintl:
298	return "rintl";	×
299	case data_flow::TaskletCode::lrint:
300	return "lrint";	×
301	case data_flow::TaskletCode::llrint:
302	return "llrint";	×
303	case data_flow::TaskletCode::lround:
304	return "lround";	×
305	case data_flow::TaskletCode::llround:
306	return "llround";	×
307	case data_flow::TaskletCode::round:
308	return "round";	×
309	case data_flow::TaskletCode::roundf:
310	return "roundf";	×
311	case data_flow::TaskletCode::roundl:
312	return "roundl";	×
313	case data_flow::TaskletCode::roundeven:
314	return "roundeven";	×
315	case data_flow::TaskletCode::roundevenf:
316	return "roundevenf";	×
317	case data_flow::TaskletCode::roundevenl:
318	return "roundevenl";	×
319	case data_flow::TaskletCode::sin:
320	return "sin";	×
321	case data_flow::TaskletCode::sinf:
322	return "sinf";	×
323	case data_flow::TaskletCode::sinl:
324	return "sinl";	×
325	case data_flow::TaskletCode::sinh:
326	return "sinh";	×
327	case data_flow::TaskletCode::sinhf:
328	return "sinhf";	×
329	case data_flow::TaskletCode::sinhl:
330	return "sinhl";	×
331	case data_flow::TaskletCode::sqrt:
332	return "sqrt";	×
333	case data_flow::TaskletCode::sqrtf:
334	return "sqrtf";	×
335	case data_flow::TaskletCode::sqrtl:
336	return "sqrtl";	×
337	case data_flow::TaskletCode::rsqrt:
338	return "rsqrt";	×
339	case data_flow::TaskletCode::rsqrtf:
340	return "rsqrtf";	×
341	case data_flow::TaskletCode::rsqrtl:
342	return "rsqrtl";	×
343	case data_flow::TaskletCode::tan:
344	return "tan";	×
345	case data_flow::TaskletCode::tanf:
346	return "tanf";	×
347	case data_flow::TaskletCode::tanl:
348	return "tanl";	×
349	case data_flow::TaskletCode::tanh:
350	return "tanh";	×
351	case data_flow::TaskletCode::tanhf:
352	return "tanhf";	×
353	case data_flow::TaskletCode::tanhl:
354	return "tanhl";	×
355	};
356	throw std::invalid_argument("Invalid tasklet code");	×
357	};	×
358
359	std::string CUDALanguageExtension::primitive_type(const types::PrimitiveType prim_type) {	20✔
360	switch (prim_type) {	20✔
361	case types::PrimitiveType::Void:
362	return "void";	1✔
363	case types::PrimitiveType::Bool:
364	return "bool";	1✔
365	case types::PrimitiveType::Int8:
366	return "signed char";	1✔
367	case types::PrimitiveType::Int16:
368	return "short";	1✔
369	case types::PrimitiveType::Int32:
370	return "int";	6✔
371	case types::PrimitiveType::Int64:
372	return "long long";	1✔
373	case types::PrimitiveType::Int128:
374	return "__int128";	×
375	case types::PrimitiveType::UInt8:
376	return "char";	3✔
377	case types::PrimitiveType::UInt16:
378	return "unsigned short";	1✔
379	case types::PrimitiveType::UInt32:
380	return "unsigned int";	1✔
381	case types::PrimitiveType::UInt64:
382	return "unsigned long long";	1✔
383	case types::PrimitiveType::UInt128:
384	return "unsigned __int128";	×
385	case types::PrimitiveType::Half:
386	return "__fp16";	×
387	case types::PrimitiveType::BFloat:
388	return "__bf16";	×
389	case types::PrimitiveType::Float:
390	return "float";	2✔
391	case types::PrimitiveType::Double:
392	return "double";	1✔
393	case types::PrimitiveType::X86_FP80:
394	return "long double";	×
395	case types::PrimitiveType::FP128:
396	return "__float128";	×
397	case types::PrimitiveType::PPC_FP128:
398	return "__float128";	×
399	}
400
401	throw std::runtime_error("Unknown primitive type");	×
402	};	20✔
403
404	std::string CUDALanguageExtension::
405	declaration(const std::string& name, const types::IType& type, bool use_initializer, bool use_alignment) {	18✔
406	std::stringstream val;	18✔
407
408	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	18✔
409	if (scalar_type->storage_type() == types::StorageType_NV_Shared) {	8✔
410	val << "__shared__ ";	×
411	} else if (scalar_type->storage_type() == types::StorageType_NV_Constant) {	8✔
412	val << "__constant__ ";	×
413	}	×
414	val << primitive_type(scalar_type->primitive_type());	8✔
415	val << " ";	8✔
416	val << name;	8✔
417	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	18✔
418	auto& element_type = array_type->element_type();	3✔
419	val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]", element_type);	3✔
420	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	10✔
421	if (pointer_type->has_pointee_type()) {	4✔
422	const types::IType& pointee = pointer_type->pointee_type();	3✔
423
424	const bool pointee_is_function_or_array = dynamic_cast<const types::Function*>(&pointee) \|\|	6✔
425	dynamic_cast<const types::Array*>(&pointee);	3✔
426
427	// Parenthesise only when it is needed to bind tighter than [] or ()
428	std::string decorated = pointee_is_function_or_array ? "(" + name + ")" : "" + name;	3✔
429
430	val << declaration(decorated, pointee);	3✔
431	} else {	3✔
432	val << "void*";	1✔
433	val << " " << name;	1✔
434	}
435	} else if (auto ref_type = dynamic_cast<const Reference*>(&type)) {	7✔
436	val << declaration("&" + name, ref_type->reference_type());	×
437	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	3✔
438	if (structure_type->storage_type() == types::StorageType_NV_Shared) {	3✔
439	val << "__shared__ ";	×
440	} else if (structure_type->storage_type() == types::StorageType_NV_Constant) {	3✔
441	val << "__constant__ ";	×
442	}	×
443	val << structure_type->name();	3✔
444	val << " ";	3✔
445	val << name;	3✔
446	} else if (auto function_type = dynamic_cast<const types::Function*>(&type)) {	3✔
447	std::stringstream params;	×
448	for (size_t i = 0; i < function_type->num_params(); ++i) {	×
449	params << declaration("", function_type->param_type(symbolic::integer(i)));	×
450	if (i + 1 < function_type->num_params()) params << ", ";	×
451	}	×
452	if (function_type->is_var_arg()) {	×
453	// ISO C++ forbids empty parameter lists before ...
NEW 454	if (function_type->num_params() > 0) {	×
NEW 455	params << ", ";	×
NEW 456	params << "...";	×
NEW 457	}	×
UNCOV 458	}	×
459
460	const std::string fun_name = name + "(" + params.str() + ")";	×
461	val << declaration(fun_name, function_type->return_type());	×
462	} else {	×
463	throw std::runtime_error("Unknown declaration type");	×
464	}
465
466	if (use_alignment && type.alignment() > 0) {	18✔
467	val << " __attribute__((aligned(" << type.alignment() << ")))";	×
468	}	×
469
470	if (use_initializer && !type.initializer().empty()) {	18✔
471	val << " = " << type.initializer();	×
472	}	×
473
474	return val.str();	18✔
475	};	18✔
476
477	std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {	1✔
478	std::stringstream val;	1✔
479
480	val << "reinterpret_cast";	1✔
481	val << "<";	1✔
482	val << declaration("", type);	1✔
483	val << ">";	1✔
484	val << "(" << name << ")";	1✔
485
486	return val.str();	1✔
487	};	1✔
488
489	std::string CUDALanguageExtension::subset(const Function& function, const types::IType& type, const data_flow::Subset& sub) {	3✔
490	if (sub.empty()) {	3✔
491	return "";	1✔
492	}
493
494	if (dynamic_cast<const types::Scalar*>(&type)) {	2✔
495	return "";	×
496	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	2✔
497	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	1✔
498
499	if (sub.size() > 1) {	1✔
500	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
501	auto& element_type = array_type->element_type();	×
502	return subset_str + subset(function, element_type, element_subset);	×
503	} else {	×
504	return subset_str;	1✔
505	}
506	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	2✔
507	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	×
508
509	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
510	auto& pointee_type = pointer_type->pointee_type();	×
511	return subset_str + subset(function, pointee_type, element_subset);	×
512	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	1✔
513	auto& definition = function.structure(structure_type->name());	1✔
514
515	std::string subset_str = ".member_" + this->expression(sub.at(0));	1✔
516	if (sub.size() > 1) {	1✔
517	auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));	×
518	auto& member_type = definition.member_type(member);	×
519	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
520	return subset_str + subset(function, member_type, element_subset);	×
521	} else {	×
522	return subset_str;	1✔
523	}
524	}	1✔
525
526	throw std::invalid_argument("Invalid subset type");	×
527	};	3✔
528
529	std::string CUDALanguageExtension::expression(const symbolic::Expression& expr) {	5✔
530	CPPSymbolicPrinter printer;	5✔
531	return printer.apply(expr);	5✔
532	};	5✔
533
534	std::string CUDALanguageExtension::access_node(const data_flow::AccessNode& node) {	×
NEW 535	if (dynamic_cast<const data_flow::ConstantNode*>(&node)) {	×
NEW 536	std::string name = node.data();	×
NEW 537	if (symbolic::is_nullptr(symbolic::symbol(name))) {	×
NEW 538	return this->expression(symbolic::__nullptr__());	×
539	}
NEW 540	return name;	×
NEW 541	} else {	×
NEW 542	std::string name = node.data();	×
NEW 543	if (this->external_variables_.find(name) != this->external_variables_.end()) {	×
NEW 544	return "(&" + name + ")";	×
545	}
NEW 546	return name;	×
547	}	×
548	};	×
549
550	std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {	×
551	std::string op = code_to_string(tasklet.code());	×
552	std::vector<std::string> arguments;	×
553	for (size_t i = 0; i < tasklet.inputs().size(); ++i) {	×
554	arguments.push_back(tasklet.input(i));	×
555	}	×
556
557	if (tasklet.code() == data_flow::TaskletCode::assign) {	×
558	return arguments.at(0);	×
559	} else if (data_flow::is_infix(tasklet.code())) {	×
560	switch (data_flow::arity(tasklet.code())) {	×
561	case 1:
562	return op + arguments.at(0);	×
563	case 2:
564	return arguments.at(0) + " " + op + " " + arguments.at(1);	×
565	default:
566	throw std::runtime_error("Unsupported arity");	×
567	}
568	} else {
569	return op + "(" + helpers::join(arguments, ", ") + ")";	×
570	}
571	};	×
572
573	std::string CUDALanguageExtension::zero(const types::PrimitiveType prim_type) {	×
574	switch (prim_type) {	×
575	case types::Void:
576	throw InvalidSDFGException("No zero for void type possible");	×
577	case types::Bool:
578	return "false";	×
579	case types::Int8:
580	return "0";	×
581	case types::Int16:
582	return "0";	×
583	case types::Int32:
584	return "0";	×
585	case types::Int64:
586	return "0ll";	×
587	case types::Int128:
588	return "0";	×
589	case types::UInt8:
590	return "0u";	×
591	case types::UInt16:
592	return "0u";	×
593	case types::UInt32:
594	return "0u";	×
595	case types::UInt64:
596	return "0ull";	×
597	case types::UInt128:
598	return "0";	×
599	case types::Half:
600	return "CUDART_ZERO_FP16";	×
601	case types::BFloat:
602	return "CUDART_ZERO_BF16";	×
603	case types::Float:
604	return "0.0f";	×
605	case types::Double:
606	return "0.0";	×
607	case types::X86_FP80:
608	return "0.0l";	×
609	case types::FP128:
610	return "0.0";	×
611	case types::PPC_FP128:
612	return "0.0";	×
613	}	×
614	}	×
615
616	} // namespace codegen
617	} // namespace sdfg

daisytuner / sdfglib / 17656823807

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous