Merge branch 'cleanup' into 'main' (b94c3ff6) · Commits · Cianciosa, Mark / graph_framework

graph_framework/arithmetic.hpp

+8 −9

Original line number	Diff line number	Diff line
		@@ -893,7 +893,6 @@ namespace graph {
		// Idenity reductions.
		auto l = constant_cast(this->left);
		if (this->left->is_match(this->right)) {
		auto l = constant_cast(this->left);
		if (l.get() && l->is(0)) {
		return this->left;
		}
		@@ -2502,21 +2501,21 @@ namespace graph {
		stream << " " << registers[this] << " = ";
		if constexpr (SAFE_MATH) {
		stream << "(" << registers[l.get()] << " == ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		stream << "0";
		}
		stream << " \|\| " << registers[r.get()] << " == ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		stream << "0";
		}
		stream << ") ? ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		@@ -3494,14 +3493,14 @@ namespace graph {
		stream << " " << registers[this] << " = ";
		if constexpr (SAFE_MATH) {
		stream << registers[l.get()] << " == ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		stream << "0";
		}
		stream << " ? ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		@@ -5069,14 +5068,14 @@ namespace graph {
		stream << " " << registers[this] << " = ";
		if constexpr (SAFE_MATH) {
		stream << "(" << registers[l.get()] << " == ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		stream << "0";
		}
		stream << " \|\| " << registers[m.get()] << " == ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(0, 0)";
		} else {
		@@ -5084,7 +5083,7 @@ namespace graph {
		}
		stream << ") ? " << registers[r.get()] << " : ";
		}
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		stream << registers[l.get()] << "*"
		<< registers[m.get()] << " + "
		<< registers[r.get()];

graph_framework/backend.hpp

+11 −12

Original line number	Diff line number	Diff line
		@@ -255,9 +255,8 @@ namespace backend {
		//------------------------------------------------------------------------------
		/// @brief Take erfi.
		//------------------------------------------------------------------------------
		template<jit::float_scalar D=T>
		typename std::enable_if<jit::is_complex<D> (), void>::type erfi() {
		for (D &d : memory) {
		void erfi() requires(jit::complex_scalar<T>) {
		for (T &d : memory) {
		d = special::erfi(d);
		}
		}
		@@ -278,7 +277,7 @@ namespace backend {
		//------------------------------------------------------------------------------
		bool is_normal() const {
		for (const T &x : memory) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		if (std::isnan(std::real(x)) \|\| std::isinf(std::real(x)) \|\|
		std::isnan(std::imag(x)) \|\| std::isinf(std::imag(x))) {
		return false;
		@@ -597,7 +596,7 @@ namespace backend {
		const size_t num_rows = x.size();
		for (size_t i = 0; i < num_rows; i++) {
		for (size_t j = 0; j < num_colmns; j++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		memory[inum_colmns + j] = std::atan(x[i]/memory[inum_colmns + j]);
		} else {
		memory[inum_colmns + j] = std::atan2(x[i], memory[inum_colmns + j]);
		@@ -613,7 +612,7 @@ namespace backend {
		const size_t num_rows = size();
		for (size_t i = 0; i < num_rows; i++) {
		for (size_t j = 0; j < num_colmns; j++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		m[inum_colmns + j] = std::atan(x[inum_colmns + j]/memory[i]);
		} else {
		m[inum_colmns + j] = std::atan2(x[inum_colmns + j], memory[i]);
		@@ -641,7 +640,7 @@ namespace backend {
		const size_t num_rows = x.size();
		for (size_t i = 0; i < num_colmns; i++) {
		for (size_t j = 0; j < num_rows; j++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		memory[inum_colmns + j] = std::atan(x[j]/memory[inum_colmns + j]);
		} else {
		memory[inum_colmns + j] = std::atan2(x[j], memory[inum_colmns + j]);
		@@ -657,7 +656,7 @@ namespace backend {
		const size_t num_rows = size();
		for (size_t i = 0; i < num_rows; i++) {
		for (size_t j = 0; j < num_colmns; j++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		m[inum_colmns + j] = std::atan(x[inum_colmns + j]/memory[j]);
		} else {
		m[inum_colmns + j] = std::atan2(x[inum_colmns + j], memory[j]);
		@@ -918,7 +917,7 @@ namespace backend {
		inline buffer<T> fma(buffer<T> &a,
		buffer<T> &b,
		buffer<T> &c) {
		constexpr bool use_fma = !jit::is_complex<T> () &&
		constexpr bool use_fma = !jit::complex_scalar<T> &&
		#ifdef FP_FAST_FMA
		true;
		#else
		@@ -1100,7 +1099,7 @@ namespace backend {
		if (y.size() == 1) {
		const T right = y.at(0);
		for (size_t i = 0, ie = x.size(); i < ie; i++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		x[i] = std::atan(right/x[i]);
		} else {
		x[i] = std::atan2(right, x[i]);
		@@ -1110,7 +1109,7 @@ namespace backend {
		} else if (x.size() == 1) {
		const T left = x.at(0);
		for (size_t i = 0, ie = y.size(); i < ie; i++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		y[i] = std::atan(y[i]/left);
		} else {
		y[i] = std::atan2(y[i], left);
		@@ -1122,7 +1121,7 @@ namespace backend {
		assert(x.size() == y.size() &&
		"Left and right sizes are incompatable.");
		for (size_t i = 0, ie = x.size(); i < ie; i++) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		x[i] = std::atan(y[i]/x[i]);
		} else {
		x[i] = std::atan2(y[i], x[i]);

graph_framework/cpu_context.hpp

+5 −5

Original line number	Diff line number	Diff line
		@@ -311,7 +311,7 @@ namespace gpu {

		return [run, begin, end] () mutable {
		run();
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		return *std::max_element(begin, end,
		[] (const T a, const T b) {
		return std::abs(a) < std::abs(b);
		@@ -346,7 +346,7 @@ namespace gpu {
		const graph::output_nodes<T, SAFE_MATH> &nodes) {
		for (auto &out : nodes) {
		const T temp = kernel_arguments[out.get()][index];
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		std::cout << std::real(temp) << " " << std::imag(temp) << " ";
		} else {
		std::cout << temp << " ";
		@@ -401,7 +401,7 @@ namespace gpu {
		void create_header(std::ostringstream &source_buffer) {
		source_buffer << "#include <map>" << std::endl
		<< "#include <array>" << std::endl;
		if (jit::is_complex<T> ()) {
		if (jit::complex_scalar<T>) {
		source_buffer << "#include <complex>" << std::endl;
		source_buffer << "#include <special_functions.hpp>" << std::endl;
		} else {
		@@ -517,7 +517,7 @@ namespace gpu {
		source_buffer << " " << jit::to_string('v', in.get());
		source_buffer << "[i] = ";
		if constexpr (SAFE_MATH) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (source_buffer);
		source_buffer << " (";
		source_buffer << "isnan(real(" << registers[a.get()]
		@@ -543,7 +543,7 @@ namespace gpu {
		source_buffer << " " << jit::to_string('o', out.get());
		source_buffer << "[i] = ";
		if constexpr (SAFE_MATH) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (source_buffer);
		source_buffer << " (";
		source_buffer << "isnan(real(" << registers[a.get()]

graph_framework/cuda_context.hpp

+14 −14

Original line number	Diff line number	Diff line
		@@ -387,16 +387,16 @@ namespace gpu {
		texture_desc.addressMode[0] = CU_TR_ADDRESS_MODE_BORDER;
		texture_desc.addressMode[1] = CU_TR_ADDRESS_MODE_BORDER;
		texture_desc.addressMode[2] = CU_TR_ADDRESS_MODE_BORDER;
		if constexpr (jit::is_float<T> ()) {
		if constexpr (jit::float_base<T>) {
		array_desc.Format = CU_AD_FORMAT_FLOAT;
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		array_desc.NumChannels = 2;
		} else {
		array_desc.NumChannels = 1;
		}
		} else {
		array_desc.Format = CU_AD_FORMAT_UNSIGNED_INT32;
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		array_desc.NumChannels = 4;
		} else {
		array_desc.NumChannels = 2;
		@@ -433,16 +433,16 @@ namespace gpu {
		texture_desc.addressMode[1] = CU_TR_ADDRESS_MODE_BORDER;
		texture_desc.addressMode[2] = CU_TR_ADDRESS_MODE_BORDER;
		const size_t total = size[0]*size[1];
		if constexpr (jit::is_float<T> ()) {
		if constexpr (jit::float_base<T>) {
		array_desc.Format = CU_AD_FORMAT_FLOAT;
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		array_desc.NumChannels = 2;
		} else {
		array_desc.NumChannels = 1;
		}
		} else {
		array_desc.Format = CU_AD_FORMAT_UNSIGNED_INT32;
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		array_desc.NumChannels = 4;
		} else {
		array_desc.NumChannels = 2;
		@@ -582,7 +582,7 @@ namespace gpu {
		wait();
		for (auto &out : nodes) {
		const T temp = reinterpret_cast<T *> (kernel_arguments[out.get()])[index];
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		std::cout << std::real(temp) << " " << std::imag(temp) << " ";
		} else {
		std::cout << temp << " ";
		@@ -651,13 +651,13 @@ namespace gpu {
		<< " return _buffer[index];" << std::endl
		<< " }" << std::endl
		<< "};" << std::endl;
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		source_buffer << "#define CUDA_DEVICE_CODE" << std::endl
		<< "#define M_PI " << M_PI << std::endl
		<< "#include <cuda/std/complex>" << std::endl
		<< "#include <special_functions.hpp>" << std::endl;
		#ifdef USE_CUDA_TEXTURES
		if constexpr (jit::is_float<T> ()) {
		if constexpr (jit::float_base<T>) {
		source_buffer << "static __inline__ __device__ complex<float> to_cmp_float(float2 p) {"
		<< std::endl
		<< " return ";
		@@ -673,7 +673,7 @@ namespace gpu {
		<< std::endl
		<< "}" << std::endl;
		}
		} else if constexpr (jit::is_double<T> ()) {
		} else if constexpr (jit::double_base<T>) {
		source_buffer << "static __inline__ __device__ double to_double(uint2 p) {"
		<< std::endl
		<< " return __hiloint2double(p.y, p.x);"
		@@ -858,7 +858,7 @@ namespace gpu {
		}
		source_buffer << "index] = ";
		if constexpr (SAFE_MATH) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (source_buffer);
		source_buffer << " (";
		source_buffer << "isnan(real(" << registers[a.get()]
		@@ -889,7 +889,7 @@ namespace gpu {
		}
		source_buffer << "index] = ";
		if constexpr (SAFE_MATH) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (source_buffer);
		source_buffer << " (";
		source_buffer << "isnan(real(" << registers[a.get()]
		@@ -932,13 +932,13 @@ namespace gpu {
		source_buffer << " const unsigned int k = threadIdx.x%32;" << std::endl;
		source_buffer << " if (i < " << size << ") {" << std::endl;
		source_buffer << " " << jit::type_to_string<T> () << " sub_max = ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		source_buffer << "abs(input[i]);" << std::endl;
		} else {
		source_buffer << "input[i];" << std::endl;
		}
		source_buffer << " for (size_t index = i + 1024; index < " << size <<"; index += 1024) {" << std::endl;
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		source_buffer << " sub_max = max(sub_max, abs(input[index]));" << std::endl;
		} else {
		source_buffer << " sub_max = max(sub_max, input[index]);" << std::endl;

graph_framework/math.hpp

+5 −9

Original line number	Diff line number	Diff line
		@@ -440,11 +440,11 @@ namespace graph {
		jit::add_type<T> (stream);
		stream << " " << registers[this] << " = ";
		if constexpr (SAFE_MATH) {
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		stream << "real(";
		}
		stream << registers[a.get()];
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		stream << ")";
		}
		stream << " < 709.8 ? ";
		@@ -452,16 +452,12 @@ namespace graph {
		stream << "exp(" << registers[a.get()] << ")";
		if constexpr (SAFE_MATH) {
		stream << " : ";
		if constexpr (jit::is_complex<T> ()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		stream << "(";
		}
		if constexpr (jit::is_float<T> ()) {
		stream << std::numeric_limits<float>::max();
		} else {
		stream << std::numeric_limits<double>::max();
		}
		if constexpr (jit::is_complex<T> ()) {
		stream << jit::max_base<T> ();
		if constexpr (jit::complex_scalar<T>) {
		stream << ")";
		}
		}