15044057891

Committed 15 May 2025 11:42AM UTC coverage: 59.37% (+1.8%) from 57.525%

Build # 15044057891

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #14 from daisytuner/sanitizers

enables sanitizer on unit tests

Run Details

63 of 67 new or added lines in 47 files covered. (94.03%)

570 existing lines in 62 files now uncovered.

7356 of 12390 relevant lines covered (59.37%)

505.93 hits per line

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

28.66

/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";
    };
    assert(false);
};

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::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::Float:
            return "float";
        case types::PrimitiveType::Double:
            return "double";
    }

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

std::string CUDALanguageExtension::declaration(const std::string& name, const types::IType& type,
                                               bool use_initializer) {
    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)) {
        auto& pointee_type = pointer_type->pointee_type();
        val << declaration("(*" + name + ")", pointee_type);
    } 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)) {
        val << declaration("", function_type->return_type());
        val << " ";
        val << name;
        val << "(";
        for (size_t i = 0; i < function_type->num_params(); ++i) {
            val << declaration("", function_type->param_type(symbolic::integer(i)));
            if (i < function_type->num_params() - 1) {
                val << ", ";
            }
        }
        if (function_type->is_var_arg()) {
            if (function_type->num_params() > 0) {
                val << ", ";
            }
            val << "...";
        }
        val << ")";
    } else {
        throw std::runtime_error("Unknown declaration type");
    }

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

    return val.str();
};

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

    if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {
        val << declaration(name, *scalar_type);
    } else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {
        val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]",
                           array_type->element_type());
    } else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {
        std::string pointee_name = name + "__daisy_nvptx_internal_";
        val << declaration(pointee_name, pointer_type->pointee_type());
        val << ";";
        val << std::endl;

        val << declaration(name, type);
        val << " = ";
        val << "&" << pointee_name;
    } else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {
        val << declaration(name, *structure_type);
    } else {
        throw std::runtime_error("Unknown allocation type");
    }

    return val.str();
};

std::string CUDALanguageExtension::deallocation(const std::string& name, const types::IType& type) {
    return "";
};

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 (auto scalar_type = 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;
        }
    } else {
        assert(false);
    }
};

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	assert(false);	×
UNCOV 343	};	×
344
345	std::string CUDALanguageExtension::primitive_type(const types::PrimitiveType prim_type) {	24✔
346	switch (prim_type) {	24✔
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";	8✔
357	case types::PrimitiveType::Int64:
358	return "long long";	1✔
359	case types::PrimitiveType::UInt8:
360	return "char";	3✔
361	case types::PrimitiveType::UInt16:
362	return "unsigned short";	1✔
363	case types::PrimitiveType::UInt32:
364	return "unsigned int";	1✔
365	case types::PrimitiveType::UInt64:
366	return "unsigned long long";	1✔
367	case types::PrimitiveType::Float:
368	return "float";	4✔
369	case types::PrimitiveType::Double:
370	return "double";	1✔
371	}
372
373	throw std::runtime_error("Unknown primitive type");	×
374	};	24✔
375
376	std::string CUDALanguageExtension::declaration(const std::string& name, const types::IType& type,	23✔
377	bool use_initializer) {
378	std::stringstream val;	23✔
379
380	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	23✔
381	if (type.address_space() == 3) {	12✔
382	val << "__shared__ ";	×
383	} else if (type.address_space() == 4) {	12✔
384	val << "__constant__ ";	×
UNCOV 385	}	×
386	val << primitive_type(scalar_type->primitive_type());	12✔
387	val << " ";	12✔
388	val << name;	12✔
389	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	23✔
390	auto& element_type = array_type->element_type();	3✔
391	val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]",	3✔
392	element_type);	3✔
393	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	11✔
394	auto& pointee_type = pointer_type->pointee_type();	4✔
395	val << declaration("(*" + name + ")", pointee_type);	4✔
396	} else if (auto ref_type = dynamic_cast<const Reference*>(&type)) {	8✔
397	val << declaration("&" + name, ref_type->reference_type());	×
398	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	4✔
399	if (type.address_space() == 3) {	4✔
400	val << "__shared__ ";	×
401	} else if (type.address_space() == 4) {	4✔
402	val << "__constant__ ";	×
UNCOV 403	}	×
404	val << structure_type->name();	4✔
405	val << " ";	4✔
406	val << name;	4✔
407	} else if (auto function_type = dynamic_cast<const types::Function*>(&type)) {	4✔
408	val << declaration("", function_type->return_type());	×
409	val << " ";	×
410	val << name;	×
411	val << "(";	×
412	for (size_t i = 0; i < function_type->num_params(); ++i) {	×
413	val << declaration("", function_type->param_type(symbolic::integer(i)));	×
414	if (i < function_type->num_params() - 1) {	×
415	val << ", ";	×
UNCOV 416	}	×
UNCOV 417	}	×
418	if (function_type->is_var_arg()) {	×
419	if (function_type->num_params() > 0) {	×
420	val << ", ";	×
UNCOV 421	}	×
422	val << "...";	×
UNCOV 423	}	×
424	val << ")";	×
UNCOV 425	} else {	×
426	throw std::runtime_error("Unknown declaration type");	×
427	}
428
429	if (use_initializer && !type.initializer().empty()) {	23✔
430	val << " = " << type.initializer();	×
UNCOV 431	}	×
432
433	return val.str();	23✔
434	};	23✔
435
436	std::string CUDALanguageExtension::allocation(const std::string& name, const types::IType& type) {	4✔
437	std::stringstream val;	4✔
438
439	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	4✔
440	val << declaration(name, *scalar_type);	1✔
441	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	4✔
442	val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]",	2✔
443	array_type->element_type());	1✔
444	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	3✔
445	std::string pointee_name = name + "__daisy_nvptx_internal_";	1✔
446	val << declaration(pointee_name, pointer_type->pointee_type());	1✔
447	val << ";";	1✔
448	val << std::endl;	1✔
449
450	val << declaration(name, type);	1✔
451	val << " = ";	1✔
452	val << "&" << pointee_name;	1✔
453	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	2✔
454	val << declaration(name, *structure_type);	1✔
455	} else {	1✔
456	throw std::runtime_error("Unknown allocation type");	×
457	}
458
459	return val.str();	4✔
460	};	4✔
461
462	std::string CUDALanguageExtension::deallocation(const std::string& name, const types::IType& type) {	4✔
463	return "";	4✔
UNCOV 464	};	×
465
466	std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {	1✔
467	std::stringstream val;	1✔
468
469	val << "reinterpret_cast";	1✔
470	val << "<";	1✔
471	val << declaration("", type);	1✔
472	val << ">";	1✔
473	val << "(" << name << ")";	1✔
474
475	return val.str();	1✔
476	};	1✔
477
478	std::string CUDALanguageExtension::subset(const Function& function, const types::IType& type,	3✔
479	const data_flow::Subset& sub) {
480	if (sub.empty()) {	3✔
481	return "";	1✔
482	}
483
484	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	2✔
485	return "";	×
486	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	2✔
487	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	1✔
488
489	if (sub.size() > 1) {	1✔
490	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
491	auto& element_type = array_type->element_type();	×
492	return subset_str + subset(function, element_type, element_subset);	×
493	} else {	×
494	return subset_str;	1✔
495	}
496	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	2✔
497	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	×
498
499	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
500	auto& pointee_type = pointer_type->pointee_type();	×
501	return subset_str + subset(function, pointee_type, element_subset);	×
502	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	1✔
503	auto& definition = function.structure(structure_type->name());	1✔
504
505	std::string subset_str = ".member_" + this->expression(sub.at(0));	1✔
506	if (sub.size() > 1) {	1✔
507	auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));	×
508	auto& member_type = definition.member_type(member);	×
509	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
510	return subset_str + subset(function, member_type, element_subset);	×
511	} else {	×
512	return subset_str;	1✔
513	}
514	} else {	1✔
515	assert(false);	×
516	}
517	};	3✔
518
519	std::string CUDALanguageExtension::expression(const symbolic::Expression& expr) {	10✔
520	CPPSymbolicPrinter printer;	10✔
521	return printer.apply(expr);	10✔
522	};	10✔
523
524	std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {	×
525	std::string op = code_to_string(tasklet.code());	×
526	std::vector<std::string> arguments;	×
527	for (size_t i = 0; i < tasklet.inputs().size(); ++i) {	×
528	std::string arg = tasklet.input(i).first;	×
529	if (!tasklet.needs_connector(i)) {	×
530	if (arg != "NAN" && arg != "INFINITY") {	×
531	if (tasklet.input(i).second.primitive_type() == types::PrimitiveType::Float) {	×
532	arg += "f";	×
UNCOV 533	}	×
UNCOV 534	}	×
UNCOV 535	}	×
536	arguments.push_back(arg);	×
537	}	×
538
539	if (tasklet.code() == data_flow::TaskletCode::assign) {	×
540	return arguments.at(0);	×
541	} else if (data_flow::is_infix(tasklet.code())) {	×
542	switch (data_flow::arity(tasklet.code())) {	×
543	case 1:
544	return op + arguments.at(0);	×
545	case 2:
546	return arguments.at(0) + " " + op + " " + arguments.at(1);	×
547	default:
548	throw std::runtime_error("Unsupported arity");	×
549	}
550	} else {
551	return op + "(" + helpers::join(arguments, ", ") + ")";	×
552	}
553	};	×
554
555	std::string CUDALanguageExtension::library_node(const data_flow::LibraryNode& libnode) {	2✔
556	data_flow::LibraryNodeType lib_node_type = libnode.call();	2✔
557	switch (lib_node_type) {	2✔
558	case sdfg::data_flow::LibraryNodeType::LocalBarrier:
559	return "__syncthreads();";	2✔
560	default:
561	throw std::runtime_error("Unsupported library node type");	×
562	}
UNCOV 563	}	×
564
565	} // namespace codegen
566	} // namespace sdfg

daisytuner / sdfglib / 15044057891

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