# ---
# jupyter:
#   jupytext:
#     text_representation:
#       extension: .py
#       format_name: light
#       format_version: '1.5'
#       jupytext_version: 1.14.4
#   kernelspec:
#     display_name: Python 3 (ipykernel)
#     language: python
#     name: python3
# ---

# # Timing OpenCL kernels
#
# Measuring the performance of OpenCL kernels is important for finding out how well they peform. The OpenCL framework provides a way to get profiling information from kernels.

# We re-use the matrix multiplication code for this example. One optimisation that can be applied is to transpose the matrix before using it in the matrix computation as shown in the following figure: 

# <figure style="float:center">
#     <img style="display:inline-block; vertical-align:top; margin:20px" src="images/transposed_matrix.png" width="95%">
#     <figcaption style= "text-align:lower; margin:2px; float:bottom; vertical-align:bottom">Figure: Matrix transpose as a potential optimisation </figcaption>
# </figure>
#
# The theory behind this is that since memory elements in dimension 0 (the row dimension) are contiguous, memory accesses to and from matrix $A$ will be more efficient if $A$ was transposed ($A^{T}$) prior to computing the matrix multiplication. Then an element of matrix C at position $(m,n)$ is computed as the dot product of column $m$ from matrix $A^T$ and column $n$ from matrix $B$.

# The code [mat_mult_transpose](code/mat_mult_transpose.cpp) has been modified to include two additional kernels that transpose matrix A and use it in the matrix multiplication. It uses the same data as the simple matrix multiplication code, however we regenerate the data here for convenience.

# + jupyter={"outputs_hidden": true}
# Code to make the test file
import numpy as np
import os

nrows=1024
ncols=1024

# Make up some arrays of random numbers
matrix_A=np.random.random((nrows, ncols)).astype(np.float32)
matrix_B=np.random.random((nrows, ncols)).astype(np.float32)

# Make up the answer, the resulting Matrix C
matrix_C=np.matmul(matrix_A, matrix_B).astype(np.float32)

# Write the files to grid, notice how I unravel in column major (f) format
# before writing to file
matrix_A.ravel(order="f").tofile(os.path.join("code","array_A_1D.dat"))
matrix_B.ravel(order="f").tofile(os.path.join("code","array_B_1D.dat"))
matrix_C.ravel(order="f").tofile(os.path.join("code","array_C_answer_1D.dat"))
# -

# After compilation you can just run the code and see how much faster the transposed version is compared to the standard matrix multiply.

# !cd code; ./mat_mult_transpose

# Try editing line 77 to switch from CPU devices to GPU devices if you have one available and recompile.
#
# from
#
# ```C
# cl_device_type target_device=CL_DEVICE_TYPE_CPU;
# ```
#
# to
#
# ```C
# cl_device_type target_device=CL_DEVICE_TYPE_GPU;
# ```

# !cd code; ./mat_mult_transpose

# For my machine and a GPU it turns out that the transposed matrix calculation is actually around 5x *slower* than the standard implementation. If I use the CPU then the transposed matrix approach is around 2.6x *faster*. I am not certain why the code on the GPU is slower, however being able to find this information out is very helpful. 

# ## Machinery to time kernels

# In order to time kernels, a few key changes must be made to the standard matrix multiplication code. Firstly, the command queue has to be enabled to measure profiling events. This is achieved by setting the flag **CL_QUEUE_PROFILING_ENABLE** in the call to the [clCreateCommandQueue](https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/clCreateCommandQueue.html) function when the command queues are created. 
#
# ```C
# command_queues_1d[queue_counter]=clCreateCommandQueue(  contexts_1d[platform_counter],
#                                                         devices_2d[platform_counter][j], 
#                                                         CL_QUEUE_PROFILING_ENABLE, 
#                                                         &errcode);
# ```
#

# OpenCL events can store timing information from commands that are enqueued. In order to time the kernels we need to associate an OpenCL **Event** with each kernel execution. Watch how I use the event **event_mat_mult** in the execution of the standard matrix multiply kernel.

# ```C
# cl_event event_mat_mult;
#
# // Now enqueue the standard matrix multiply kernel
# errchk(clEnqueueNDRangeKernel(  command_queue,
#                                 kernel_mat_mult,
#                                 work_dim,
#                                 NULL,
#                                 global_size_mat_mult,
#                                 NULL,
#                                 1,
#                                 &event_mat_transpose,
#                                 &event_mat_mult), "Running the kernel");
#
# ```

# This kernel enqueue waits for **event_mat_transpose** from the matrix transpose kernel and associates **event_mat_mult**  with the running of that kernel. 

# At the end of running the kernels it is important to make sure they are finished so we can get proper timing information from the events. I add a **clFinish** to make sure the code finishes all work with the command queue.

# ```C
# clFinish(command_queue);
# ```

# Now to extract the timing information from the completed events we use [clGetEventProfilingInfo](https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/clGetEventProfilingInfo.html) to extract the integer start and stop times (in nanoseconds) of the enqueued kernels. The code below extracts the start and stop times and converts the elapsed time to milliseconds. 

# ```C
# // Get the timing information from each event
# cl_ulong start_counter=0, end_counter=0;
#
# // next the standard matrix multiply
# clGetEventProfilingInfo(    event_mat_mult,
#                             CL_PROFILING_COMMAND_START,
#                             sizeof(cl_ulong),
#                             &start_counter,
#                             NULL);
# clGetEventProfilingInfo(    event_mat_mult,
#                             CL_PROFILING_COMMAND_END,
#                             sizeof(cl_ulong),
#                             &end_counter,
#                             NULL);
#
# // This should give the time in milliseconds
# cl_double time_mat_mult=(cl_double)(end_counter-start_counter)*(cl_double)1.0e-6;
#
# printf("Standard matrix multiply took %f ms\n", time_mat_mult);
# ```

# <address>
# &copy; 2018 by Dr. <NAME><br>
# email: <a href="mailto:<EMAIL>"><EMAIL></a><br>
# Visit us at: <a href="https://www.pelagos-consulting.com">www.pelagos-consulting.com</a><br>
# </address>

StarcoderdataJupyterScriptsDedupFiltered