Immediate mode#

MIOpen v2.0 introduces immediate mode, which removes the requirement for miopenFindConvolution*() calls, thereby reducing runtime costs. In this mode, you can query the MIOpen runtime for all of the supported solutions for a given convolution configuration. The sequence of operations for immediate mode is similar to launching regular convolutions in MIOpen, for instance, through the miopenFindConvolution*() API. However, in this case, the different APIs have a lower runtime cost.

A typical convolution call is similar to the following sequence:

• Construct the MIOpen handle and relevant descriptors, such as the convolution descriptor.

• With the above data structures, call miopenConvolution*GetSolutionCount to get the maximum number of supported solutions for the convolution descriptor.

• Use the obtained count to allocate memory for the miopenConvSolution_t structure (introduced in MIOpen v2.0).

• Call miopenConvolution*GetSolution to populate the miopenConvSolution_t structures allocated above. The returned list is sorted in order of best performance, where the first element is the fastest.

• While the above structure returns the amount of workspace required for an algorithm, you can query the amount of a workspace required for a known solution ID using miopenConvolution*GetSolutionWorkspaceSize. However, this is not a requirement because the structure returned by miopenConvolution*GetSolution already has this information.

• Initiate the convolution operation in immediate mode by calling miopenConvolution*Immediate. This populates the output tensor descriptor with the respective convolution result. However, the first call to miopenConvolution*Immediate might take more time because the kernel must be compiled if it isn’t present in the kernel cache.

• Optionally, you can compile the solution of choice by calling miopenConvolution*CompileSolution. This ensures that the kernel represented by the chosen solution is populated in the kernel cache, removing the need to compile it.

miopenConvolutionForwardGetSolutionCount(handle,
                                        weightTensorDesc,
                                        inputTensorDesc,
                                        convDesc,
                                        outputTensorDesc,
                                        &solutionCount);


// < allocate an array of miopenConvSolution_t of size solutionCount >


miopenConvolutionForwardGetSolution(handle,
                                    weightTensorDesc,
                                    inputTensorDesc,
                                    convDesc,
                                    outputTensorDesc,
                                    solutionCount,
                                    &actualCount,
                                    solutions);

// < select a solution from solutions array >

miopenConvolutionForwardGetSolutionWorkspaceSize(handle,
                                                weightTensorDesc,
                                                inputTensorDesc,
                                                convDesc,
                                                outputTensorDesc,
                                                selected->solution_id,
                                                &ws_size);

// < allocate solution workspace of size ws_size >


// This stage is optional.
miopenConvolutionForwardCompileSolution(handle,
                                        weightTensorDesc,
                                        inputTensorDesc,
                                        convDesc,
                                        outputTensorDesc,
                                        selected->solution_id);



miopenConvolutionForwardImmediate(handle,
                                  weightTensor,
                                  weight_device_mem,
                                  inputTensorDesc,
                                  input_device_mem,
                                  convDesc,
                                  outputTensorDesc,
                                  output_device_mem,
                                  workspace_device_mem,
                                  ws_size,
                                  selected->solution_id);

Find modes#

MIOpen provides a set of find modes that are used to accelerate find API calls. Set the different modes by using the MIOPEN_FIND_MODE environment variable with one of these values:

• NORMAL/1 (normal find): This is the full find mode call, which benchmarks all the solvers and returns a list.

• FAST/2 (fast find): Checks FindDb for an entry. If there’s a FindDb hit, it uses that entry. If there’s a miss, it uses the immediate mode fallback. This mode offers fast start-up times at the cost of GPU performance.

• HYBRID/3 or unset MIOPEN_FIND_MODE (hybrid find): Checks FindDb for an entry. If there’s a FindDb hit, it uses that entry. If there’s a miss, it uses the existing find machinery. This mode offers slower start-up times than fast find without the GPU performance drop.

• 4: This value is reserved and should not be used.

• DYNAMIC_HYBRID/5 (dynamic hybrid find): Checks FindDb for an entry. If there’s a FindDb hit, it uses that entry. If there’s a miss, it uses the existing find machinery, skipping non-dynamic kernels. It offers faster start-up times than hybrid find, but GPU performance might decrease.

The default find mode is DYNAMIC_HYBRID. To run the full NORMAL find mode, use export MIOPEN_FIND_MODE=NORMAL or export MIOPEN_FIND_MODE=1.