29934626444

Committed 22 Jul 2026 03:42PM UTC coverage: 63.883% (+0.1%) from 63.782%

Build # 29934626444

Build Type

Pull #866

github

Committed by

web-flow

Commit Message

Merge 425767288 into d833d292d

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

Coverage Stats

244 of 298 new or added lines in 11 files covered. (81.88%)

1 existing line in 1 file now uncovered.

42410 of 66387 relevant lines covered (63.88%)

742.79 hits per line

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

83.02

/sdfg/src/codegen/utils.cpp

#include "sdfg/codegen/utils.h"

#include "sdfg/data_flow/data_flow_graph.h"
#include "sdfg/data_flow/tasklet.h"
#include "sdfg/exceptions.h"
#include "sdfg/function.h"

namespace sdfg {
namespace codegen {

// Constructor
PrettyPrinter::PrettyPrinter(int indent, bool frozen)
    : owned_stream(std::make_unique<std::stringstream>()), stream(*owned_stream.get()), indentSize(indent),
      frozen_(frozen) {}

PrettyPrinter::PrettyPrinter(std::ostream& stream, int indent, bool frozen)
    : stream(stream), indentSize(indent), frozen_(frozen) {}

// Set the indentation level
void PrettyPrinter::setIndent(int indent) { indentSize = indent; };

int PrettyPrinter::indent() const { return indentSize; };

int PrettyPrinter::changeIndent(int delta) { return indentSize += delta; };

// Get the underlying string
std::string PrettyPrinter::str() const { return owned_stream->str(); };

// Clear the stringstream content
void PrettyPrinter::clear() {
    owned_stream->str("");
    owned_stream->clear();
}

// Overload for manipulators (like std::endl)
PrettyPrinter& PrettyPrinter::operator<<(std::ostream& (*manip)(std::ostream&) ) {
    if (frozen_) {
        throw std::runtime_error("PrettyPrinter is frozen");
    }
    stream << manip;
    // Reset indent application on new lines
    if (manip == static_cast<std::ostream& (*) (std::ostream&)>(std::endl)) {
        isNewLine = true;
    }
    return *this;
};

// Apply indentation only at the beginning of a new line
void PrettyPrinter::applyIndent() {
    if (isNewLine && indentSize > 0) {
        stream << std::setw(indentSize) << "";
        isNewLine = false;
    }
};

Reference::Reference(const types::IType& reference_) : reference_(reference_.clone()) {};

Reference::Reference(
    types::StorageType storage_type, size_t alignment, const std::string& initializer, const types::IType& reference_
)
    : IType(storage_type, alignment, initializer), reference_(reference_.clone()) {};

std::unique_ptr<types::IType> Reference::clone() const {
    return std::make_unique<Reference>(this->storage_type(), this->alignment(), this->initializer(), *this->reference_);
};

types::TypeID Reference::type_id() const { return types::TypeID::Reference; };

types::PrimitiveType Reference::primitive_type() const { return this->reference_->primitive_type(); };

bool Reference::is_symbol() const { return false; };

const types::IType& Reference::reference_type() const { return *this->reference_; };

bool Reference::operator==(const types::IType& other) const {
    if (auto reference = dynamic_cast<const Reference*>(&other)) {
        return *(this->reference_) == *reference->reference_ && this->alignment_ == reference->alignment_;
    } else {
        return false;
    }
};

std::string Reference::print() const { return "Reference(" + this->reference_->print() + ")"; };


std::string complex_type_name(types::PrimitiveType prim_type) {
    switch (prim_type) {
        case types::PrimitiveType::CHalf:
            return "half2";
        case types::PrimitiveType::CBFloat:
            return "bfloat162";
        case types::PrimitiveType::CFloat:
            return "float2";
        case types::PrimitiveType::CDouble:
            return "double2";
        case types::PrimitiveType::CFP128:
            return "fp128_2";
        default:
            throw InvalidSDFGException("complex_type_name: not a complex primitive type");
    }
};

std::string complex_op_suffix(types::PrimitiveType prim_type) {
    switch (prim_type) {
        case types::PrimitiveType::CHalf:
            return "h";
        case types::PrimitiveType::CBFloat:
            return "bf";
        case types::PrimitiveType::CFloat:
            return "f";
        case types::PrimitiveType::CDouble:
            return "d";
        case types::PrimitiveType::CFP128:
            return "q";
        default:
            throw InvalidSDFGException("complex_op_suffix: not a complex primitive type");
    }
};

std::string complex_tasklet(sdfg::Function& function, const data_flow::Tasklet& tasklet) {
    // Determine the element suffix from the (complex) operand type.
    auto& graph = tasklet.get_parent();
    auto in_edges = graph.in_edges(tasklet);
    if (in_edges.begin() == in_edges.end()) {
        throw InvalidSDFGException("complex_tasklet: tasklet has no inputs");
    }
    auto operand_type = (*in_edges.begin()).result_type(function);
    std::string suffix = complex_op_suffix(operand_type->primitive_type());

    auto& inputs = tasklet.inputs();
    switch (tasklet.code()) {
        case data_flow::TaskletCode::complex_neg:
            return "__daisy_cneg_" + suffix + "(" + inputs.at(0) + ")";
        case data_flow::TaskletCode::complex_real:
            return "__daisy_creal_" + suffix + "(" + inputs.at(0) + ")";
        case data_flow::TaskletCode::complex_imag:
            return "__daisy_cimag_" + suffix + "(" + inputs.at(0) + ")";
        case data_flow::TaskletCode::complex_add:
            return "__daisy_cadd_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";
        case data_flow::TaskletCode::complex_sub:
            return "__daisy_csub_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";
        case data_flow::TaskletCode::complex_mul:
            return "__daisy_cmul_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";
        case data_flow::TaskletCode::complex_div:
            return "__daisy_cdiv_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";
        case data_flow::TaskletCode::complex_eq:
            return "__daisy_ceq_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";
        case data_flow::TaskletCode::complex_ne:
            return "__daisy_cne_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";
        default:
            throw InvalidSDFGException("complex_tasklet: not a complex tasklet code");
    }
};

std::string complex_support_preamble(bool device) {
    const std::string qual = device ? "__host__ __device__ static inline" : "static inline";

    std::stringstream out;
    out << "/* Complex type support (generated by sdfglib) */" << std::endl;

    // Component vector types. On GPU, float2/double2/half2/__nv_bfloat162 are provided by the
    // toolchain and reused; on CPU they are defined here. fp128_2 has no native type anywhere.
    if (device) {
        out << "typedef __nv_bfloat162 bfloat162;" << std::endl;
    } else {
        out << "typedef struct { float x; float y; } float2;" << std::endl;
        out << "typedef struct { double x; double y; } double2;" << std::endl;
        out << "typedef struct { _Float16 x; _Float16 y; } half2;" << std::endl;
        out << "typedef struct { __bf16 x; __bf16 y; } bfloat162;" << std::endl;
    }
    out << "typedef struct { __float128 x; __float128 y; } fp128_2;" << std::endl;

    // Element-wise helper functions. Arithmetic is carried out in a wide type (WT) so that
    // narrow element types (half/bfloat) that lack native arithmetic are computed via float.
    out << "#define __DAISY_DEFINE_COMPLEX(CT, ST, WT, SUF, QUAL) \\" << std::endl;
    out << "  QUAL CT __daisy_cadd_##SUF(CT a, CT b) { CT r; r.x = (ST)((WT)a.x + (WT)b.x); r.y = "
           "(ST)((WT)a.y + (WT)b.y); return r; } \\"
        << std::endl;
    out << "  QUAL CT __daisy_csub_##SUF(CT a, CT b) { CT r; r.x = (ST)((WT)a.x - (WT)b.x); r.y = "
           "(ST)((WT)a.y - (WT)b.y); return r; } \\"
        << std::endl;
    out << "  QUAL CT __daisy_cmul_##SUF(CT a, CT b) { CT r; r.x = (ST)((WT)a.x * (WT)b.x - (WT)a.y * (WT)b.y); "
           "r.y = (ST)((WT)a.x * (WT)b.y + (WT)a.y * (WT)b.x); return r; } \\"
        << std::endl;
    out << "  QUAL CT __daisy_cdiv_##SUF(CT a, CT b) { CT r; WT d = (WT)b.x * (WT)b.x + (WT)b.y * (WT)b.y; "
           "r.x = (ST)(((WT)a.x * (WT)b.x + (WT)a.y * (WT)b.y) / d); "
           "r.y = (ST)(((WT)a.y * (WT)b.x - (WT)a.x * (WT)b.y) / d); return r; } \\"
        << std::endl;
    out << "  QUAL CT __daisy_cneg_##SUF(CT a) { CT r; r.x = (ST)(-(WT)a.x); r.y = (ST)(-(WT)a.y); return r; } \\"
        << std::endl;
    out << "  QUAL ST __daisy_creal_##SUF(CT a) { return a.x; } \\" << std::endl;
    out << "  QUAL ST __daisy_cimag_##SUF(CT a) { return a.y; } \\" << std::endl;
    out << "  QUAL bool __daisy_ceq_##SUF(CT a, CT b) { return (WT)a.x == (WT)b.x && (WT)a.y == (WT)b.y; } \\"
        << std::endl;
    out << "  QUAL bool __daisy_cne_##SUF(CT a, CT b) { return !((WT)a.x == (WT)b.x && (WT)a.y == (WT)b.y); }"
        << std::endl;

    if (device) {
        out << "__DAISY_DEFINE_COMPLEX(float2, float, float, f, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(double2, double, double, d, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(half2, __half, float, h, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(bfloat162, __nv_bfloat16, float, bf, " << qual << ")" << std::endl;
        // __float128 is a host-only type; keep fp128 helpers off the device.
        out << "__DAISY_DEFINE_COMPLEX(fp128_2, __float128, __float128, q, __host__ inline)" << std::endl;
    } else {
        out << "__DAISY_DEFINE_COMPLEX(float2, float, float, f, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(double2, double, double, d, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(half2, _Float16, float, h, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(bfloat162, __bf16, float, bf, " << qual << ")" << std::endl;
        out << "__DAISY_DEFINE_COMPLEX(fp128_2, __float128, __float128, q, " << qual << ")" << std::endl;
    }
    out << "#undef __DAISY_DEFINE_COMPLEX" << std::endl;

    return out.str();
};


} // namespace codegen
} // namespace sdfg

1	#include "sdfg/codegen/utils.h"
2
3	#include "sdfg/data_flow/data_flow_graph.h"
4	#include "sdfg/data_flow/tasklet.h"
5	#include "sdfg/exceptions.h"
6	#include "sdfg/function.h"
7
8	namespace sdfg {
9	namespace codegen {
10
11	// Constructor
12	PrettyPrinter::PrettyPrinter(int indent, bool frozen)
13	: owned_stream(std::make_unique<std::stringstream>()), stream(*owned_stream.get()), indentSize(indent),	235✔
14	frozen_(frozen) {}	235✔
15
16	PrettyPrinter::PrettyPrinter(std::ostream& stream, int indent, bool frozen)
17	: stream(stream), indentSize(indent), frozen_(frozen) {}	×
18
19	// Set the indentation level
20	void PrettyPrinter::setIndent(int indent) { indentSize = indent; };	528✔
21
22	int PrettyPrinter::indent() const { return indentSize; };	360✔
23
24	int PrettyPrinter::changeIndent(int delta) { return indentSize += delta; };	×
25
26	// Get the underlying string
27	std::string PrettyPrinter::str() const { return owned_stream->str(); };	121✔
28
29	// Clear the stringstream content
30	void PrettyPrinter::clear() {	14✔
31	owned_stream->str("");	14✔
32	owned_stream->clear();	14✔
33	}	14✔
34
35	// Overload for manipulators (like std::endl)
36	PrettyPrinter& PrettyPrinter::operator<<(std::ostream& (*manip)(std::ostream&) ) {	1,640✔
37	if (frozen_) {	1,640✔
38	throw std::runtime_error("PrettyPrinter is frozen");	×
39	}	×
40	stream << manip;	1,640✔
41	// Reset indent application on new lines
42	if (manip == static_cast<std::ostream& (*) (std::ostream&)>(std::endl)) {	1,640✔
43	isNewLine = true;	1,640✔
44	}	1,640✔
45	return *this;	1,640✔
46	};	1,640✔
47
48	// Apply indentation only at the beginning of a new line
49	void PrettyPrinter::applyIndent() {	4,862✔
50	if (isNewLine && indentSize > 0) {	4,862✔
51	stream << std::setw(indentSize) << "";	1,121✔
52	isNewLine = false;	1,121✔
53	}	1,121✔
54	};	4,862✔
55
56	Reference::Reference(const types::IType& reference_) : reference_(reference_.clone()) {};	16✔
57
58	Reference::Reference(
59	types::StorageType storage_type, size_t alignment, const std::string& initializer, const types::IType& reference_
60	)
61	: IType(storage_type, alignment, initializer), reference_(reference_.clone()) {};	6✔
62
63	std::unique_ptr<types::IType> Reference::clone() const {	5✔
64	return std::make_unique<Reference>(this->storage_type(), this->alignment(), this->initializer(), *this->reference_);	5✔
65	};	5✔
66
67	types::TypeID Reference::type_id() const { return types::TypeID::Reference; };	30✔
68
69	types::PrimitiveType Reference::primitive_type() const { return this->reference_->primitive_type(); };	×
70
71	bool Reference::is_symbol() const { return false; };	1✔
72
73	const types::IType& Reference::reference_type() const { return *this->reference_; };	22✔
74
75	bool Reference::operator==(const types::IType& other) const {	4✔
76	if (auto reference = dynamic_cast<const Reference*>(&other)) {	4✔
77	return (this->reference_) == reference->reference_ && this->alignment_ == reference->alignment_;	4✔
78	} else {	4✔
79	return false;	×
80	}	×
81	};	4✔
82
83	std::string Reference::print() const { return "Reference(" + this->reference_->print() + ")"; };	1✔
84
85
86	std::string complex_type_name(types::PrimitiveType prim_type) {	28✔
87	switch (prim_type) {	28✔
88	case types::PrimitiveType::CHalf:	4✔
89	return "half2";	4✔
90	case types::PrimitiveType::CBFloat:	4✔
91	return "bfloat162";	4✔
92	case types::PrimitiveType::CFloat:	9✔
93	return "float2";	9✔
94	case types::PrimitiveType::CDouble:	6✔
95	return "double2";	6✔
96	case types::PrimitiveType::CFP128:	5✔
97	return "fp128_2";	5✔
NEW 98	default:	×
NEW 99	throw InvalidSDFGException("complex_type_name: not a complex primitive type");	×
100	}	28✔
101	};	28✔
102
103	std::string complex_op_suffix(types::PrimitiveType prim_type) {	9✔
104	switch (prim_type) {	9✔
NEW 105	case types::PrimitiveType::CHalf:	×
NEW 106	return "h";	×
NEW 107	case types::PrimitiveType::CBFloat:	×
NEW 108	return "bf";	×
109	case types::PrimitiveType::CFloat:	7✔
110	return "f";	7✔
111	case types::PrimitiveType::CDouble:	2✔
112	return "d";	2✔
NEW 113	case types::PrimitiveType::CFP128:	×
NEW 114	return "q";	×
NEW 115	default:	×
NEW 116	throw InvalidSDFGException("complex_op_suffix: not a complex primitive type");	×
117	}	9✔
118	};	9✔
119
120	std::string complex_tasklet(sdfg::Function& function, const data_flow::Tasklet& tasklet) {	9✔
121	// Determine the element suffix from the (complex) operand type.
122	auto& graph = tasklet.get_parent();	9✔
123	auto in_edges = graph.in_edges(tasklet);	9✔
124	if (in_edges.begin() == in_edges.end()) {	9✔
NEW 125	throw InvalidSDFGException("complex_tasklet: tasklet has no inputs");	×
NEW 126	}	×
127	auto operand_type = (*in_edges.begin()).result_type(function);	9✔
128	std::string suffix = complex_op_suffix(operand_type->primitive_type());	9✔
129
130	auto& inputs = tasklet.inputs();	9✔
131	switch (tasklet.code()) {	9✔
132	case data_flow::TaskletCode::complex_neg:	1✔
133	return "__daisy_cneg_" + suffix + "(" + inputs.at(0) + ")";	1✔
134	case data_flow::TaskletCode::complex_real:	1✔
135	return "__daisy_creal_" + suffix + "(" + inputs.at(0) + ")";	1✔
NEW 136	case data_flow::TaskletCode::complex_imag:	×
NEW 137	return "__daisy_cimag_" + suffix + "(" + inputs.at(0) + ")";	×
138	case data_flow::TaskletCode::complex_add:	3✔
139	return "__daisy_cadd_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";	3✔
NEW 140	case data_flow::TaskletCode::complex_sub:	×
NEW 141	return "__daisy_csub_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";	×
142	case data_flow::TaskletCode::complex_mul:	2✔
143	return "__daisy_cmul_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";	2✔
144	case data_flow::TaskletCode::complex_div:	1✔
145	return "__daisy_cdiv_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";	1✔
NEW 146	case data_flow::TaskletCode::complex_eq:	×
NEW 147	return "__daisy_ceq_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";	×
148	case data_flow::TaskletCode::complex_ne:	1✔
149	return "__daisy_cne_" + suffix + "(" + inputs.at(0) + ", " + inputs.at(1) + ")";	1✔
NEW 150	default:	×
NEW 151	throw InvalidSDFGException("complex_tasklet: not a complex tasklet code");	×
152	}	9✔
153	};	9✔
154
155	std::string complex_support_preamble(bool device) {	22✔
156	const std::string qual = device ? "__host__ __device__ static inline" : "static inline";	22✔
157
158	std::stringstream out;	22✔
159	out << "/* Complex type support (generated by sdfglib) */" << std::endl;	22✔
160
161	// Component vector types. On GPU, float2/double2/half2/__nv_bfloat162 are provided by the
162	// toolchain and reused; on CPU they are defined here. fp128_2 has no native type anywhere.
163	if (device) {	22✔
164	out << "typedef __nv_bfloat162 bfloat162;" << std::endl;	3✔
165	} else {	19✔
166	out << "typedef struct { float x; float y; } float2;" << std::endl;	19✔
167	out << "typedef struct { double x; double y; } double2;" << std::endl;	19✔
168	out << "typedef struct { _Float16 x; _Float16 y; } half2;" << std::endl;	19✔
169	out << "typedef struct { __bf16 x; __bf16 y; } bfloat162;" << std::endl;	19✔
170	}	19✔
171	out << "typedef struct { __float128 x; __float128 y; } fp128_2;" << std::endl;	22✔
172
173	// Element-wise helper functions. Arithmetic is carried out in a wide type (WT) so that
174	// narrow element types (half/bfloat) that lack native arithmetic are computed via float.
175	out << "#define __DAISY_DEFINE_COMPLEX(CT, ST, WT, SUF, QUAL) \\" << std::endl;	22✔
176	out << " QUAL CT __daisy_cadd_##SUF(CT a, CT b) { CT r; r.x = (ST)((WT)a.x + (WT)b.x); r.y = "	22✔
177	"(ST)((WT)a.y + (WT)b.y); return r; } \\"	22✔
178	<< std::endl;	22✔
179	out << " QUAL CT __daisy_csub_##SUF(CT a, CT b) { CT r; r.x = (ST)((WT)a.x - (WT)b.x); r.y = "	22✔
180	"(ST)((WT)a.y - (WT)b.y); return r; } \\"	22✔
181	<< std::endl;	22✔
182	out << " QUAL CT __daisy_cmul_##SUF(CT a, CT b) { CT r; r.x = (ST)((WT)a.x * (WT)b.x - (WT)a.y * (WT)b.y); "	22✔
183	"r.y = (ST)((WT)a.x * (WT)b.y + (WT)a.y * (WT)b.x); return r; } \\"	22✔
184	<< std::endl;	22✔
185	out << " QUAL CT __daisy_cdiv_##SUF(CT a, CT b) { CT r; WT d = (WT)b.x * (WT)b.x + (WT)b.y * (WT)b.y; "	22✔
186	"r.x = (ST)(((WT)a.x * (WT)b.x + (WT)a.y * (WT)b.y) / d); "	22✔
187	"r.y = (ST)(((WT)a.y * (WT)b.x - (WT)a.x * (WT)b.y) / d); return r; } \\"	22✔
188	<< std::endl;	22✔
189	out << " QUAL CT __daisy_cneg_##SUF(CT a) { CT r; r.x = (ST)(-(WT)a.x); r.y = (ST)(-(WT)a.y); return r; } \\"	22✔
190	<< std::endl;	22✔
191	out << " QUAL ST __daisy_creal_##SUF(CT a) { return a.x; } \\" << std::endl;	22✔
192	out << " QUAL ST __daisy_cimag_##SUF(CT a) { return a.y; } \\" << std::endl;	22✔
193	out << " QUAL bool __daisy_ceq_##SUF(CT a, CT b) { return (WT)a.x == (WT)b.x && (WT)a.y == (WT)b.y; } \\"	22✔
194	<< std::endl;	22✔
195	out << " QUAL bool __daisy_cne_##SUF(CT a, CT b) { return !((WT)a.x == (WT)b.x && (WT)a.y == (WT)b.y); }"	22✔
196	<< std::endl;	22✔
197
198	if (device) {	22✔
199	out << "__DAISY_DEFINE_COMPLEX(float2, float, float, f, " << qual << ")" << std::endl;	3✔
200	out << "__DAISY_DEFINE_COMPLEX(double2, double, double, d, " << qual << ")" << std::endl;	3✔
201	out << "__DAISY_DEFINE_COMPLEX(half2, __half, float, h, " << qual << ")" << std::endl;	3✔
202	out << "__DAISY_DEFINE_COMPLEX(bfloat162, __nv_bfloat16, float, bf, " << qual << ")" << std::endl;	3✔
203	// __float128 is a host-only type; keep fp128 helpers off the device.
204	out << "__DAISY_DEFINE_COMPLEX(fp128_2, __float128, __float128, q, __host__ inline)" << std::endl;	3✔
205	} else {	19✔
206	out << "__DAISY_DEFINE_COMPLEX(float2, float, float, f, " << qual << ")" << std::endl;	19✔
207	out << "__DAISY_DEFINE_COMPLEX(double2, double, double, d, " << qual << ")" << std::endl;	19✔
208	out << "__DAISY_DEFINE_COMPLEX(half2, _Float16, float, h, " << qual << ")" << std::endl;	19✔
209	out << "__DAISY_DEFINE_COMPLEX(bfloat162, __bf16, float, bf, " << qual << ")" << std::endl;	19✔
210	out << "__DAISY_DEFINE_COMPLEX(fp128_2, __float128, __float128, q, " << qual << ")" << std::endl;	19✔
211	}	19✔
212	out << "#undef __DAISY_DEFINE_COMPLEX" << std::endl;	22✔
213
214	return out.str();	22✔
215	};	22✔
216
217
218	} // namespace codegen
219	} // namespace sdfg

daisytuner / docc / 29934626444

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