20781520363

Committed 07 Jan 2026 10:50AM UTC coverage: 61.994% (-0.2%) from 62.168%

Build # 20781520363

Build Type

Pull #438

github

Committed by

web-flow

Commit Message

Merge af6583671 into 0c34ccd02

Pull Request Pull Request #438: Cuda tiling

Run Details

8 of 51 new or added lines in 6 files covered. (15.69%)

2 existing lines in 1 file now uncovered.

14904 of 24041 relevant lines covered (61.99%)

88.13 hits per line

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

29.13

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

daisytuner / sdfglib / 20781520363

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