19224974600

Committed 10 Nov 2025 08:11AM UTC coverage: 61.223% (+0.02%) from 61.206%

Build # 19224974600

Build Type

Pull #331

github

Committed by

web-flow

Commit Message

Merge 66fefbaff into d01b8b39b

Pull Request Pull Request #331: allow interpretation of pointers as ints in symbolic expressions

Run Details

54 of 95 new or added lines in 19 files covered. (56.84%)

3 existing lines in 3 files now uncovered.

10321 of 16858 relevant lines covered (61.22%)

107.0 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()) {
            // ISO C++ forbids empty parameter lists before ...
            if (function_type->num_params() > 0) {
                params << ", ";
                params << "...";
            }
        }

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

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

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

    return val.str();
};

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

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

    return val.str();
};

std::string CUDALanguageExtension::subset(const 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()) {	×
105	// ISO C++ forbids empty parameter lists before ...
106	if (function_type->num_params() > 0) {	×
107	params << ", ";	×
108	params << "...";	×
109	}	×
110	}	×
111
112	const std::string fun_name = name + "(" + params.str() + ")";	×
113	val << declaration(fun_name, function_type->return_type());	×
114	} else {	×
115	throw std::runtime_error("Unknown declaration type");	×
116	}
117
118	if (use_alignment && type.alignment() > 0) {	18✔
119	val << " __attribute__((aligned(" << type.alignment() << ")))";	×
120	}	×
121
122	if (use_initializer && !type.initializer().empty()) {	18✔
123	val << " = " << type.initializer();	×
124	}	×
125
126	return val.str();	18✔
127	};	18✔
128
129	std::string CUDALanguageExtension::type_cast(const std::string& name, const types::IType& type) {	1✔
130	std::stringstream val;	1✔
131
132	val << "reinterpret_cast";	1✔
133	val << "<";	1✔
134	val << declaration("", type);	1✔
135	val << ">";	1✔
136	val << "(" << name << ")";	1✔
137
138	return val.str();	1✔
139	};	1✔
140
141	std::string CUDALanguageExtension::subset(const types::IType& type, const data_flow::Subset& sub) {	3✔
142	if (sub.empty()) {	3✔
143	return "";	1✔
144	}
145
146	if (dynamic_cast<const types::Scalar*>(&type)) {	2✔
147	return "";	×
148	} else if (auto array_type = dynamic_cast<const types::Array*>(&type)) {	2✔
149	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	1✔
150
151	if (sub.size() > 1) {	1✔
152	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
153	auto& element_type = array_type->element_type();	×
NEW 154	return subset_str + subset(element_type, element_subset);	×
155	} else {	×
156	return subset_str;	1✔
157	}
158	} else if (auto pointer_type = dynamic_cast<const types::Pointer*>(&type)) {	2✔
159	std::string subset_str = "[" + this->expression(sub.at(0)) + "]";	×
160
161	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
162	auto& pointee_type = pointer_type->pointee_type();	×
NEW 163	return subset_str + subset(pointee_type, element_subset);	×
164	} else if (auto structure_type = dynamic_cast<const types::Structure*>(&type)) {	1✔
165	auto& definition = this->function_.structure(structure_type->name());	1✔
166
167	std::string subset_str = ".member_" + this->expression(sub.at(0));	1✔
168	if (sub.size() > 1) {	1✔
169	auto member = SymEngine::rcp_dynamic_cast<const SymEngine::Integer>(sub.at(0));	×
170	auto& member_type = definition.member_type(member);	×
171	data_flow::Subset element_subset(sub.begin() + 1, sub.end());	×
NEW 172	return subset_str + subset(member_type, element_subset);	×
173	} else {	×
174	return subset_str;	1✔
175	}
176	}	1✔
177
178	throw std::invalid_argument("Invalid subset type");	×
179	};	3✔
180
181	std::string CUDALanguageExtension::expression(const symbolic::Expression expr) {	5✔
182	CPPSymbolicPrinter printer(this->function_, this->external_prefix_);	5✔
183	return printer.apply(expr);	5✔
184	};	5✔
185
186	std::string CUDALanguageExtension::access_node(const data_flow::AccessNode& node) {	×
187	if (dynamic_cast<const data_flow::ConstantNode*>(&node)) {	×
188	std::string name = node.data();	×
189	if (symbolic::is_nullptr(symbolic::symbol(name))) {	×
NEW 190	return "nullptr";	×
191	}
192	return name;	×
193	} else {	×
194	std::string name = node.data();	×
NEW 195	if (this->function_.is_external(name)) {	×
196	return "(&" + name + ")";	×
197	}
198	return name;	×
199	}	×
200	};	×
201
202	std::string CUDALanguageExtension::tasklet(const data_flow::Tasklet& tasklet) {	×
203	switch (tasklet.code()) {	×
204	case data_flow::TaskletCode::assign:
205	return tasklet.inputs().at(0);	×
206	case data_flow::TaskletCode::fp_neg:
207	return "-" + tasklet.inputs().at(0);	×
208	case data_flow::TaskletCode::fp_add:
209	return tasklet.inputs().at(0) + " + " + tasklet.inputs().at(1);	×
210	case data_flow::TaskletCode::fp_sub:
211	return tasklet.inputs().at(0) + " - " + tasklet.inputs().at(1);	×
212	case data_flow::TaskletCode::fp_mul:
213	return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1);	×
214	case data_flow::TaskletCode::fp_div:
215	return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);	×
216	case data_flow::TaskletCode::fp_rem:
217	return "remainder(" + tasklet.inputs().at(0) + ", " + tasklet.inputs().at(1) + ")";	×
218	case data_flow::TaskletCode::fp_fma:
219	return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1) + " + " + tasklet.inputs().at(2);	×
220	case data_flow::TaskletCode::fp_oeq:
221	return tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);	×
222	case data_flow::TaskletCode::fp_one:
223	return tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);	×
224	case data_flow::TaskletCode::fp_ogt:
225	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
226	case data_flow::TaskletCode::fp_oge:
227	return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
228	case data_flow::TaskletCode::fp_olt:
229	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
230	case data_flow::TaskletCode::fp_ole:
231	return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
232	case data_flow::TaskletCode::fp_ord:
233	return "std::isnan(" + tasklet.inputs().at(0) + ") && std::isnan(" + tasklet.inputs().at(1) + ")";	×
234	case data_flow::TaskletCode::fp_ueq:
235	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
236	tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);	×
237	case data_flow::TaskletCode::fp_une:
238	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
239	tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);	×
240	case data_flow::TaskletCode::fp_ugt:
241	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
242	tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
243	case data_flow::TaskletCode::fp_uge:
244	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
245	tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
246	case data_flow::TaskletCode::fp_ult:
247	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
248	tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
249	case data_flow::TaskletCode::fp_ule:
250	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")" + " \|\| " +	×
251	tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
252	case data_flow::TaskletCode::fp_uno:
253	return "std::isnan(" + tasklet.inputs().at(0) + ") \|\| std::isnan(" + tasklet.inputs().at(1) + ")";	×
254	case data_flow::TaskletCode::int_add:
255	return tasklet.inputs().at(0) + " + " + tasklet.inputs().at(1);	×
256	case data_flow::TaskletCode::int_sub:
257	return tasklet.inputs().at(0) + " - " + tasklet.inputs().at(1);	×
258	case data_flow::TaskletCode::int_mul:
259	return tasklet.inputs().at(0) + " * " + tasklet.inputs().at(1);	×
260	case data_flow::TaskletCode::int_sdiv:
261	return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);	×
262	case data_flow::TaskletCode::int_srem:
263	return tasklet.inputs().at(0) + " % " + tasklet.inputs().at(1);	×
264	case data_flow::TaskletCode::int_udiv:
265	return tasklet.inputs().at(0) + " / " + tasklet.inputs().at(1);	×
266	case data_flow::TaskletCode::int_urem:
267	return tasklet.inputs().at(0) + " % " + tasklet.inputs().at(1);	×
268	case data_flow::TaskletCode::int_and:
269	return tasklet.inputs().at(0) + " & " + tasklet.inputs().at(1);	×
270	case data_flow::TaskletCode::int_or:
271	return tasklet.inputs().at(0) + " \| " + tasklet.inputs().at(1);	×
272	case data_flow::TaskletCode::int_xor:
273	return tasklet.inputs().at(0) + " ^ " + tasklet.inputs().at(1);	×
274	case data_flow::TaskletCode::int_shl:
275	return tasklet.inputs().at(0) + " << " + tasklet.inputs().at(1);	×
276	case data_flow::TaskletCode::int_lshr:
277	return tasklet.inputs().at(0) + " >> " + tasklet.inputs().at(1);	×
278	case data_flow::TaskletCode::int_ashr:
279	return tasklet.inputs().at(0) + " >> " + tasklet.inputs().at(1);	×
280	case data_flow::TaskletCode::int_smin:
281	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
282	tasklet.inputs().at(1);	×
283	case data_flow::TaskletCode::int_smax:
284	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
285	tasklet.inputs().at(1);	×
286	case data_flow::TaskletCode::int_scmp:
287	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? -1 : (" + tasklet.inputs().at(0) +	×
288	" > " + tasklet.inputs().at(1) + " ? 1 : 0)";	×
289	case data_flow::TaskletCode::int_umin:
290	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
291	tasklet.inputs().at(1);	×
292	case data_flow::TaskletCode::int_umax:
293	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1) + " ? " + tasklet.inputs().at(0) + " : " +	×
294	tasklet.inputs().at(1);	×
295	case data_flow::TaskletCode::int_ucmp:
296	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1) + " ? -1 : (" + tasklet.inputs().at(0) +	×
297	" > " + tasklet.inputs().at(1) + " ? 1 : 0)";	×
298	case data_flow::TaskletCode::int_abs:
299	return "(" + tasklet.inputs().at(0) + " < 0 ? -" + tasklet.inputs().at(0) + " : " + tasklet.inputs().at(0) +	×
300	")";
301	case data_flow::TaskletCode::int_eq:
302	return tasklet.inputs().at(0) + " == " + tasklet.inputs().at(1);	×
303	case data_flow::TaskletCode::int_ne:
304	return tasklet.inputs().at(0) + " != " + tasklet.inputs().at(1);	×
305	case data_flow::TaskletCode::int_sgt:
306	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
307	case data_flow::TaskletCode::int_sge:
308	return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
309	case data_flow::TaskletCode::int_slt:
310	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
311	case data_flow::TaskletCode::int_sle:
312	return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
313	case data_flow::TaskletCode::int_ugt:
314	return tasklet.inputs().at(0) + " > " + tasklet.inputs().at(1);	×
315	case data_flow::TaskletCode::int_uge:
316	return tasklet.inputs().at(0) + " >= " + tasklet.inputs().at(1);	×
317	case data_flow::TaskletCode::int_ult:
318	return tasklet.inputs().at(0) + " < " + tasklet.inputs().at(1);	×
319	case data_flow::TaskletCode::int_ule:
320	return tasklet.inputs().at(0) + " <= " + tasklet.inputs().at(1);	×
321	};
322	throw std::invalid_argument("Invalid tasklet code");	×
323	};	×
324
325	std::string CUDALanguageExtension::zero(const types::PrimitiveType prim_type) {	×
326	switch (prim_type) {	×
327	case types::Void:
328	throw InvalidSDFGException("No zero for void type possible");	×
329	case types::Bool:
330	return "false";	×
331	case types::Int8:
332	return "0";	×
333	case types::Int16:
334	return "0";	×
335	case types::Int32:
336	return "0";	×
337	case types::Int64:
338	return "0ll";	×
339	case types::Int128:
340	return "0";	×
341	case types::UInt8:
342	return "0u";	×
343	case types::UInt16:
344	return "0u";	×
345	case types::UInt32:
346	return "0u";	×
347	case types::UInt64:
348	return "0ull";	×
349	case types::UInt128:
350	return "0";	×
351	case types::Half:
352	return "CUDART_ZERO_FP16";	×
353	case types::BFloat:
354	return "CUDART_ZERO_BF16";	×
355	case types::Float:
356	return "0.0f";	×
357	case types::Double:
358	return "0.0";	×
359	case types::X86_FP80:
360	return "0.0l";	×
361	case types::FP128:
362	return "0.0";	×
363	case types::PPC_FP128:
364	return "0.0";	×
365	}	×
366	}	×
367
368	} // namespace codegen
369	} // namespace sdfg

daisytuner / sdfglib / 19224974600

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