| 2861 | } |
| 2862 | |
| 2863 | port::Status MIOpenSupport::DoConvolve( |
| 2864 | dnn::ConvolutionKind kind, dnn::DataType element_type, |
| 2865 | dnn::DataType output_type, Stream* stream, |
| 2866 | const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data, |
| 2867 | const dnn::FilterDescriptor& filter_descriptor, |
| 2868 | DeviceMemoryBase filter_data, const dnn::BatchDescriptor& output_descriptor, |
| 2869 | DeviceMemoryBase output_data, |
| 2870 | const dnn::ConvolutionDescriptor& convolution_descriptor, |
| 2871 | dnn::AlgorithmDesc algorithm_desc, DeviceMemory<uint8> scratch_memory, |
| 2872 | dnn::ProfileResult* output_profile_result) { |
| 2873 | auto miopen = miopen_->GetHandle(parent_, stream); |
| 2874 | ScopedTensorDescriptor input_nd{input_descriptor, |
| 2875 | ToMIOpenDataType(element_type)}; |
| 2876 | ScopedTensorDescriptor output_nd{output_descriptor, |
| 2877 | ToMIOpenDataType(element_type)}; |
| 2878 | ScopedFilterDescriptor filter{filter_descriptor, input_descriptor, |
| 2879 | ToMIOpenDataType(element_type)}; |
| 2880 | ScopedConvolutionDescriptor conv{convolution_descriptor, |
| 2881 | ToMIOpenDataType(element_type)}; |
| 2882 | |
| 2883 | // Alpha is the scaling factor for input. |
| 2884 | float alpha = 1.0; |
| 2885 | // Beta is the scaling factor for output. |
| 2886 | float beta = 0.0; |
| 2887 | |
| 2888 | const bool is_profiling = output_profile_result != nullptr; |
| 2889 | |
| 2890 | std::unique_ptr<GpuTimer> timer; |
| 2891 | if (is_profiling) { |
| 2892 | timer.reset(new GpuTimer(parent_)); |
| 2893 | if (!timer->Init()) { |
| 2894 | return port::Status(port::error::INTERNAL, "Failed to init timer"); |
| 2895 | } |
| 2896 | // The start and stop of the timer should be as close to the MIOpen call as |
| 2897 | // possible. It is still possible for other threads to issue workload on |
| 2898 | // to this stream. So it could take multiple profiling measurements. |
| 2899 | if (!timer->Start(AsGpuStream(stream))) { |
| 2900 | timer->Destroy(); |
| 2901 | return port::Status(port::error::INTERNAL, "Failed to start timer"); |
| 2902 | } |
| 2903 | } |
| 2904 | |
| 2905 | miopenStatus_t status = miopenStatusSuccess; |
| 2906 | switch (kind) { |
| 2907 | case dnn::ConvolutionKind::FORWARD: { |
| 2908 | status = wrap::miopenConvolutionForward( |
| 2909 | miopen.handle(), |
| 2910 | /*alpha=*/&alpha, /*srcDesc=*/input_nd.handle(), |
| 2911 | /*srcData=*/input_data.opaque(), /*filterDesc=*/filter.handle(), |
| 2912 | /*filterData=*/filter_data.opaque(), /*convDesc=*/conv.handle(), |
| 2913 | /*algo=*/ |
| 2914 | static_cast<miopenConvFwdAlgorithm_t>(algorithm_desc.algo_id()), |
| 2915 | /*beta=*/&beta, /*destDesc=*/output_nd.handle(), |
| 2916 | /*destData=*/output_data.opaque(), |
| 2917 | /*workSpace=*/scratch_memory.opaque(), |
| 2918 | /*workSpaceSizeInBytes=*/scratch_memory.size()); |
| 2919 | break; |
| 2920 | } |
nothing calls this directly
no test coverage detected