19276018096

Committed 11 Nov 2025 07:13PM UTC coverage: 61.767%. Remained the same

Build # 19276018096

Build Type

push

github

Committed by

web-flow

Commit Message

Merge pull request #343 from daisytuner/var-arg-funcs

fixes codegen for var arg functions without fixed params

Run Details

0 of 2 new or added lines in 2 files covered. (0.0%)

4 existing lines in 2 files now uncovered.

10538 of 17061 relevant lines covered (61.77%)

107.11 hits per line

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

33.48

/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 {

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().is_nv_shared()) {
            val << "__shared__ ";
        } else if (scalar_type->storage_type().is_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().is_nv_shared()) {
            val << "__shared__ ";
        } else if (structure_type->storage_type().is_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 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(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(pointee_type, element_subset);
    } else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {
        auto& definition = this->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(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(this->function_, this->external_prefix_);
    return printer.apply(expr);
};

std::string CUDALanguageExtension::access_node(const data_flow::AccessNode& node) {
    if (dynamic_cast<const data_flow::ConstantNode*>(&node)) {
        std::string name = node.data();
        if (symbolic::is_nullptr(symbolic::symbol(name))) {
            return "nullptr";
        }
        return name;
    } else {
        std::string name = node.data();
        if (this->function_.is_external(name)) {
            return "(&" + name + ")";
        }
        return name;
    }
};

std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {
    switch (tasklet.code()) {
        case data_flow::TaskletCode::assign:
            return tasklet.inputs().at(0);
        case data_flow::TaskletCode::fp_neg:
            return "-" + tasklet.inputs().at(0);
        case data_flow::TaskletCode::fp_add:
            return tasklet.inputs().at(0) + " + " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_sub:
            return tasklet.inputs().at(0) + " - " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_mul:
            return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_div:
            return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_rem:
            return "remainder(" + tasklet.inputs().at(0) + ", " + tasklet.inputs().at(1) + ")";
        case data_flow::TaskletCode::fp_fma:
            return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1) + " + " + tasklet.inputs().at(2);
        case data_flow::TaskletCode::fp_oeq:
            return tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_one:
            return tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_ogt:
            return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_oge:
            return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_olt:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_ole:
            return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_ord:
            return "std::isnan(" + tasklet.inputs().at(0) + ") && std::isnan(" + tasklet.inputs().at(1) + ")";
        case data_flow::TaskletCode::fp_ueq:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")" + " || " +
                   tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_une:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")" + " || " +
                   tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_ugt:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")" + " || " +
                   tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_uge:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")" + " || " +
                   tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_ult:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")" + " || " +
                   tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_ule:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")" + " || " +
                   tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::fp_uno:
            return "std::isnan(" + tasklet.inputs().at(0) + ") || std::isnan(" + tasklet.inputs().at(1) + ")";
        case data_flow::TaskletCode::int_add:
            return tasklet.inputs().at(0) + " + " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_sub:
            return tasklet.inputs().at(0) + " - " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_mul:
            return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_sdiv:
            return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_srem:
            return tasklet.inputs().at(0) + " % " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_udiv:
            return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_urem:
            return tasklet.inputs().at(0) + " % " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_and:
            return tasklet.inputs().at(0) + " & " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_or:
            return tasklet.inputs().at(0) + " | " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_xor:
            return tasklet.inputs().at(0) + " ^ " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_shl:
            return tasklet.inputs().at(0) + " << " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_lshr:
            return tasklet.inputs().at(0) + " >> " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_ashr:
            return tasklet.inputs().at(0) + " >> " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_smin:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +
                   tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_smax:
            return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +
                   tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_scmp:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? -1 : (" + tasklet.inputs().at(0) +
                   " > " + tasklet.inputs().at(1) + " ? 1 : 0)";
        case data_flow::TaskletCode::int_umin:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +
                   tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_umax:
            return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +
                   tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_ucmp:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? -1 : (" + tasklet.inputs().at(0) +
                   " > " + tasklet.inputs().at(1) + " ? 1 : 0)";
        case data_flow::TaskletCode::int_abs:
            return "(" + tasklet.inputs().at(0) + " < 0 ? -" + tasklet.inputs().at(0) + " : " + tasklet.inputs().at(0) +
                   ")";
        case data_flow::TaskletCode::int_eq:
            return tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_ne:
            return tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_sgt:
            return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_sge:
            return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_slt:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_sle:
            return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_ugt:
            return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_uge:
            return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_ult:
            return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);
        case data_flow::TaskletCode::int_ule:
            return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);
    };
    throw std::invalid_argument("Invalid tasklet code");
};

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	std::string CUDALanguageExtension::primitive_type(const types::PrimitiveType prim_type) {	20✔
12	switch (prim_type) {	20✔
13	case types::PrimitiveType::Void:
14	return "void";	1✔
15	case types::PrimitiveType::Bool:
16	return "bool";	1✔
17	case types::PrimitiveType::Int8:
18	return "signed char";	1✔
19	case types::PrimitiveType::Int16:
20	return "short";	1✔
21	case types::PrimitiveType::Int32:
22	return "int";	6✔
23	case types::PrimitiveType::Int64:
24	return "long long";	1✔
25	case types::PrimitiveType::Int128:
26	return "__int128";	×
27	case types::PrimitiveType::UInt8:
28	return "char";	3✔
29	case types::PrimitiveType::UInt16:
30	return "unsigned short";	1✔
31	case types::PrimitiveType::UInt32:
32	return "unsigned int";	1✔
33	case types::PrimitiveType::UInt64:
34	return "unsigned long long";	1✔
35	case types::PrimitiveType::UInt128:
36	return "unsigned __int128";	×
37	case types::PrimitiveType::Half:
38	return "__fp16";	×
39	case types::PrimitiveType::BFloat:
40	return "__bf16";	×
41	case types::PrimitiveType::Float:
42	return "float";	2✔
43	case types::PrimitiveType::Double:
44	return "double";	1✔
45	case types::PrimitiveType::X86_FP80:
46	return "long double";	×
47	case types::PrimitiveType::FP128:
48	return "__float128";	×
49	case types::PrimitiveType::PPC_FP128:
50	return "__float128";	×
51	}
52
53	throw std::runtime_error("Unknown primitive type");	×
54	};	20✔
55
56	std::string CUDALanguageExtension::
57	declaration(const std::string& name, const types::IType& type, bool use_initializer, bool use_alignment) {	18✔
58	std::stringstream val;	18✔
59
60	if (auto scalar_type = dynamic_cast<const types::Scalar*>(&type)) {	18✔
61	if (scalar_type->storage_type().is_nv_shared()) {	8✔
62	val << "__shared__ ";	×
63	} else if (scalar_type->storage_type().is_nv_constant()) {	8✔
64	val << "__constant__ ";	×
65	}	×
66	val << primitive_type(scalar_type->primitive_type());	8✔
67	val << " ";	8✔
68	val << name;	8✔
69	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	18✔
70	auto& element_type = array_type->element_type();	3✔
71	val << declaration(name + "[" + this->expression(array_type->num_elements()) + "]", element_type);	3✔
72	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	10✔
73	if (pointer_type->has_pointee_type()) {	4✔
74	const types::IType& pointee = pointer_type->pointee_type();	3✔
75
76	const bool pointee_is_function_or_array = dynamic_cast<const types::Function*>(&pointee) \|\|	6✔
77	dynamic_cast<const types::Array*>(&pointee);	3✔
78
79	// Parenthesise only when it is needed to bind tighter than [] or ()
80	std::string decorated = pointee_is_function_or_array ? "(" + name + ")" : "" + name;	3✔
81
82	val << declaration(decorated, pointee);	3✔
83	} else {	3✔
84	val << "void*";	1✔
85	val << " " << name;	1✔
86	}
87	} else if (auto ref_type = dynamic_cast<const Reference*>(&type)) {	7✔
88	val << declaration("&" + name, ref_type->reference_type());	×
89	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	3✔
90	if (structure_type->storage_type().is_nv_shared()) {	3✔
91	val << "__shared__ ";	×
92	} else if (structure_type->storage_type().is_nv_constant()) {	3✔
93	val << "__constant__ ";	×
94	}	×
95	val << structure_type->name();	3✔
96	val << " ";	3✔
97	val << name;	3✔
98	} else if (auto function_type = dynamic_cast<const types::Function*>(&type)) {	3✔
99	std::stringstream params;	×
100	for (size_t i = 0; i < function_type->num_params(); ++i) {	×
101	params << declaration("", function_type->param_type(symbolic::integer(i)));	×
102	if (i + 1 < function_type->num_params()) params << ", ";	×
103	}	×
104	if (function_type->is_var_arg()) {	×
UNCOV 105	if (function_type->num_params() > 0) {	×
106	params << ", ";	×
107	}	×
NEW 108	params << "...";	×
UNCOV 109	}	×
110
111	const std::string fun_name = name + "(" + params.str() + ")";	×
112	val << declaration(fun_name, function_type->return_type());	×
113	} else {	×
114	throw std::runtime_error("Unknown declaration type");	×
115	}
116
117	if (use_alignment && type.alignment() > 0) {	18✔
118	val << " __attribute__((aligned(" << type.alignment() << ")))";	×
119	}	×
120
121	if (use_initializer && !type.initializer().empty()) {	18✔
122	val << " = " << type.initializer();	×
123	}	×
124
125	return val.str();	18✔
126	};	18✔
127
128	std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {	1✔
129	std::stringstream val;	1✔
130
131	val << "reinterpret_cast";	1✔
132	val << "<";	1✔
133	val << declaration("", type);	1✔
134	val << ">";	1✔
135	val << "(" << name << ")";	1✔
136
137	return val.str();	1✔
138	};	1✔
139
140	std::string CUDALanguageExtension::subset(const types::IType& type, const data_flow::Subset& sub) {	3✔
141	if (sub.empty()) {	3✔
142	return "";	1✔
143	}
144
145	if (dynamic_cast<const types::Scalar*>(&type)) {	2✔
146	return "";	×
147	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	2✔
148	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	1✔
149
150	if (sub.size() > 1) {	1✔
151	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
152	auto& element_type = array_type->element_type();	×
153	return subset_str + subset(element_type, element_subset);	×
154	} else {	×
155	return subset_str;	1✔
156	}
157	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	2✔
158	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	×
159
160	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
161	auto& pointee_type = pointer_type->pointee_type();	×
162	return subset_str + subset(pointee_type, element_subset);	×
163	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	1✔
164	auto& definition = this->function_.structure(structure_type->name());	1✔
165
166	std::string subset_str = ".member_" + this->expression(sub.at(0));	1✔
167	if (sub.size() > 1) {	1✔
168	auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));	×
169	auto& member_type = definition.member_type(member);	×
170	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
171	return subset_str + subset(member_type, element_subset);	×
172	} else {	×
173	return subset_str;	1✔
174	}
175	}	1✔
176
177	throw std::invalid_argument("Invalid subset type");	×
178	};	3✔
179
180	std::string CUDALanguageExtension::expression(const symbolic::Expression expr) {	5✔
181	CPPSymbolicPrinter printer(this->function_, this->external_prefix_);	5✔
182	return printer.apply(expr);	5✔
183	};	5✔
184
185	std::string CUDALanguageExtension::access_node(const data_flow::AccessNode& node) {	×
186	if (dynamic_cast<const data_flow::ConstantNode*>(&node)) {	×
187	std::string name = node.data();	×
188	if (symbolic::is_nullptr(symbolic::symbol(name))) {	×
189	return "nullptr";	×
190	}
191	return name;	×
192	} else {	×
193	std::string name = node.data();	×
194	if (this->function_.is_external(name)) {	×
195	return "(&" + name + ")";	×
196	}
197	return name;	×
198	}	×
199	};	×
200
201	std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {	×
202	switch (tasklet.code()) {	×
203	case data_flow::TaskletCode::assign:
204	return tasklet.inputs().at(0);	×
205	case data_flow::TaskletCode::fp_neg:
206	return "-" + tasklet.inputs().at(0);	×
207	case data_flow::TaskletCode::fp_add:
208	return tasklet.inputs().at(0) + " + " + tasklet.inputs().at(1);	×
209	case data_flow::TaskletCode::fp_sub:
210	return tasklet.inputs().at(0) + " - " + tasklet.inputs().at(1);	×
211	case data_flow::TaskletCode::fp_mul:
212	return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1);	×
213	case data_flow::TaskletCode::fp_div:
214	return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);	×
215	case data_flow::TaskletCode::fp_rem:
216	return "remainder(" + tasklet.inputs().at(0) + ", " + tasklet.inputs().at(1) + ")";	×
217	case data_flow::TaskletCode::fp_fma:
218	return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1) + " + " + tasklet.inputs().at(2);	×
219	case data_flow::TaskletCode::fp_oeq:
220	return tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);	×
221	case data_flow::TaskletCode::fp_one:
222	return tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);	×
223	case data_flow::TaskletCode::fp_ogt:
224	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
225	case data_flow::TaskletCode::fp_oge:
226	return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
227	case data_flow::TaskletCode::fp_olt:
228	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
229	case data_flow::TaskletCode::fp_ole:
230	return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
231	case data_flow::TaskletCode::fp_ord:
232	return "std::isnan(" + tasklet.inputs().at(0) + ") && std::isnan(" + tasklet.inputs().at(1) + ")";	×
233	case data_flow::TaskletCode::fp_ueq:
234	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
235	tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);	×
236	case data_flow::TaskletCode::fp_une:
237	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
238	tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);	×
239	case data_flow::TaskletCode::fp_ugt:
240	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
241	tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
242	case data_flow::TaskletCode::fp_uge:
243	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
244	tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
245	case data_flow::TaskletCode::fp_ult:
246	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
247	tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
248	case data_flow::TaskletCode::fp_ule:
249	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
250	tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
251	case data_flow::TaskletCode::fp_uno:
252	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")";	×
253	case data_flow::TaskletCode::int_add:
254	return tasklet.inputs().at(0) + " + " + tasklet.inputs().at(1);	×
255	case data_flow::TaskletCode::int_sub:
256	return tasklet.inputs().at(0) + " - " + tasklet.inputs().at(1);	×
257	case data_flow::TaskletCode::int_mul:
258	return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1);	×
259	case data_flow::TaskletCode::int_sdiv:
260	return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);	×
261	case data_flow::TaskletCode::int_srem:
262	return tasklet.inputs().at(0) + " % " + tasklet.inputs().at(1);	×
263	case data_flow::TaskletCode::int_udiv:
264	return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);	×
265	case data_flow::TaskletCode::int_urem:
266	return tasklet.inputs().at(0) + " % " + tasklet.inputs().at(1);	×
267	case data_flow::TaskletCode::int_and:
268	return tasklet.inputs().at(0) + " & " + tasklet.inputs().at(1);	×
269	case data_flow::TaskletCode::int_or:
270	return tasklet.inputs().at(0) + " \| " + tasklet.inputs().at(1);	×
271	case data_flow::TaskletCode::int_xor:
272	return tasklet.inputs().at(0) + " ^ " + tasklet.inputs().at(1);	×
273	case data_flow::TaskletCode::int_shl:
274	return tasklet.inputs().at(0) + " << " + tasklet.inputs().at(1);	×
275	case data_flow::TaskletCode::int_lshr:
276	return tasklet.inputs().at(0) + " >> " + tasklet.inputs().at(1);	×
277	case data_flow::TaskletCode::int_ashr:
278	return tasklet.inputs().at(0) + " >> " + tasklet.inputs().at(1);	×
279	case data_flow::TaskletCode::int_smin:
280	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
281	tasklet.inputs().at(1);	×
282	case data_flow::TaskletCode::int_smax:
283	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
284	tasklet.inputs().at(1);	×
285	case data_flow::TaskletCode::int_scmp:
286	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? -1 : (" + tasklet.inputs().at(0) +	×
287	" > " + tasklet.inputs().at(1) + " ? 1 : 0)";	×
288	case data_flow::TaskletCode::int_umin:
289	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
290	tasklet.inputs().at(1);	×
291	case data_flow::TaskletCode::int_umax:
292	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
293	tasklet.inputs().at(1);	×
294	case data_flow::TaskletCode::int_ucmp:
295	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? -1 : (" + tasklet.inputs().at(0) +	×
296	" > " + tasklet.inputs().at(1) + " ? 1 : 0)";	×
297	case data_flow::TaskletCode::int_abs:
298	return "(" + tasklet.inputs().at(0) + " < 0 ? -" + tasklet.inputs().at(0) + " : " + tasklet.inputs().at(0) +	×
299	")";
300	case data_flow::TaskletCode::int_eq:
301	return tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);	×
302	case data_flow::TaskletCode::int_ne:
303	return tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);	×
304	case data_flow::TaskletCode::int_sgt:
305	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
306	case data_flow::TaskletCode::int_sge:
307	return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
308	case data_flow::TaskletCode::int_slt:
309	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
310	case data_flow::TaskletCode::int_sle:
311	return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
312	case data_flow::TaskletCode::int_ugt:
313	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
314	case data_flow::TaskletCode::int_uge:
315	return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
316	case data_flow::TaskletCode::int_ult:
317	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
318	case data_flow::TaskletCode::int_ule:
319	return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
320	};
321	throw std::invalid_argument("Invalid tasklet code");	×
322	};	×
323
324	std::string CUDALanguageExtension::zero(const types::PrimitiveType prim_type) {	×
325	switch (prim_type) {	×
326	case types::Void:
327	throw InvalidSDFGException("No zero for void type possible");	×
328	case types::Bool:
329	return "false";	×
330	case types::Int8:
331	return "0";	×
332	case types::Int16:
333	return "0";	×
334	case types::Int32:
335	return "0";	×
336	case types::Int64:
337	return "0ll";	×
338	case types::Int128:
339	return "0";	×
340	case types::UInt8:
341	return "0u";	×
342	case types::UInt16:
343	return "0u";	×
344	case types::UInt32:
345	return "0u";	×
346	case types::UInt64:
347	return "0ull";	×
348	case types::UInt128:
349	return "0";	×
350	case types::Half:
351	return "CUDART_ZERO_FP16";	×
352	case types::BFloat:
353	return "CUDART_ZERO_BF16";	×
354	case types::Float:
355	return "0.0f";	×
356	case types::Double:
357	return "0.0";	×
358	case types::X86_FP80:
359	return "0.0l";	×
360	case types::FP128:
361	return "0.0";	×
362	case types::PPC_FP128:
363	return "0.0";	×
364	}	×
365	}	×
366
367	} // namespace codegen
368	} // namespace sdfg

daisytuner / sdfglib / 19276018096

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