16779684622

Committed 06 Aug 2025 02:21PM UTC coverage: 64.3% (-1.0%) from 65.266%

Build # 16779684622

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #172 from daisytuner/opaque-pointers

Opaque pointers, typed memlets, untyped tasklet connectors

Run Details

330 of 462 new or added lines in 38 files covered. (71.43%)

382 existing lines in 30 files now uncovered.

8865 of 13787 relevant lines covered (64.3%)

116.73 hits per line

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

22.65

/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::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()) {
            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) {
    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::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::lrint:
NEW 294	return "lrint";	×
295	case data_flow::TaskletCode::llrint:
NEW 296	return "llrint";	×
297	case data_flow::TaskletCode::lround:
NEW 298	return "lround";	×
299	case data_flow::TaskletCode::llround:
NEW 300	return "llround";	×
301	case data_flow::TaskletCode::round:
302	return "round";	×
303	case data_flow::TaskletCode::roundf:
304	return "roundf";	×
305	case data_flow::TaskletCode::roundl:
306	return "roundl";	×
307	case data_flow::TaskletCode::roundeven:
308	return "roundeven";	×
309	case data_flow::TaskletCode::roundevenf:
310	return "roundevenf";	×
311	case data_flow::TaskletCode::roundevenl:
312	return "roundevenl";	×
313	case data_flow::TaskletCode::sin:
314	return "sin";	×
315	case data_flow::TaskletCode::sinf:
316	return "sinf";	×
317	case data_flow::TaskletCode::sinl:
318	return "sinl";	×
319	case data_flow::TaskletCode::sinh:
320	return "sinh";	×
321	case data_flow::TaskletCode::sinhf:
322	return "sinhf";	×
323	case data_flow::TaskletCode::sinhl:
324	return "sinhl";	×
325	case data_flow::TaskletCode::sqrt:
326	return "sqrt";	×
327	case data_flow::TaskletCode::sqrtf:
328	return "sqrtf";	×
329	case data_flow::TaskletCode::sqrtl:
330	return "sqrtl";	×
331	case data_flow::TaskletCode::rsqrt:
332	return "rsqrt";	×
333	case data_flow::TaskletCode::rsqrtf:
334	return "rsqrtf";	×
335	case data_flow::TaskletCode::rsqrtl:
336	return "rsqrtl";	×
337	case data_flow::TaskletCode::tan:
338	return "tan";	×
339	case data_flow::TaskletCode::tanf:
340	return "tanf";	×
341	case data_flow::TaskletCode::tanl:
342	return "tanl";	×
343	case data_flow::TaskletCode::tanh:
344	return "tanh";	×
345	case data_flow::TaskletCode::tanhf:
346	return "tanhf";	×
347	case data_flow::TaskletCode::tanhl:
348	return "tanhl";	×
349	};
350	throw std::invalid_argument("Invalid tasklet code");	×
351	};	×
352
353	std::string CUDALanguageExtension::primitive_type(const types::PrimitiveType prim_type) {	20✔
354	switch (prim_type) {	20✔
355	case types::PrimitiveType::Void:
356	return "void";	1✔
357	case types::PrimitiveType::Bool:
358	return "bool";	1✔
359	case types::PrimitiveType::Int8:
360	return "signed char";	1✔
361	case types::PrimitiveType::Int16:
362	return "short";	1✔
363	case types::PrimitiveType::Int32:
364	return "int";	6✔
365	case types::PrimitiveType::Int64:
366	return "long long";	1✔
367	case types::PrimitiveType::Int128:
368	return "__int128";	×
369	case types::PrimitiveType::UInt8:
370	return "char";	3✔
371	case types::PrimitiveType::UInt16:
372	return "unsigned short";	1✔
373	case types::PrimitiveType::UInt32:
374	return "unsigned int";	1✔
375	case types::PrimitiveType::UInt64:
376	return "unsigned long long";	1✔
377	case types::PrimitiveType::UInt128:
378	return "unsigned __int128";	×
379	case types::PrimitiveType::Half:
380	return "__fp16";	×
381	case types::PrimitiveType::BFloat:
382	return "__bf16";	×
383	case types::PrimitiveType::Float:
384	return "float";	2✔
385	case types::PrimitiveType::Double:
386	return "double";	1✔
387	case types::PrimitiveType::X86_FP80:
388	return "long double";	×
389	case types::PrimitiveType::FP128:
390	return "__float128";	×
391	case types::PrimitiveType::PPC_FP128:
392	return "__float128";	×
393	}
394
395	throw std::runtime_error("Unknown primitive type");	×
396	};	20✔
397
398	std::string CUDALanguageExtension::
399	declaration(const std::string& name, const types::IType& type, bool use_initializer, bool use_alignment) {	18✔
400	std::stringstream val;	18✔
401
402	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	18✔
403	if (scalar_type->storage_type() == types::StorageType_NV_Shared) {	8✔
404	val << "__shared__ ";	×
405	} else if (scalar_type->storage_type() == types::StorageType_NV_Constant) {	8✔
406	val << "__constant__ ";	×
407	}	×
408	val << primitive_type(scalar_type->primitive_type());	8✔
409	val << " ";	8✔
410	val << name;	8✔
411	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	18✔
412	auto& element_type = array_type->element_type();	3✔
413	val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]", element_type);	3✔
414	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	10✔
415	if (pointer_type->has_pointee_type()) {	4✔
416	const types::IType& pointee = pointer_type->pointee_type();	3✔
417
418	const bool pointee_is_function_or_array = dynamic_cast<const types::Function*>(&pointee) \|\|	6✔
419	dynamic_cast<const types::Array*>(&pointee);	3✔
420
421	// Parenthesise only when it is needed to bind tighter than [] or ()
422	std::string decorated = pointee_is_function_or_array ? "(" + name + ")" : "" + name;	3✔
423
424	val << declaration(decorated, pointee);	3✔
425	} else {	3✔
426	val << "void*";	1✔
427	val << " " << name;	1✔
428	}
429	} else if (auto ref_type = dynamic_cast<const Reference*>(&type)) {	7✔
430	val << declaration("&" + name, ref_type->reference_type());	×
431	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	3✔
432	if (structure_type->storage_type() == types::StorageType_NV_Shared) {	3✔
433	val << "__shared__ ";	×
434	} else if (structure_type->storage_type() == types::StorageType_NV_Constant) {	3✔
435	val << "__constant__ ";	×
436	}	×
437	val << structure_type->name();	3✔
438	val << " ";	3✔
439	val << name;	3✔
440	} else if (auto function_type = dynamic_cast<const types::Function*>(&type)) {	3✔
441	std::stringstream params;	×
442	for (size_t i = 0; i < function_type->num_params(); ++i) {	×
443	params << declaration("", function_type->param_type(symbolic::integer(i)));	×
444	if (i + 1 < function_type->num_params()) params << ", ";	×
445	}	×
446	if (function_type->is_var_arg()) {	×
447	if (function_type->num_params() > 0) params << ", ";	×
448	params << "...";	×
449	}	×
450
451	const std::string fun_name = name + "(" + params.str() + ")";	×
452	val << declaration(fun_name, function_type->return_type());	×
453	} else {	×
454	throw std::runtime_error("Unknown declaration type");	×
455	}
456
457	if (use_alignment && type.alignment() > 0) {	18✔
458	val << " __attribute__((aligned(" << type.alignment() << ")))";	×
459	}	×
460
461	if (use_initializer && !type.initializer().empty()) {	18✔
462	val << " = " << type.initializer();	×
463	}	×
464
465	return val.str();	18✔
466	};	18✔
467
468	std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {	1✔
469	std::stringstream val;	1✔
470
471	val << "reinterpret_cast";	1✔
472	val << "<";	1✔
473	val << declaration("", type);	1✔
474	val << ">";	1✔
475	val << "(" << name << ")";	1✔
476
477	return val.str();	1✔
478	};	1✔
479
480	std::string CUDALanguageExtension::subset(const Function& function, const types::IType& type, const data_flow::Subset& sub) {	3✔
481	if (sub.empty()) {	3✔
482	return "";	1✔
483	}
484
485	if (dynamic_cast<const types::Scalar*>(&type)) {	2✔
486	return "";	×
487	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	2✔
488	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	1✔
489
490	if (sub.size() > 1) {	1✔
491	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
492	auto& element_type = array_type->element_type();	×
493	return subset_str + subset(function, element_type, element_subset);	×
494	} else {	×
495	return subset_str;	1✔
496	}
497	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	2✔
498	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	×
499
500	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
501	auto& pointee_type = pointer_type->pointee_type();	×
502	return subset_str + subset(function, pointee_type, element_subset);	×
503	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	1✔
504	auto& definition = function.structure(structure_type->name());	1✔
505
506	std::string subset_str = ".member_" + this->expression(sub.at(0));	1✔
507	if (sub.size() > 1) {	1✔
508	auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));	×
509	auto& member_type = definition.member_type(member);	×
510	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
511	return subset_str + subset(function, member_type, element_subset);	×
512	} else {	×
513	return subset_str;	1✔
514	}
515	}	1✔
516
517	throw std::invalid_argument("Invalid subset type");	×
518	};	3✔
519
520	std::string CUDALanguageExtension::expression(const symbolic::Expression& expr) {	5✔
521	CPPSymbolicPrinter printer;	5✔
522	return printer.apply(expr);	5✔
523	};	5✔
524
NEW 525	std::string CUDALanguageExtension::access_node(const data_flow::AccessNode& node) {	×
NEW 526	std::string name = node.data();	×
NEW 527	if (this->external_variables_.find(name) != this->external_variables_.end()) {	×
NEW 528	return "(&" + name + ")";	×
529	}
NEW 530	return name;	×
NEW 531	};	×
532
533	std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {	×
534	std::string op = code_to_string(tasklet.code());	×
535	std::vector<std::string> arguments;	×
536	for (size_t i = 0; i < tasklet.inputs().size(); ++i) {	×
NEW 537	arguments.push_back(tasklet.input(i));	×
538	}	×
539
540	if (tasklet.code() == data_flow::TaskletCode::assign) {	×
541	return arguments.at(0);	×
542	} else if (data_flow::is_infix(tasklet.code())) {	×
543	switch (data_flow::arity(tasklet.code())) {	×
544	case 1:
545	return op + arguments.at(0);	×
546	case 2:
547	return arguments.at(0) + " " + op + " " + arguments.at(1);	×
548	default:
549	throw std::runtime_error("Unsupported arity");	×
550	}
551	} else {
552	return op + "(" + helpers::join(arguments, ", ") + ")";	×
553	}
554	};	×
555
556	std::string CUDALanguageExtension::zero(const types::PrimitiveType prim_type) {	×
557	switch (prim_type) {	×
558	case types::Void:
559	throw InvalidSDFGException("No zero for void type possible");	×
560	case types::Bool:
561	return "false";	×
562	case types::Int8:
563	return "0";	×
564	case types::Int16:
565	return "0";	×
566	case types::Int32:
567	return "0";	×
568	case types::Int64:
569	return "0ll";	×
570	case types::Int128:
571	return "0";	×
572	case types::UInt8:
573	return "0u";	×
574	case types::UInt16:
575	return "0u";	×
576	case types::UInt32:
577	return "0u";	×
578	case types::UInt64:
579	return "0ull";	×
580	case types::UInt128:
581	return "0";	×
582	case types::Half:
583	return "CUDART_ZERO_FP16";	×
584	case types::BFloat:
585	return "CUDART_ZERO_BF16";	×
586	case types::Float:
587	return "0.0f";	×
588	case types::Double:
589	return "0.0";	×
590	case types::X86_FP80:
591	return "0.0l";	×
592	case types::FP128:
593	return "0.0";	×
594	case types::PPC_FP128:
595	return "0.0";	×
596	}	×
597	}	×
598
599	} // namespace codegen
600	} // namespace sdfg

daisytuner / sdfglib / 16779684622

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