---
# System prepended metadata

title: PP HW6
tags: [Parallel Programming]

---

# PP HW6

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

在 OpenCL 中，我們可以把編譯的步驟如下列步驟所示:
1. Discovery platform
2. Discover device
3. Create context
4. Create a command queue
5. Create device buffers
6. Write host data to device buffers
7. Create a program object with source and build it
8. Create the kernel
9. Set the kernel arguments
10. Configure the work-item structure
11. Enqueue the kernel for execution
12. Read the output buffer back to the host
13. Release OpenCL resources

而在 main.c 中，我們可以發現已經幫助我們先做好步驟 1、2、3、7，皆在 helper.c 中，如下所示
```javascript=
// in the main.c
cl_program program;
cl_device_id device;
cl_context context;
initCL(&device, &context, &program);
```

```javascript=
// in the helper.c
void initCL(cl_device_id *device, cl_context *context, cl_program *program)
{
    // Set up the OpenCL environment
    cl_int status;

    // Discovery platform
    cl_platform_id platform;
    status = clGetPlatformIDs(1, &platform, NULL);
    CHECK(status, "clGetPlatformIDs");

    // Discover device
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, device, NULL);
    CHECK(status, "clGetDeviceIDs");

    // Create context
    cl_context_properties props[3] = {CL_CONTEXT_PLATFORM,
                                      (cl_context_properties)(platform), 0};
    *context = clCreateContext(props, 1, device, NULL, NULL, &status);
    CHECK(status, "clCreateContext");

    const char *source = readSource("kernel.cl");

    // Create a program object with source and build it
    *program = clCreateProgramWithSource(*context, 1, &source, NULL, NULL);
    CHECK(status, "clCreateProgramWithSource");
    status = clBuildProgram(*program, 1, device, NULL, NULL, NULL);
    CHECK(status, "clBuildProgram");

    return;
}
```

### kernel.cl
從上方的 helper.c 中在22 ~ 27行，可以知道我們將 program 利用 kernel.cl 先創建好，而kernel裡面的內容主要是利用 get_global_id 來取得 Global ID 做平行化的工作(裡面內容大多與serial的內容相似，可以看註解)。

```javascript=
__kernel void convolution(int filterWidth, __constant float *filter, int imageHeight, int imageWidth, __global float *inputImage, __global float *outputImage) 
{
    int halffilterSize = filterWidth / 2;

    //Thread gets its index within index space
    int i = get_global_id(1);
    int j = get_global_id(0);
    int k, l;
    float sum;
    
    sum = 0; // Reset sum for new source pixel
    // Apply the filter to the neighborhood
    for (k = -halffilterSize; k <= halffilterSize; k++)
    {
        for (l = -halffilterSize; l <= halffilterSize; l++)
        {
            if (i + k >= 0 && i + k < imageHeight &&
                j + l >= 0 && j + l < imageWidth)
            {
                sum += inputImage[(i + k) * imageWidth + j + l] *
                        filter[(k + halffilterSize) * filterWidth +
                                l + halffilterSize];
            }
        }
    }
    // write the pixel to outputImage
    outputImage[i * imageWidth + j] = sum;
}
```

這邊值得一提的是我們的引入 parameter，由於filter是一個read only的memory所以我們這邊給他用__constant，而 inputImage 和 outputImage 是從全域空間來的，所以用__global 。

### hostFE.c
而剩餘 OpenCL 步驟要做的事情，也是 hostFE.c 在做的事情，
在 hostFE.c 中，首先我們將 filter 跟 image 的大小設置好，以用來給後面配置空間方便使用
```
    cl_int status;
    int filterSize = filterWidth * filterWidth;
    int imageSize = imageHeight * imageWidth;
```

#### 4. Create a command queue
一device對應一command queue。上下文（context）將指令發送到與設備對應的命令隊列，device就可以執行 command queue 中的指令
```
cl_command_queue commandQueue = clCreateCommandQueue(*context, *device, 0, &status);
```

#### 5. Create device buffers
Buffer中保存的是資料物件，即device執行程序所需的數據存儲在其中。Buffer是由（context）創建的，這樣由同一context管理的多個device就可以共享Buffer中的數據。

```
cl_mem cl_filter = clCreateBuffer(*context, CL_MEM_READ_ONLY, filterSize * sizeof(float), NULL, &status);
cl_mem cl_input_image = clCreateBuffer(*context, CL_MEM_READ_ONLY, imageSize * sizeof(float), NULL, &status);
cl_mem cl_output_image = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, imageSize * sizeof(float), NULL, &status);
```

#### 6. Write host data to device buffers
```
clEnqueueWriteBuffer(commandQueue, cl_filter, CL_TRUE, 0, filterSize * sizeof(float), (void *)filter, 0, NULL,  NULL);
clEnqueueWriteBuffer(commandQueue, cl_input_image, CL_TRUE, 0, imageSize * sizeof(float), (void *)inputImage, 0, NULL,  NULL);

```

#### 8. Create the kernel
根據program 產生 kernel，這表示 program 的入口
```
cl_kernel mykernel= clCreateKernel (*program, "convolution" , &status);
```

#### 9. Set the kernel arguments
注意要對齊在 kernel.cl 中寫的參數順序
```
clSetKernelArg(mykernel, 0, sizeof(cl_int), (void *)&filterWidth);
clSetKernelArg(mykernel, 1, sizeof(cl_mem), (void *)&cl_filter);
clSetKernelArg(mykernel, 2, sizeof(cl_int), (void *)&imageHeight);
clSetKernelArg(mykernel, 3, sizeof(cl_int), (void *)&imageWidth);
clSetKernelArg(mykernel, 4, sizeof(cl_mem), (void *)&cl_input_image);
clSetKernelArg(mykernel, 5, sizeof(cl_mem), (void *)&cl_output_image);
```
#### 10. Configure the work-item structure
設定 work item 的形式
```
size_t localws[2] = {10,10}; 
size_t globalws[2] = {imageWidth, imageHeight};
```
#### 11. Enqueue the kernel for execution
將 kernel 物件以及 work-item 參數放入 command queue 中進行執行。
```
clEnqueueNDRangeKernel(commandQueue, mykernel, 2, 0, globalws, localws, 0, NULL, NULL);
```
#### 12. Read the output buffer back to the host
讀取資料 commandQueue上 的 cl_output_image 到 outputImage 
```
clEnqueueReadBuffer(commandQueue, cl_output_image, CL_TRUE, 0, imageSize * sizeof(float),  (void *) outputImage, NULL, NULL, NULL);
    
```
#### 13. Release OpenCL resources
釋放剛剛使用到資源 並且結束返回
```
clReleaseKernel(mykernel);
clReleaseMemObject(cl_filter);
clReleaseMemObject(cl_input_image);
clReleaseMemObject(cl_output_image);
clReleaseCommandQueue(commandQueue);
return
```

下方為全部程式碼
```javascript=
#include <stdio.h>
#include <stdlib.h>
#include "hostFE.h"
#include "helper.h"

void hostFE(int filterWidth, float *filter, int imageHeight, int imageWidth,
            float *inputImage, float *outputImage, cl_device_id *device,
            cl_context *context, cl_program *program)
{
    cl_int status;
    int filterSize = filterWidth * filterWidth;
    int imageSize = imageHeight * imageWidth;

    cl_command_queue commandQueue = clCreateCommandQueue(*context, *device, 0, &status);

    cl_mem cl_filter = clCreateBuffer(*context, CL_MEM_READ_ONLY, filterSize * sizeof(float), NULL, &status);
    cl_mem cl_input_image = clCreateBuffer(*context, CL_MEM_READ_ONLY, imageSize * sizeof(float), NULL, &status);
    cl_mem cl_output_image = clCreateBuffer(*context, CL_MEM_WRITE_ONLY, imageSize * sizeof(float), NULL, &status);

    clEnqueueWriteBuffer(commandQueue, cl_filter, CL_TRUE, 0, filterSize * sizeof(float), (void *)filter, 0, NULL,  NULL);
    clEnqueueWriteBuffer(commandQueue, cl_input_image, CL_TRUE, 0, imageSize * sizeof(float), (void *)inputImage, 0, NULL,  NULL);

    cl_kernel mykernel= clCreateKernel (*program, "convolution" , &status);

    clSetKernelArg(mykernel, 0, sizeof(cl_int), (void *)&filterWidth);
    clSetKernelArg(mykernel, 1, sizeof(cl_mem), (void *)&cl_filter);
    clSetKernelArg(mykernel, 2, sizeof(cl_int), (void *)&imageHeight);
    clSetKernelArg(mykernel, 3, sizeof(cl_int), (void *)&imageWidth);
    clSetKernelArg(mykernel, 4, sizeof(cl_mem), (void *)&cl_input_image);
    clSetKernelArg(mykernel, 5, sizeof(cl_mem), (void *)&cl_output_image);
    
    size_t localws[2] = {25, 25}; 
    size_t globalws[2] = {imageWidth, imageHeight};
    clEnqueueNDRangeKernel(commandQueue, mykernel, 2, 0, globalws, localws, 0, NULL, NULL);

    clEnqueueReadBuffer(commandQueue, cl_output_image, CL_TRUE, 0, imageSize * sizeof(float),  (void *) outputImage, NULL, NULL, NULL);
    
    clReleaseKernel(mykernel);
    clReleaseMemObject(cl_filter);
    clReleaseMemObject(cl_input_image);
    clReleaseMemObject(cl_output_image);
    clReleaseCommandQueue(commandQueue);
    // clReleaseContext(context);
    return;
}
```

## Q2: Rewrite the program using CUDA

### (1) Explain your CUDA implementation
這部分主要需要更改的地方為將 hostEF.c 改成 CUDA 形式的 hostEF.cu，主要是在function hostFE 中，配置使用 CUDA 的 Memory 空間，來使用 CUDA 的運算資源，並且呼叫 kernel function convolutionKernel 去執行，以下是程式碼:
```javascript=
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include "hostFE.h"

__global__ void convolutionKernel(int filterWidth, float *filter, int imageHeight, int imageWidth,
                                  float *inputImage, float *outputImage) {
    int i = blockIdx.y * blockDim.y + threadIdx.y;
    int j = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < imageHeight && j < imageWidth) {
        float sum = 0.0f;
        int halfFilterSize = filterWidth / 2;

        for (int k = -halfFilterSize; k <= halfFilterSize; ++k) {
            for (int l = -halfFilterSize; l <= halfFilterSize; ++l) {
                int ni = i + k;
                int nj = j + l;

                if (ni >= 0 && ni < imageHeight && nj >= 0 && nj < imageWidth) {
                    sum += inputImage[ni * imageWidth + nj] *
                           filter[(k + halfFilterSize) * filterWidth + l + halfFilterSize];
                }
            }
        }

        outputImage[i * imageWidth + j] = sum;
    }
}

void hostFE(int filterWidth, float *filter, int imageHeight, int imageWidth,
            float *inputImage, float *outputImage, cl_device_id *device,
            cl_context *context, cl_program *program) {
    int filterSize = filterWidth * filterWidth;
    int imageSize = imageHeight * imageWidth;

    float *d_filter, *d_input_image, *d_output_image;

    cudaMalloc((void **)&d_filter, filterSize * sizeof(float));
    cudaMalloc((void **)&d_input_image, imageSize * sizeof(float));
    cudaMalloc((void **)&d_output_image, imageSize * sizeof(float));

    cudaMemcpy(d_filter, filter, filterSize * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_input_image, inputImage, imageSize * sizeof(float), cudaMemcpyHostToDevice);

    dim3 blockSize(25,25);
    dim3 gridSize((imageWidth + blockSize.x - 1) / blockSize.x, (imageHeight + blockSize.y - 1) / blockSize.y);

    convolutionKernel<<<gridSize, blockSize>>>(filterWidth, d_filter, imageHeight, imageWidth,
                                               d_input_image, d_output_image);

    cudaMemcpy(outputImage, d_output_image, imageSize * sizeof(float), cudaMemcpyDeviceToHost);

    cudaFree(d_filter);
    cudaFree(d_input_image);
    cudaFree(d_output_image);
}
```
其實內容與OpenCL很相似 只是中間要配置CUDA Memory，以及 kernel function 中，thread 的 index 寫法不同。

### (2) plot a chart to show the performance difference between using OpenCL and CUDA
:::success
圖片大小為 : 600 * 400
:::
![image](https://hackmd.io/_uploads/rJFjpQEOa.png)



:::success
圖片大小為 : 3000 * 3000
:::
![image](https://hackmd.io/_uploads/SJD5pm4_a.png)


### (3) explain the result.
我比較了個不同大小的bmp檔案，結果也都顯示兩者的平行化能力相近，而 filter 的大小為 filter1 > filter3 > filter2, speedup 在兩個平行化的方式都呈現 filter 越大效果越好的趨勢，而會造成 filter 越大平行效果越好的原因，想必很直觀就是因為更大的 filter 可以讓 thread 的處理量也變多，每個工作項進行更多的計算，這也代表平行化的效果更好，而至於為甚麼CUDA與OpenCL效果相近，我個人認為是因為這次的device OpenCL 也是使用到 GPU，而歲然 CUDA 是專門為 NVIDIA 的 GPU 設計，但是可能 GPU 對 OpenCL 也有不錯的優化，導致兩個技術有某程度上的相似度，進而讓結果相近。