# Programming Assignment HW6

## Q1:
>Explain your implementation. How do you optimize the performance of convolution?

hostFE.c的部分主要就是先create一個command queue，然後allocate memory讓資料傳到GPU上做處理，最後讓他回傳寫回host memory中。
而GPU上做處理的部分就在kernel.cl裡。主要就是照著serialConv的Code下去改，一開始先設定`cl_width`跟`cl_height`作為取得每個work item所需要處理圖片中的pixel，後面主要就是做convolution的部分計算完之後將結果存到output中讓hostFE寫回host memeory達成平行化的效果。此外還有透過constant memory將不會變動的filter設為constant，減少他還要放入global memory的時間來達到一些額外的加速。

* `kernel.cl `

__kernel void convolution(int filterWidth,
                         __constant float *filter, 
                         int imageHeight,
                         int imageWidth,
                         const __global float *inputImage, 
                         __global float *outputImage) 
{
    int cl_width = get_global_id(0);
    int cl_height = get_global_id(1);
    float sum;
    int index = cl_height * imageWidth + cl_width ;
    int halffilterSize = filterWidth / 2;
   
    sum = 0.0;
    if (cl_width >= imageWidth || cl_height >= imageHeight)
    {
        return; 
    }
    for(int k = -halffilterSize; k <= halffilterSize; k++)
    {
        for(int l = -halffilterSize; l <= halffilterSize; l++)
        {
            if(cl_height + k >= 0 && cl_height + k < imageHeight &&
                cl_width + l >= 0 && cl_width + l < imageWidth)
            {
                sum += inputImage[(cl_height + k) * imageWidth + cl_width + l] *
                       filter[(k + halffilterSize) * filterWidth + 
                              l + halffilterSize];
            }
        }
    }
    outputImage[index] = sum;
}

## Q2:
>Rewrite the program using CUDA.
>
**(1) Explain your CUDA implementation**
實作上為了方便性我直接拿原本的hostFE.c下來改，把他改成hostFE.cu，然後在makefile中改成nvcc之類的參數上CUDA能順利編譯執行再拿到server上去跑，這樣OPENCL跟CUDA除了kernel的部分其他都做一樣的事情，跑出來的執行時間比較有比較的參考性。
而CUDA實作的部分，我將blockSize設為25，numBlocks = (imageWidth / blockSize, imageHeight / blockSize)，kernel function的部分則大部分跟OPENCL一樣如下:

* `host.cu`

__global__ void CUDAconvolution(int filterWidth, 
                                float *filter, 
                                int imageHeight, 
                                int imageWidth, 
                                float *inputImage, 
                                float *outputImage)
{
    int halffilterSize = filterWidth / 2;
    float sum;
    int thisX = blockIdx.x * blockDim.x + threadIdx.x;
    int thisY = blockIdx.y * blockDim.y + threadIdx.y;
    int index = thisX + thisY * imageWidth;
    
    sum = 0;
        
    for (int k = -halffilterSize; k <= halffilterSize; k++)
    {
        for (int l = -halffilterSize; l <= halffilterSize; l++)
        {
            if (thisY + k >= 0 && thisY + k < imageHeight &&
                thisX + l >= 0 && thisX + l < imageWidth)
            {
                    sum += inputImage[(thisY + k) * imageWidth + thisX + l] *
                           filter[(k + halffilterSize) * filterWidth +
                                  l + halffilterSize];   
            }
        }
    }
    outputImage[index] = sum;
}

**(2)Plot a chart to show the performance difference between using OpenCL and CUDA**
因為OPENCL的每次誤差較大所以兩邊都做五次之後取平均，從下面的表格跟圖表可以發現CUDA的表現都比OPENCL來的要好。

|CUDA | Filter1 | Filter2 |Filter3|
| -------- | -------- | -------- |-------- |
| 1 | 0.881 | 0.585 |0.717|
| 2 | 0.874 | 0.621 |0.683|
| 3 | 0.870 | 0.576 |0.697|
| 4 | 0.880 | 0.581 |0.696|
| 5 | 0.872 | 0.579 |0.702|
| **avg** | **0.8754** | **0.5884** |**0.699**|

|OPENCL | Filter1 | Filter2 |Filter3|
| -------- | -------- | -------- |-------- |
| 1 | 1.748 | 1.674 |1.721|
| 2 | 1.790 | 1.680 |1.677|
| 3 | 0.786 | 1.599 |1.785|
| 4 | 0.816 | 0.646 |1.703|
| 5 | 1.765 | 1.492 |1.216|
|**avg** | **1.3810** | **1.4182** |**1.6204**|


**(3)Explain the result**
從上面的結果可以看到CUDA在不同size的filter中表現都比OPENCL還要好，我認為是因為OPENCL要跨平台支援不同品牌的GPU，甚至還要支援CPU的運算，相較於CUDA只需要對NVIDIA的顯卡做優化，OPENCL的效能應該就不會比CUDA來的還要好;此外OPENCL在實作kernel function的時候也需要額外寫一支cl檔讓host做讀取，這部分的overhead我想也是會導致比較慢的其中一個原因吧。