void sparse_fully_connected_1x1_layer_tester_cuda::enqueue_forward_propagation(
cudaStream_t stream_id,
cuda_linear_buffer_device::ptr output_buffer,
const std::vector<cuda_linear_buffer_device::const_ptr>& schema_data,
const std::vector<cuda_linear_buffer_device::const_ptr>& data,
const std::vector<cuda_linear_buffer_device::const_ptr>& data_custom,
const std::vector<cuda_linear_buffer_device::const_ptr>& input_buffers,
const std::vector<cuda_linear_buffer_device::const_ptr>& persistent_working_data,
cuda_linear_buffer_device::ptr temporary_working_fixed_buffer,
cuda_linear_buffer_device::ptr temporary_working_per_entry_buffer,
unsigned int entry_count)
{
{
cusparse_safe_call(cusparseSetStream(cuda_config->get_cusparse_handle(), stream_id));
float alpha = 1.0F;
float beta = 0.0F;
cusparseMatDescr_t mat_descr;
cusparse_safe_call(cusparseCreateMatDescr(&mat_descr));
cusparse_safe_call(cusparseScsrmm(
cuda_config->get_cusparse_handle(),
CUSPARSE_OPERATION_NON_TRANSPOSE,
output_elem_count_per_entry,
entry_count,
input_elem_count_per_entry_list[0],
feature_map_connection_count,
&alpha,
mat_descr,
*data[0],
*data_custom[1],
*data_custom[0],
*input_buffers[0],
input_elem_count_per_entry_list[0],
&beta,
*output_buffer,
output_elem_count_per_entry));
}
// Add bias
{
cudnn_safe_call(cudnnSetStream(cuda_config->get_cudnn_handle(), stream_id));
cudnn_util::set_tensor_descriptor(
output_data_desc,
output_configuration_specific,
entry_count);
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnAddTensor(
cuda_config->get_cudnn_handle(),
&alpha,
bias_desc,
*data[1],
&beta,
output_data_desc,
*output_buffer));
}
}
void cuda_running_configuration::update_parameters()
{
cuda_safe_call(cudaDriverGetVersion(&driver_version));
cuda_safe_call(cudaRuntimeGetVersion(&runtime_version));
int device_count;
cuda_safe_call(cudaGetDeviceCount(&device_count));
if (device_count <= 0)
throw neural_network_exception("No CUDA capable devices are found");
if (device_id >= device_count)
throw neural_network_exception((boost::format("Device ID %1% specified while %2% devices are available") % device_id % device_count).str());
cudaDeviceProp device_prop;
cuda_safe_call(cudaGetDeviceProperties(&device_prop, device_id));
device_name = device_prop.name;
compute_capability_major = device_prop.major;
compute_capability_minor = device_prop.minor;
clock_rate = device_prop.clockRate;
memory_clock_rate = device_prop.memoryClockRate;
memory_bus_width = device_prop.memoryBusWidth;
global_memory_size = device_prop.totalGlobalMem;
ecc_enabled = (device_prop.ECCEnabled != 0);
l2_cache_size = device_prop.l2CacheSize;
multiprocessor_count = device_prop.multiProcessorCount;
smem_per_block = device_prop.sharedMemPerBlock;
max_threads_per_multiprocessor = device_prop.maxThreadsPerMultiProcessor;
max_threads_per_block = device_prop.maxThreadsPerBlock;
for(int i = 0; i < sizeof(max_threads_dim) / sizeof(max_threads_dim[0]); ++i)
max_threads_dim[i] = device_prop.maxThreadsDim[i];
for(int i = 0; i < sizeof(max_grid_size) / sizeof(max_grid_size[0]); ++i)
max_grid_size[i] = device_prop.maxGridSize[i];
max_texture_1d_linear = device_prop.maxTexture1DLinear;
texture_alignment = device_prop.textureAlignment;
pci_bus_id = device_prop.pciBusID;
pci_device_id = device_prop.pciDeviceID;
#ifdef _WIN32
tcc_mode = (device_prop.tccDriver != 0);
#endif
cuda_safe_call(cudaSetDevice(device_id));
cublas_safe_call(cublasCreate(&cublas_handle));
cusparse_safe_call(cusparseCreate(&cusparse_handle));
}
void sparse_1x1_layer_tester_cuda::enqueue_forward_propagation(
cudaStream_t stream_id,
cuda_linear_buffer_device::ptr output_buffer,
const std::vector<cuda_linear_buffer_device::const_ptr>& schema_data,
const std::vector<cuda_linear_buffer_device::const_ptr>& data,
const std::vector<cuda_linear_buffer_device::const_ptr>& data_custom,
const std::vector<cuda_linear_buffer_device::const_ptr>& input_buffers,
const std::vector<cuda_linear_buffer_device::const_ptr>& persistent_working_data,
cuda_linear_buffer_device::ptr temporary_working_fixed_buffer,
cuda_linear_buffer_device::ptr temporary_working_per_entry_buffer,
unsigned int entry_count)
{
// Convert input data strided NCHW to packed CNHW format
if (unit_stride)
{
cuda_util::transpose23(
*cuda_config,
*input_buffers[0],
*temporary_working_per_entry_buffer,
input_elem_count_per_feature_map_list[0],
input_configuration_specific_list[0].feature_map_count,
entry_count,
stream_id);
}
else
{
std::vector<unsigned int> input_converted_CNHW_strides = input_converted_CNHW_strides_base;
input_converted_CNHW_strides[input_converted_CNHW_strides.size() - 2] = input_converted_CNHW_strides[input_converted_CNHW_strides.size() - 1] * entry_count;
cudnn_safe_call(cudnnSetStream(cuda_config->get_cudnn_handle(), stream_id));
cudnn_util::set_tensor_descriptor(
input_strided_data_desc,
input_strided_config,
entry_count,
input_strides);
cudnn_util::set_tensor_descriptor(
input_converted_CNHW_data_desc,
input_strided_config,
entry_count,
input_converted_CNHW_strides);
float alpha = 1.0F;
float beta = 0.0F;
cudnn_safe_call(cudnnAddTensor(
cuda_config->get_cudnn_handle(),
&alpha,
input_strided_data_desc,
*input_buffers[0],
&beta,
input_converted_CNHW_data_desc,
*temporary_working_per_entry_buffer));
}
{
cusparse_safe_call(cusparseSetStream(cuda_config->get_cusparse_handle(), stream_id));
float alpha = 1.0F;
float beta = 0.0F;
cusparseMatDescr_t mat_descr;
cusparse_safe_call(cusparseCreateMatDescr(&mat_descr));
cusparse_safe_call(cusparseScsrmm2(
cuda_config->get_cusparse_handle(),
CUSPARSE_OPERATION_NON_TRANSPOSE,
CUSPARSE_OPERATION_TRANSPOSE,
output_configuration_specific.feature_map_count,
entry_count * output_elem_count_per_feature_map,
input_strided_config.feature_map_count,
feature_map_connection_count,
&alpha,
mat_descr,
*data[0],
*data_custom[1],
*data_custom[0],
*temporary_working_per_entry_buffer,
entry_count * output_elem_count_per_feature_map,
&beta,
((float *)*temporary_working_per_entry_buffer) + input_converted_elem_count_per_entry_aligned * entry_count,
output_configuration_specific.feature_map_count));
}
// Convert output from NHWC to NCHW
{
cuda_util::transpose(
*cuda_config,
((float *)*temporary_working_per_entry_buffer) + input_converted_elem_count_per_entry_aligned * entry_count,
*output_buffer,
output_configuration_specific.feature_map_count,
output_elem_count_per_feature_map,
entry_count,
stream_id);
}
// Add bias
if (bias)
{
cudnn_safe_call(cudnnSetStream(cuda_config->get_cudnn_handle(), stream_id));
cudnn_util::set_tensor_descriptor(
output_data_desc,
output_configuration_specific,
entry_count);
float alpha = 1.0F;
float beta = 1.0F;
cudnn_safe_call(cudnnAddTensor(
cuda_config->get_cudnn_handle(),
&alpha,
bias_desc,
*data[1],
&beta,
output_data_desc,
*output_buffer));
}
}
