Архитектурная акустика/Реализация/CUDA - История изменений

Андрей Чусов: /* По реализации */

2019-06-22T08:39:52Z

По реализации

← Предыдущая		Версия 18:39, 22 июня 2019
Строка 48:		Строка 48:

	==По реализации==		==По реализации==

			По автоматизации ввода-вывода см. статью [[Автоматизация ввода-вывода CUDA]]. Представленные ниже черновые наброски устарели, но пока оставлены на всякий случай.

	Пусть X - некоторый класс, который реализуется для CUDA устройства и для CPU. Поскольку не всегда есть гарантия совместимости представлений класса для CPU и GPU ([http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#windows-specific]), такие классы должны быть приводимы к формату, который можно передать ядру CUDA, с помощью метода <source lang = "cpp" inline>X::kernel_data get_device_object(bool async = false, cudaStream_t = 0) const;</source>, который должен возвращать объект вложенного типа <tt>X::kernel_data</tt>, удовлетворяющий требованиям {{cpp_concept\|StandardLayoutType}}. Этот метод, возможно перегруженный, должен		Пусть X - некоторый класс, который реализуется для CUDA устройства и для CPU. Поскольку не всегда есть гарантия совместимости представлений класса для CPU и GPU ([http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#windows-specific]), такие классы должны быть приводимы к формату, который можно передать ядру CUDA, с помощью метода <source lang = "cpp" inline>X::kernel_data get_device_object(bool async = false, cudaStream_t = 0) const;</source>, который должен возвращать объект вложенного типа <tt>X::kernel_data</tt>, удовлетворяющий требованиям {{cpp_concept\|StandardLayoutType}}. Этот метод, возможно перегруженный, должен

Alinap95: /* Балансировка нагрузки между CUDA и CPU */

2019-05-14T02:34:54Z

Балансировка нагрузки между CUDA и CPU

← Предыдущая		Версия 12:34, 14 мая 2019
Строка 44:		Строка 44:
	===Второй способ===		===Второй способ===
	Реализовать для второй пул потоков, которые передают задачу на CUDA и блокируются в ожидании завершения. В этом случае требуется также изменение класса [[Архитектурная акустика/Реализация/CLocalLoadBalancer\|CLocalLoadBalancer]] так, чтобы эти, дополнительные, потоки также были включены в общий пул - тогда шаблон [[Архитектурная акустика/Реализация/CLocalLoadBalancer\|CLocalLoadBalancer]] также будет зависить от политики выполнения.		Реализовать для второй пул потоков, которые передают задачу на CUDA и блокируются в ожидании завершения. В этом случае требуется также изменение класса [[Архитектурная акустика/Реализация/CLocalLoadBalancer\|CLocalLoadBalancer]] так, чтобы эти, дополнительные, потоки также были включены в общий пул - тогда шаблон [[Архитектурная акустика/Реализация/CLocalLoadBalancer\|CLocalLoadBalancer]] также будет зависить от политики выполнения.
			[[Файл:Cuda_node.jpg\|thumb\|500px\|Модификация node-ов]]
			Нужно реализовать Config::CudaCPUThreads().

	==По реализации==		==По реализации==

Андрей Чусов: /* По реализации */

2018-04-08T14:03:13Z

По реализации

← Предыдущая		Версия 00:03, 9 апреля 2018
Строка 89:		Строка 89:
	//...		//...
	};		};
	__host__ X(~~kernel_data~~&&, bool async = false, cudaStream_t = 0);		__host__ X(output_kernel_data&&, bool async = false, cudaStream_t = 0);
	};		};
	</source>		</source>

Андрей Чусов: kernel_type -> kernel_data

2018-04-08T12:13:41Z

kernel_type -> kernel_data

← Предыдущая		Версия 22:13, 8 апреля 2018
Строка 79:		Строка 79:
	//...		//...
	};		};
	__device__ X(~~kernel_type~~&&);		__device__ X(kernel_data&&);

	//Тип для считывания X из устройства		//Тип для считывания X из устройства
Строка 89:		Строка 89:
	//...		//...
	};		};
	__host__ X(~~output_kernel_type~~&&, bool async = false, cudaStream_t = 0);		__host__ X(kernel_data&&, bool async = false, cudaStream_t = 0);
	};		};
	</source>		</source>

Андрей Чусов: Скрыт cuda_error_info, т.к. это реализация; класс конструируется из rvalue kernel_data и rvalue output_kernel_data

2018-03-27T11:34:34Z

Скрыт cuda_error_info, т.к. это реализация; класс конструируется из rvalue kernel_data и rvalue output_kernel_data

← Предыдущая		Версия 21:34, 27 марта 2018
Строка 79:		Строка 79:
	//...		//...
	};		};
	__device__ X(~~const~~ kernel_type&);		__device__ X(kernel_type&&);

	//Тип для считывания X из устройства		//Тип для считывания X из устройства
Строка 89:		Строка 89:
	//...		//...
	};		};
	__host__ X(~~const~~ output_kernel_type&, bool async = false, cudaStream_t = 0);		__host__ X(output_kernel_type&&, bool async = false, cudaStream_t = 0);
	};		};
	</source>		</source>
Строка 97:		Строка 97:
	{{todo\|Описать из CudaUtils.h}}		{{todo\|Описать из CudaUtils.h}}

	* [[cuda_error_info]];
	* [[CUDA_CS_LOCK]] и [[CUDA_CS_UNLOCK]];		* [[CUDA_CS_LOCK]] и [[CUDA_CS_UNLOCK]];
	* [[cuda_abort_with_error]];		* [[cuda_abort_with_error]];

Андрей Чусов: /* По реализации */

2018-03-20T06:00:54Z

По реализации

← Предыдущая		Версия 16:00, 20 марта 2018
Строка 104:		Строка 104:

	{{example_begin}}		{{example_begin}}
	<source lang="~~cpp~~">		<source lang="cu">
	__global__ void my_kernel(bool fail, std::size_t host_thread_id, _Out_ unsigned* pVector)		__global__ void my_kernel(bool fail, std::size_t host_thread_id, _Out_ unsigned* pVector)
	{		{
	auto iThread = threadIdx.x + blockIdx.x * blockDim.x;		auto iThread = threadIdx.x + blockIdx.x * blockDim.x;
	if (fail && (iThread & 1))		if (fail && (iThread & 1))
	cuda_abort_with_error(-1, "CUDA error");		cuda_abort_with_error(-1, "CUDA error");
	pVector[iThread] = iThread;		pVector[iThread] = iThread;
	}		}

	void test(cudaStream_t stream)		void test(cudaStream_t stream)
	{		{

Андрей Чусов: /* По реализации */

2018-03-19T01:18:08Z

По реализации

← Предыдущая		Версия 11:18, 19 марта 2018
Строка 102:		Строка 102:
	* [[cuda_kernel_call]];		* [[cuda_kernel_call]];
	* [[cuda_stream_synchronize]].		* [[cuda_stream_synchronize]].

			{{example_begin}}
			<source lang="cpp">
			__global__ void my_kernel(bool fail, std::size_t host_thread_id, _Out_ unsigned* pVector)
			{
			auto iThread = threadIdx.x + blockIdx.x * blockDim.x;
			if (fail && (iThread & 1))
			cuda_abort_with_error(-1, "CUDA error");
			pVector[iThread] = iThread;
			}

			void test(cudaStream_t stream)
			{
			try
			{
			unsigned pVector[100];
			auto pDeviceVector = make_cuda_unique_ptr<unsigned>(sizeof(pVector));
			cuda_kernel_call(my_kernel_2, 10, 10, 0, stream, true, thread_id, pDeviceVector.get());
			cuda_runtime_call(cudaMemcpyAsync, pVector, pDeviceVector.get(), sizeof(pVector), cudaMemcpyDeviceToHost, stream);
			cuda_stream_synchronize(stream);
			}
			catch(cuda_user_exception& ex)
			{
			std::cerr << "Encountered an expected user-defined exception #" << std::hex << ex.code() << ": " << ex.what() << "\n";
			}
			catch(cuda_exception& ex)
			{
			std::cerr << "CUDA exception (code 0x" << std::hex << ex.code << ", description: " << ex.what() << ").\n";
			}
			catch (std::exception& ex)
			{
			std::cerr << "std::exception: " << ex.what() << "\n";
			}
			}
			</source>
			{{example_end}}

Андрей Чусов: /* По реализации */

2018-03-18T20:00:04Z

По реализации

← Предыдущая		Версия 06:00, 19 марта 2018
Строка 94:		Строка 94:

	Если класс <tt>X_dev</tt>, заданный для устройства, отличен от соответствующего ему класса <tt>X_host</TT>, заданного для хоста, то...		Если класс <tt>X_dev</tt>, заданный для устройства, отличен от соответствующего ему класса <tt>X_host</TT>, заданного для хоста, то...

			{{todo\|Описать из CudaUtils.h}}

			* [[cuda_error_info]];
			* [[CUDA_CS_LOCK]] и [[CUDA_CS_UNLOCK]];
			* [[cuda_abort_with_error]];
			* [[cuda_kernel_call]];
			* [[cuda_stream_synchronize]].

Андрей Чусов в 00:48, 15 марта 2018

2018-03-15T00:48:05Z

← Предыдущая		Версия 10:48, 15 марта 2018
Строка 71:		Строка 71:
	{		{
	//...		//...
			//Тип для записи X в устройство
	struct kernel_data		struct kernel_data
	{		{
Строка 79:		Строка 80:
	};		};
	__device__ X(const kernel_type&);		__device__ X(const kernel_type&);

			//Тип для считывания X из устройства
	struct output_kernel_data		struct output_kernel_data
	{		{
Строка 89:		Строка 92:
	};		};
	</source>		</source>

			Если класс <tt>X_dev</tt>, заданный для устройства, отличен от соответствующего ему класса <tt>X_host</TT>, заданного для хоста, то...

Андрей Чусов в 19:22, 14 марта 2018

2018-03-14T19:22:48Z

← Предыдущая		Версия 05:22, 15 марта 2018
Строка 64:		Строка 64:
	</source>		</source>
	Вложенный тип <tt>X::kernel_data</tt> должен определять тип <tt>X::kernel_data::class_type</tt>, эквивалентный <tt>X</tt>, а также выполняемый устройством метод <source lang = "cpp" inline>X::kernel_data::get_object</source> без параметров, который возвращает восстановленное для устройства представление исходного экземпляра X.		Вложенный тип <tt>X::kernel_data</tt> должен определять тип <tt>X::kernel_data::class_type</tt>, эквивалентный <tt>X</tt>, а также выполняемый устройством метод <source lang = "cpp" inline>X::kernel_data::get_object</source> без параметров, который возвращает восстановленное для устройства представление исходного экземпляра X.

			''Замена''

			<source lang = "cu">
			struct X
			{
			//...
			struct kernel_data
			{
			typedef X class_type;
			__host__ kernel_data(const class_type&, bool async = false, cudaStream_t = 0);
			__host__ kernel_data(class_type&&, bool async = false, cudaStream_t = 0); //Опционально
			//...
			};
			__device__ X(const kernel_type&);
			struct output_kernel_data
			{
			typedef X class_type;
			__device__ output_kernel_data(const class_type&);
			__device__ output_kernel_data(class_type&&); //Опционально
			//...
			};
			__host__ X(const output_kernel_type&, bool async = false, cudaStream_t = 0);
			};
			</source>