15470021361

Committed 05 Jun 2025 02:42PM UTC coverage: 57.727% (-0.2%) from 57.883%

Build # 15470021361

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #57 from daisytuner/alignments

adds explicit alignment to types

Run Details

44 of 61 new or added lines in 12 files covered. (72.13%)

7 existing lines in 5 files now uncovered.

7953 of 13777 relevant lines covered (57.73%)

108.58 hits per line

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

24.53

/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::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::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 (type.address_space() == 3) {
            val << "__shared__ ";
        } else if (type.address_space() == 4) {
            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)) {
        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 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 (type.address_space() == 3) {
            val << "__shared__ ";
        } else if (type.address_space() == 4) {
            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()) {
            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::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) {
        std::string arg = tasklet.input(i).first;
        if (!tasklet.needs_connector(i)) {
            if (arg != "NAN" && arg != "INFINITY") {
                if (tasklet.input(i).second.primitive_type() == types::PrimitiveType::Float) {
                    arg += "f";
                }
            }
        }
        arguments.push_back(arg);
    }

    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::library_node(const data_flow::LibraryNode& libnode) {
    data_flow::LibraryNodeType lib_node_type = libnode.call();
    switch (lib_node_type) {
        case sdfg::data_flow::LibraryNodeType::LocalBarrier:
            return "__syncthreads();";
        default:
            throw std::runtime_error("Unsupported library node type");
    }
}

}  // 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::exp10:
168	return "exp10";	×
169	case data_flow::TaskletCode::exp10f:
170	return "exp10f";	×
171	case data_flow::TaskletCode::exp10l:
172	return "exp10l";	×
173	case data_flow::TaskletCode::exp2:
174	return "exp2";	×
175	case data_flow::TaskletCode::exp2f:
176	return "exp2f";	×
177	case data_flow::TaskletCode::exp2l:
178	return "exp2l";	×
179	case data_flow::TaskletCode::exp:
180	return "exp";	×
181	case data_flow::TaskletCode::expf:
182	return "expf";	×
183	case data_flow::TaskletCode::expl:
184	return "expl";	×
185	case data_flow::TaskletCode::expm1:
186	return "expm1";	×
187	case data_flow::TaskletCode::expm1f:
188	return "expm1f";	×
189	case data_flow::TaskletCode::expm1l:
190	return "expm1l";	×
191	case data_flow::TaskletCode::fabs:
192	return "fabs";	×
193	case data_flow::TaskletCode::fabsf:
194	return "fabsf";	×
195	case data_flow::TaskletCode::fabsl:
196	return "fabsl";	×
197	case data_flow::TaskletCode::floor:
198	return "floor";	×
199	case data_flow::TaskletCode::floorf:
200	return "floorf";	×
201	case data_flow::TaskletCode::floorl:
202	return "floorl";	×
203	case data_flow::TaskletCode::fls:
204	return "fls";	×
205	case data_flow::TaskletCode::flsl:
206	return "flsl";	×
207	case data_flow::TaskletCode::fmax:
208	return "fmax";	×
209	case data_flow::TaskletCode::fmaxf:
210	return "fmaxf";	×
211	case data_flow::TaskletCode::fmaxl:
212	return "fmaxl";	×
213	case data_flow::TaskletCode::fmin:
214	return "fmin";	×
215	case data_flow::TaskletCode::fminf:
216	return "fminf";	×
217	case data_flow::TaskletCode::fminl:
218	return "fminl";	×
219	case data_flow::TaskletCode::fmod:
220	return "fmod";	×
221	case data_flow::TaskletCode::fmodf:
222	return "fmodf";	×
223	case data_flow::TaskletCode::fmodl:
224	return "fmodl";	×
225	case data_flow::TaskletCode::frexp:
226	return "frexp";	×
227	case data_flow::TaskletCode::frexpf:
228	return "frexpf";	×
229	case data_flow::TaskletCode::frexpl:
230	return "frexpl";	×
231	case data_flow::TaskletCode::labs:
232	return "labs";	×
233	case data_flow::TaskletCode::ldexp:
234	return "ldexp";	×
235	case data_flow::TaskletCode::ldexpf:
236	return "ldexpf";	×
237	case data_flow::TaskletCode::ldexpl:
238	return "ldexpl";	×
239	case data_flow::TaskletCode::log10:
240	return "log10";	×
241	case data_flow::TaskletCode::log10f:
242	return "log10f";	×
243	case data_flow::TaskletCode::log10l:
244	return "log10l";	×
245	case data_flow::TaskletCode::log2:
246	return "log2";	×
247	case data_flow::TaskletCode::log2f:
248	return "log2f";	×
249	case data_flow::TaskletCode::log2l:
250	return "log2l";	×
251	case data_flow::TaskletCode::log:
252	return "log";	×
253	case data_flow::TaskletCode::logf:
254	return "logf";	×
255	case data_flow::TaskletCode::logl:
256	return "logl";	×
257	case data_flow::TaskletCode::logb:
258	return "logb";	×
259	case data_flow::TaskletCode::logbf:
260	return "logbf";	×
261	case data_flow::TaskletCode::logbl:
262	return "logbl";	×
263	case data_flow::TaskletCode::log1p:
264	return "log1p";	×
265	case data_flow::TaskletCode::log1pf:
266	return "log1pf";	×
267	case data_flow::TaskletCode::log1pl:
268	return "log1pl";	×
269	case data_flow::TaskletCode::modf:
270	return "modf";	×
271	case data_flow::TaskletCode::modff:
272	return "modff";	×
273	case data_flow::TaskletCode::modfl:
274	return "modfl";	×
275	case data_flow::TaskletCode::nearbyint:
276	return "nearbyint";	×
277	case data_flow::TaskletCode::nearbyintf:
278	return "nearbyintf";	×
279	case data_flow::TaskletCode::nearbyintl:
280	return "nearbyintl";	×
281	case data_flow::TaskletCode::pow:
282	return "pow";	×
283	case data_flow::TaskletCode::powf:
284	return "powf";	×
285	case data_flow::TaskletCode::powl:
286	return "powl";	×
287	case data_flow::TaskletCode::rint:
288	return "rint";	×
289	case data_flow::TaskletCode::rintf:
290	return "rintf";	×
291	case data_flow::TaskletCode::rintl:
292	return "rintl";	×
293	case data_flow::TaskletCode::round:
294	return "round";	×
295	case data_flow::TaskletCode::roundf:
296	return "roundf";	×
297	case data_flow::TaskletCode::roundl:
298	return "roundl";	×
299	case data_flow::TaskletCode::roundeven:
300	return "roundeven";	×
301	case data_flow::TaskletCode::roundevenf:
302	return "roundevenf";	×
303	case data_flow::TaskletCode::roundevenl:
304	return "roundevenl";	×
305	case data_flow::TaskletCode::sin:
306	return "sin";	×
307	case data_flow::TaskletCode::sinf:
308	return "sinf";	×
309	case data_flow::TaskletCode::sinl:
310	return "sinl";	×
311	case data_flow::TaskletCode::sinh:
312	return "sinh";	×
313	case data_flow::TaskletCode::sinhf:
314	return "sinhf";	×
315	case data_flow::TaskletCode::sinhl:
316	return "sinhl";	×
317	case data_flow::TaskletCode::sqrt:
318	return "sqrt";	×
319	case data_flow::TaskletCode::sqrtf:
320	return "sqrtf";	×
321	case data_flow::TaskletCode::sqrtl:
322	return "sqrtl";	×
323	case data_flow::TaskletCode::rsqrt:
324	return "rsqrt";	×
325	case data_flow::TaskletCode::rsqrtf:
326	return "rsqrtf";	×
327	case data_flow::TaskletCode::rsqrtl:
328	return "rsqrtl";	×
329	case data_flow::TaskletCode::tan:
330	return "tan";	×
331	case data_flow::TaskletCode::tanf:
332	return "tanf";	×
333	case data_flow::TaskletCode::tanl:
334	return "tanl";	×
335	case data_flow::TaskletCode::tanh:
336	return "tanh";	×
337	case data_flow::TaskletCode::tanhf:
338	return "tanhf";	×
339	case data_flow::TaskletCode::tanhl:
340	return "tanhl";	×
341	};
342	throw std::invalid_argument("Invalid tasklet code");	×
343	};	×
344
345	std::string CUDALanguageExtension::primitive_type(const types::PrimitiveType prim_type) {	20✔
346	switch (prim_type) {	20✔
347	case types::PrimitiveType::Void:
348	return "void";	1✔
349	case types::PrimitiveType::Bool:
350	return "bool";	1✔
351	case types::PrimitiveType::Int8:
352	return "signed char";	1✔
353	case types::PrimitiveType::Int16:
354	return "short";	1✔
355	case types::PrimitiveType::Int32:
356	return "int";	6✔
357	case types::PrimitiveType::Int64:
358	return "long long";	1✔
359	case types::PrimitiveType::Int128:
360	return "__int128";	×
361	case types::PrimitiveType::UInt8:
362	return "char";	3✔
363	case types::PrimitiveType::UInt16:
364	return "unsigned short";	1✔
365	case types::PrimitiveType::UInt32:
366	return "unsigned int";	1✔
367	case types::PrimitiveType::UInt64:
368	return "unsigned long long";	1✔
369	case types::PrimitiveType::UInt128:
370	return "unsigned __int128";	×
371	case types::PrimitiveType::Half:
372	return "__fp16";	×
373	case types::PrimitiveType::BFloat:
374	return "__bf16";	×
375	case types::PrimitiveType::Float:
376	return "float";	2✔
377	case types::PrimitiveType::Double:
378	return "double";	1✔
379	case types::PrimitiveType::X86_FP80:
380	return "long double";	×
381	case types::PrimitiveType::FP128:
382	return "__float128";	×
383	case types::PrimitiveType::PPC_FP128:
384	return "__float128";	×
385	}
386
387	throw std::runtime_error("Unknown primitive type");	×
388	};	20✔
389
390	std::string CUDALanguageExtension::declaration(const std::string& name, const types::IType& type,	17✔
391	bool use_initializer, bool use_alignment) {
392	std::stringstream val;	17✔
393
394	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	17✔
395	if (type.address_space() == 3) {	8✔
396	val << "__shared__ ";	×
397	} else if (type.address_space() == 4) {	8✔
398	val << "__constant__ ";	×
399	}	×
400	val << primitive_type(scalar_type->primitive_type());	8✔
401	val << " ";	8✔
402	val << name;	8✔
403	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	17✔
404	auto& element_type = array_type->element_type();	3✔
405	val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]",	3✔
406	element_type);	3✔
407	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	9✔
408	const types::IType& pointee = pointer_type->pointee_type();	3✔
409
410	const bool pointee_is_function_or_array = dynamic_cast<const types::Function*>(&pointee) \|\|	6✔
411	dynamic_cast<const types::Array*>(&pointee);	3✔
412
413	// Parenthesise only when it is needed to bind tighter than [] or ()
414	std::string decorated = pointee_is_function_or_array ? "(" + name + ")" : "" + name;	3✔
415
416	val << declaration(decorated, pointee);	3✔
417	} else if (auto ref_type = dynamic_cast<const Reference*>(&type)) {	6✔
418	val << declaration("&" + name, ref_type->reference_type());	×
419	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	3✔
420	if (type.address_space() == 3) {	3✔
421	val << "__shared__ ";	×
422	} else if (type.address_space() == 4) {	3✔
423	val << "__constant__ ";	×
424	}	×
425	val << structure_type->name();	3✔
426	val << " ";	3✔
427	val << name;	3✔
428	} else if (auto function_type = dynamic_cast<const types::Function*>(&type)) {	3✔
429	std::stringstream params;	×
430	for (size_t i = 0; i < function_type->num_params(); ++i) {	×
431	params << declaration("", function_type->param_type(symbolic::integer(i)));	×
432	if (i + 1 < function_type->num_params()) params << ", ";	×
433	}	×
434	if (function_type->is_var_arg()) {	×
435	if (function_type->num_params() > 0) params << ", ";	×
436	params << "...";	×
437	}	×
438
439	const std::string fun_name = name + "(" + params.str() + ")";	×
440	val << declaration(fun_name, function_type->return_type());	×
441	} else {	×
442	throw std::runtime_error("Unknown declaration type");	×
443	}
444
445	if (use_alignment && type.alignment() > 0) {	17✔
NEW 446	val << " __attribute__((aligned(" << type.alignment() << ")))";	×
UNCOV 447	}	×
448
449	if (use_initializer && !type.initializer().empty()) {	17✔
NEW 450	val << " = " << type.initializer();	×
UNCOV 451	}	×
452
453	return val.str();	17✔
454	};	17✔
455
456	std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {	1✔
457	std::stringstream val;	1✔
458
459	val << "reinterpret_cast";	1✔
460	val << "<";	1✔
461	val << declaration("", type);	1✔
462	val << ">";	1✔
463	val << "(" << name << ")";	1✔
464
465	return val.str();	1✔
466	};	1✔
467
468	std::string CUDALanguageExtension::subset(const Function& function, const types::IType& type,	3✔
469	const data_flow::Subset& sub) {
470	if (sub.empty()) {	3✔
471	return "";	1✔
472	}
473
474	if (dynamic_cast<const types::Scalar*>(&type)) {	2✔
475	return "";	×
476	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	2✔
477	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	1✔
478
479	if (sub.size() > 1) {	1✔
480	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
481	auto& element_type = array_type->element_type();	×
482	return subset_str + subset(function, element_type, element_subset);	×
483	} else {	×
484	return subset_str;	1✔
485	}
486	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	2✔
487	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	×
488
489	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
490	auto& pointee_type = pointer_type->pointee_type();	×
491	return subset_str + subset(function, pointee_type, element_subset);	×
492	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	1✔
493	auto& definition = function.structure(structure_type->name());	1✔
494
495	std::string subset_str = ".member_" + this->expression(sub.at(0));	1✔
496	if (sub.size() > 1) {	1✔
497	auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));	×
498	auto& member_type = definition.member_type(member);	×
499	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
500	return subset_str + subset(function, member_type, element_subset);	×
501	} else {	×
502	return subset_str;	1✔
503	}
504	}	1✔
505
506	throw std::invalid_argument("Invalid subset type");	×
507	};	3✔
508
509	std::string CUDALanguageExtension::expression(const symbolic::Expression& expr) {	9✔
510	CPPSymbolicPrinter printer;	9✔
511	return printer.apply(expr);	9✔
512	};	9✔
513
514	std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {	×
515	std::string op = code_to_string(tasklet.code());	×
516	std::vector<std::string> arguments;	×
517	for (size_t i = 0; i < tasklet.inputs().size(); ++i) {	×
518	std::string arg = tasklet.input(i).first;	×
519	if (!tasklet.needs_connector(i)) {	×
520	if (arg != "NAN" && arg != "INFINITY") {	×
521	if (tasklet.input(i).second.primitive_type() == types::PrimitiveType::Float) {	×
522	arg += "f";	×
523	}	×
524	}	×
525	}	×
526	arguments.push_back(arg);	×
527	}	×
528
529	if (tasklet.code() == data_flow::TaskletCode::assign) {	×
530	return arguments.at(0);	×
531	} else if (data_flow::is_infix(tasklet.code())) {	×
532	switch (data_flow::arity(tasklet.code())) {	×
533	case 1:
534	return op + arguments.at(0);	×
535	case 2:
536	return arguments.at(0) + " " + op + " " + arguments.at(1);	×
537	default:
538	throw std::runtime_error("Unsupported arity");	×
539	}
540	} else {
541	return op + "(" + helpers::join(arguments, ", ") + ")";	×
542	}
543	};	×
544
545	std::string CUDALanguageExtension::library_node(const data_flow::LibraryNode& libnode) {	2✔
546	data_flow::LibraryNodeType lib_node_type = libnode.call();	2✔
547	switch (lib_node_type) {	2✔
548	case sdfg::data_flow::LibraryNodeType::LocalBarrier:
549	return "__syncthreads();";	2✔
550	default:
551	throw std::runtime_error("Unsupported library node type");	×
552	}
553	}	×
554
555	} // namespace codegen
556	} // namespace sdfg

daisytuner / sdfglib / 15470021361

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