29992851043

Committed 23 Jul 2026 08:52AM UTC coverage: 64.1% (+0.3%) from 63.787%

Build # 29992851043

Build Type

Pull #866

github

Committed by

web-flow

Commit Message

Merge f4354fb28 into 2e568810b

Pull Request Pull Request #866: Add Support for Complex Types

Coverage Stats

188 of 296 new or added lines in 12 files covered. (63.51%)

272 existing lines in 10 files now uncovered.

42854 of 66855 relevant lines covered (64.1%)

737.97 hits per line

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

37.18

/sdfg/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";
        case types::PrimitiveType::CHalf:
        case types::PrimitiveType::CBFloat:
        case types::PrimitiveType::CFloat:
        case types::PrimitiveType::CDouble:
        case types::PrimitiveType::CFP128:
            return complex_type_name(prim_type);
    }

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

daisytuner / docc / 29992851043

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