Partial functionality for hip complex types. WIP. (f1346ebc) · Commits · Cianciosa, Mark / graph_framework

graph_framework/hip_context.hpp

+1 −3

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

		std::array<const char *, 2> options({
		"-std=c++20",
		"-std=c++17",
		"-I" HEADER_DIR
		});

		@@ -462,7 +462,6 @@ namespace gpu {

		source_buffer << " const int index = blockIdx.x*blockDim.x + threadIdx.x;"
		<< std::endl;
		source_buffer << " printf(%i %i %i %i, index, blockIdx.x, blockDim.x, threadIdx.x);" << std::endl;
		source_buffer << " if (index < " << size << ") {" << std::endl;

		for (auto &input : inputs) {
		@@ -557,7 +556,6 @@ namespace gpu {
		source_buffer << " const unsigned int i = threadIdx.x;" << std::endl;
		source_buffer << " const unsigned int j = threadIdx.x/64;" << std::endl;
		source_buffer << " const unsigned int k = threadIdx.x%64;" << std::endl;
		source_buffer << " printf(%i %i %i %i %i %i, i, j, k, blockIdx.x, blockDim.x, threadIdx.x);" << std::endl;
		source_buffer << " if (i < " << size << ") {" << std::endl;
		source_buffer << " " << jit::type_to_string<T> () << " sub_max = ";
		if constexpr (jit::is_complex<T> ()) {

graph_framework/piecewise.hpp

+53 −13

Original line number	Diff line number	Diff line
		@@ -33,7 +33,11 @@ void compile_index(std::ostringstream &stream,
		<< type
		<< ")";
		if constexpr (jit::complex_scalar<T>) {
		stream << "real(";
		stream << "real";
		if constexpr (jit::use_hip()) {
		stream << "<" << jit::type_to_string<T> () << "> ";
		}
		stream << "(";
		}
		stream << "((" << register_name << " - ";
		if constexpr (jit::complex_scalar<T>) {
		@@ -255,20 +259,38 @@ void compile_index(std::ostringstream &stream,
		avail_const_mem -= buffer_size;
		stream << "__constant__ ";
		}
		} else if (jit::use_hip()) {
		stream << "__device__ ";
		}
		stream << "const ";
		jit::add_type<T> (stream);
		stream << " " << registers[leaf_node<T, SAFE_MATH>::caches.backends[data_hash].data()] << "[] = {";
		if constexpr (!jit::use_hip()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		}
		stream << leaf_node<T, SAFE_MATH>::caches.backends[data_hash][0];
		} else {
		stream << "{"
		<< std::real(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][0])
		<< ","
		<< std::imag(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][0])
		<< "}";
		}
		for (size_t i = 1; i < length; i++) {
		stream << ", ";
		if constexpr (!jit::use_hip()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		}
		stream << leaf_node<T, SAFE_MATH>::caches.backends[data_hash][i];
		} else {
		stream << "{"
		<< std::real(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][i])
		<< ","
		<< std::imag(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][i])
		<< "}";
		}
		}
		stream << "};" << std::endl;
		}
		@@ -919,20 +941,38 @@ void compile_index(std::ostringstream &stream,
		avail_const_mem -= buffer_size;
		stream << "__constant__ ";
		}
		} else if (jit::use_hip()) {
		stream << "__device__ ";
		}
		stream << "const ";
		jit::add_type<T> (stream);
		stream << " " << registers[leaf_node<T, SAFE_MATH>::caches.backends[data_hash].data()] << "[] = {";
		if (!jit::use_hip()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		}
		stream << leaf_node<T, SAFE_MATH>::caches.backends[data_hash][0];
		} else {
		stream << "{"
		<< std::real(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][0])
		<< ","
		<< std::imag(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][0])
		<< "}";
		}
		for (size_t i = 1; i < length; i++) {
		stream << ", ";
		if constexpr (!jit::use_hip()) {
		if constexpr (jit::complex_scalar<T>) {
		jit::add_type<T> (stream);
		}
		stream << leaf_node<T, SAFE_MATH>::caches.backends[data_hash][i];
		} else {
		stream << "{"
		<< std::real(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][i])
		<< ","
		<< std::imag(leaf_node<T, SAFE_MATH>::caches.backends[data_hash][i])
		<< "}";
		}
		}
		stream << "};" << std::endl;
		}

graph_framework/register.hpp

+1 −1

Original line number	Diff line number	Diff line
		@@ -157,7 +157,7 @@ namespace jit {
		if constexpr (use_cuda()) {
		return "cuda::std::complex<" + type_to_string<T> () + ">";
		} else {
		return "std::complex<" + type_to_string<T> () + ">";
		return "complex<" + type_to_string<T> () + ">";
		}
		} else {
		return type_to_string<T> ();

graph_framework/special_functions.hpp

+526 −84

File changed.

Preview size limit exceeded, changes collapsed.

graph_tests/jit_test.cpp

+2 −2

Original line number	Diff line number	Diff line
		@@ -271,7 +271,7 @@ template<jit::float_scalar T> void run_math_tests() {
		{exp_node_dfdv2}, {},
		exp_node_dfdv2->evaluate().at(0), 2.0E-6);

		if constexpr (jit::use_cuda() \|\| !jit::use_gpu<T> ()) {
		if constexpr (jit::use_cuda() \|\| jit::use_hip() \|\| !jit::use_gpu<T> ()) {
		result = 1.8E-15;
		} else {
		result = 0.0;
		@@ -296,7 +296,7 @@ template<jit::float_scalar T> void run_math_tests() {
		{pow_node_dfdv1}, {},
		pow_node_dfdv1->evaluate().at(0), 0.0);

		if constexpr (jit::use_cuda() \|\| !jit::use_gpu<T> ()) {
		if constexpr (jit::use_cuda() \|\| jit::use_hip() \|\| !jit::use_gpu<T> ()) {
		result = 8.9E-16;
		} else {
		result = 0.0;