SHARE
TWEET

Untitled

a guest Sep 22nd, 2019 105 Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. // MatrixMultiplication.cpp: 콘솔 응용 프로그램의 진입점을 정의합니다.
  2. //
  3. // CPU              - Intel(R) Core(TM) i7-6700 CPU @ 3.40GHz
  4. // Memory           - 16.0 GB
  5. //
  6. // Platform Number  - 1
  7. // Platform Name    - AMD Accelerated Parallel Processing
  8. // Platform vender  - Advanced Micro Devices, Inc.
  9. //
  10. // Device Type      - CL_DEVICE_TYPE_GPU
  11. // Device Name      - Ellesmere (Radeon RX480)
  12. //      CL_DEVICE_MAX_WORK_GROUP_SIZE   - 256
  13. //      CL_DEVICE_GLOBAL_MEM_SIZE       - 3221225472
  14. //      CL_DEVICE_LOCAL_MEM_SIZE        - 32768
  15. //      CL_DEVICE_MAX_MEM_ALLOC_SIZE    - 3221225472
  16. //
  17.  
  18.  
  19. #include <CL/cl.h>
  20. #include <stdio.h>
  21. #include <stdlib.h>
  22. #include <time.h>
  23. #include <string.h>
  24.  
  25. #define N 1024
  26.  
  27. int ROW_A = N;
  28. int COL_A = N;
  29. int ROW_B = COL_A;
  30. int COL_B = N;
  31.  
  32. clock_t start, end;
  33. char *get_source_code(const char *file_name, size_t *len);
  34.  
  35. int mat_equal(int* A, int* B);
  36.  
  37. int* mat_mul_seq(int* A, int* B);
  38. int* mat_mul_seq_2(int* A, int* B);
  39. int* mat_mul_opencl(int* A, int* B);
  40. int* mat_mul_opencl_2(int* A, int* B);
  41. int* mat_mul_opencl_3(int* A, int* B);
  42. int* mat_mul_opencl_4(int* A, int* B);
  43. int* mat_mul_opencl_5(int* A, int* B);
  44. int* mat_mul_opencl_6(int* A, int* B);
  45. int* mat_mul_strassen(int* A, int* B);
  46.  
  47. int *strassen(int *C, int *A, int *B, int n);
  48. void strassen_result(int *C, int *c11, int *c12, int *c21, int *c22, int n);
  49. int *strassen_add(int *res, int *x, int *y, int n);
  50. int *strassen_sub(int *res, int *x, int *y, int n);
  51. int *strassen_mul(int *res, int *x, int *y, int n);
  52.  
  53. int main()
  54. {
  55.     int* A = (int*)malloc(sizeof(int) * ROW_A * COL_A);
  56.     int* B = (int*)malloc(sizeof(int) * ROW_B * COL_B);
  57.     int *C, *D, *E, *F, *G, *H, *I, *J, *K;
  58.     int i, j;
  59.  
  60.     printf("Make matrix A...\n");
  61.     for (i = 0; i < ROW_A * COL_A; i++) A[i] = rand() % 10;
  62.     printf("Make matrix B...\n\n");
  63.     for (i = 0; i < ROW_B * COL_B; i++) B[i] = rand() % 10;
  64.  
  65.     printf("Sequential version 1...\n");
  66.     C = mat_mul_seq(A, B);
  67.     printf("\n");
  68.  
  69.     printf("Sequential version 2...\n");
  70.     D = mat_mul_seq_2(A, B);
  71.     mat_equal(C, D);
  72.     printf("\n");
  73.  
  74.     printf("OpenCL version 1...\n");
  75.     E = mat_mul_opencl(A, B);
  76.     mat_equal(C, E);
  77.     printf("\n");
  78.  
  79.     printf("OpenCL version 2...\n");
  80.     F = mat_mul_opencl_2(A, B);
  81.     mat_equal(C, F);
  82.     printf("\n");
  83.  
  84.     printf("OpenCL version 3...\n");
  85.     G = mat_mul_opencl_3(A, B);
  86.     mat_equal(C, G);
  87.     printf("\n");
  88.  
  89.     printf("OpenCL version 4...\n");
  90.     H = mat_mul_opencl_4(A, B);
  91.     mat_equal(C, H);
  92.     printf("\n");
  93.  
  94.     printf("OpenCL version 5...\n");
  95.     I = mat_mul_opencl_5(A, B);
  96.     mat_equal(C, I);
  97.     printf("\n");
  98.  
  99.     printf("OpenCL version 6...\n");
  100.     J = mat_mul_opencl_6(A, B);
  101.     mat_equal(C, J);
  102.     printf("\n");
  103.  
  104.     printf("Strassen version...\n");
  105.     K = mat_mul_strassen(A, B);
  106.     mat_equal(C, K);
  107.     printf("\n");
  108.  
  109.     free(A);
  110.     free(B);
  111.     free(C);
  112.     free(D);
  113.     free(E);
  114.     free(F);
  115.     free(G);
  116.     free(H);
  117.     free(I);
  118.     free(J);
  119.     free(K);
  120.  
  121.     return 0;
  122. }
  123.  
  124. int mat_equal(int* A, int* B) {
  125.     int i;
  126.  
  127.     for (i = 0; i < ROW_A * COL_B; i++)
  128.         if (A[i] != B[i]) {
  129.             printf("Calculation Fail!\n");
  130.             return 0;
  131.         }
  132.  
  133.     printf("Calculation Succeed!\n");
  134.  
  135.     return 1;
  136. }
  137.  
  138. int* mat_mul_seq(int* A, int* B) {
  139.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  140.     int i, j, k;
  141.  
  142.     start = clock();
  143.  
  144.     for (i = 0; i < ROW_A; i++)
  145.         for (j = 0; j < COL_B; j++) {
  146.             C[i * COL_B + j] = 0;
  147.             for (k = 0; k < COL_A; k++)
  148.                 C[i * COL_B + j] += A[i * COL_A + k] * B[k * COL_B + j];
  149.         }
  150.  
  151.     end = clock();
  152.  
  153.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  154.  
  155.     return C;
  156. }
  157.  
  158. int* mat_mul_seq_2(int* A, int* B) {
  159.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  160.     int i, j, k, tmp;
  161.  
  162.     for (i = 0; i < ROW_A * COL_B; i++) C[i] = 0;
  163.  
  164.     start = clock();
  165.  
  166.     for (i = 0; i < ROW_A; i++)
  167.         for (j = 0; j < COL_B; j++) {
  168.             tmp = 0;
  169.             for (k = 0; k < COL_A; k++)
  170.                 tmp += A[i * COL_A + k] * B[k * COL_B + j];
  171.             C[i * COL_B + j] = tmp;
  172.         }
  173.  
  174.     end = clock();
  175.  
  176.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  177.  
  178.     return C;
  179. }
  180.  
  181. #define CHECK_ERROR(err) \
  182.     if(err != CL_SUCCESS) { \
  183.         printf("[%s:%d] OpenCL error %d\n", __FILE__, __LINE__, err); \
  184.         exit(EXIT_FAILURE); \
  185.     }
  186.  
  187. char *get_source_code(const char *file_name, size_t *len) {
  188.     char *source_code;
  189.     char buf[2] = "\0";
  190.     int cnt = 0;
  191.     size_t length;
  192.  
  193.     FILE *file = fopen(file_name, "r");
  194.  
  195.     if (file == NULL) {
  196.         printf("[%s:%d] Failed to open %s ", __FILE__, __LINE__, file_name);
  197.         exit(EXIT_FAILURE);
  198.     }
  199.  
  200.     fseek(file, 0, SEEK_END);
  201.     length = (size_t)ftell(file);
  202.     rewind(file);
  203.  
  204.     source_code = (char*)malloc(length + 1);
  205.     fread(source_code, length, 1, file);
  206.  
  207.     for (int i = 0; i < length; i++) {
  208.         buf[0] = source_code[i];
  209.  
  210.         if (buf[0] == '\n')
  211.             cnt++;
  212.     }
  213.  
  214.     source_code[length - cnt] = '\0';
  215.  
  216.     fclose(file);
  217.  
  218.     *len = length - cnt;
  219.  
  220.     return source_code;
  221. }
  222.  
  223. int* mat_mul_opencl(int* A, int* B) {
  224.     cl_uint num_platforms;
  225.     cl_platform_id *platforms;
  226.     cl_uint num_devices;
  227.     char str[1024];
  228.     cl_device_type device_type;
  229.     cl_device_id device;
  230.     cl_context context;
  231.     cl_command_queue queue;
  232.     cl_program program;
  233.     char *kernel_source;
  234.     size_t kernel_source_size;
  235.     cl_kernel kernel_mat_mul;
  236.     cl_mem bufA, bufB, bufC;
  237.     cl_int err;
  238.  
  239.     size_t global_size[2] = { ROW_A, COL_B };
  240.     size_t local_size[2] = { 1, 1 };
  241.  
  242.     int i;
  243.  
  244.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  245.  
  246.     err = clGetPlatformIDs(0, NULL, &num_platforms);
  247.     CHECK_ERROR(err);
  248.  
  249.     platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  250.     err = clGetPlatformIDs(num_platforms, platforms, NULL);
  251.     CHECK_ERROR(err);
  252.  
  253.     err = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  254.     CHECK_ERROR(err);
  255.  
  256.     // Create Context
  257.     context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  258.     CHECK_ERROR(err);
  259.  
  260.     // Create Command Queue
  261.     queue = clCreateCommandQueue(context, device, 0, &err);
  262.     CHECK_ERROR(err);
  263.  
  264.     // Create Program Object
  265.     kernel_source = get_source_code("kernel.cl", &kernel_source_size);
  266.     program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
  267.     CHECK_ERROR(err);
  268.  
  269.     // Build Program
  270.     err = clBuildProgram(program, 1, &device, "", NULL, NULL);
  271.     CHECK_ERROR(err);
  272.  
  273.     // Create Kernel
  274.     kernel_mat_mul = clCreateKernel(program, "mat_mul", &err);
  275.     CHECK_ERROR(err);
  276.  
  277.     // Write Buffer
  278.     bufA = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_A, A, &err);
  279.     CHECK_ERROR(err);
  280.     bufB = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_B * COL_B, B, &err);
  281.     CHECK_ERROR(err);
  282.     bufC = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_B, C, &err);
  283.     CHECK_ERROR(err);
  284.  
  285.     // Set Kernel Arg
  286.     err = clSetKernelArg(kernel_mat_mul, 0, sizeof(cl_mem), &bufA);
  287.     CHECK_ERROR(err);
  288.     err = clSetKernelArg(kernel_mat_mul, 1, sizeof(cl_mem), &bufB);
  289.     CHECK_ERROR(err);
  290.     err = clSetKernelArg(kernel_mat_mul, 2, sizeof(cl_mem), &bufC);
  291.     CHECK_ERROR(err);
  292.     err = clSetKernelArg(kernel_mat_mul, 3, sizeof(int), &ROW_A);
  293.     CHECK_ERROR(err);
  294.     err = clSetKernelArg(kernel_mat_mul, 4, sizeof(int), &COL_B);
  295.     CHECK_ERROR(err);
  296.     err = clSetKernelArg(kernel_mat_mul, 5, sizeof(int), &COL_A);
  297.     CHECK_ERROR(err);
  298.  
  299.     start = clock();
  300.  
  301.     // Execute Kernel
  302.     clEnqueueNDRangeKernel(queue, kernel_mat_mul, 2, NULL, global_size, local_size, 0, NULL, NULL);
  303.  
  304.     err = clFinish(queue);
  305.     CHECK_ERROR(err);
  306.  
  307.     end = clock();
  308.  
  309.     // Read Buffer
  310.     err = clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(int) * ROW_A * COL_B, C, 0, NULL, NULL);
  311.     CHECK_ERROR(err);
  312.  
  313.     err = clReleaseKernel(kernel_mat_mul);
  314.     CHECK_ERROR(err);
  315.     err = clReleaseProgram(program);
  316.     CHECK_ERROR(err);
  317.     err = clReleaseMemObject(bufA);
  318.     CHECK_ERROR(err);
  319.     err = clReleaseMemObject(bufB);
  320.     CHECK_ERROR(err);
  321.     err = clReleaseMemObject(bufC);
  322.     CHECK_ERROR(err);
  323.     err = clReleaseCommandQueue(queue);
  324.     CHECK_ERROR(err);
  325.     err = clReleaseContext(context);
  326.     CHECK_ERROR(err);
  327.  
  328.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  329.     return C;
  330. }
  331.  
  332. int* mat_mul_opencl_2(int* A, int* B) {
  333.     cl_uint num_platforms;
  334.     cl_platform_id *platforms;
  335.     cl_uint num_devices;
  336.     char str[1024];
  337.     cl_device_type device_type;
  338.     cl_device_id device;
  339.     cl_context context;
  340.     cl_command_queue queue;
  341.     cl_program program;
  342.     char *kernel_source;
  343.     size_t kernel_source_size;
  344.     cl_kernel kernel_mat_mul;
  345.     cl_mem bufA, bufB, bufC;
  346.     cl_int err;
  347.  
  348.     size_t global_size[2] = { ROW_A, COL_B };
  349.     size_t local_size[2] = { 1, 1 };
  350.  
  351.     int i;
  352.  
  353.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  354.     for (i = 0; i < ROW_A * COL_B; i++) C[i] = 0;
  355.  
  356.     err = clGetPlatformIDs(0, NULL, &num_platforms);
  357.     CHECK_ERROR(err);
  358.  
  359.     platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  360.     err = clGetPlatformIDs(num_platforms, platforms, NULL);
  361.     CHECK_ERROR(err);
  362.  
  363.     err = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  364.     CHECK_ERROR(err);
  365.  
  366.     // Create Context
  367.     context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  368.     CHECK_ERROR(err);
  369.  
  370.     // Create Command Queue
  371.     queue = clCreateCommandQueue(context, device, 0, &err);
  372.     CHECK_ERROR(err);
  373.  
  374.     // Create Program Object
  375.     kernel_source = get_source_code("kernel2.cl", &kernel_source_size);
  376.     program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
  377.     CHECK_ERROR(err);
  378.  
  379.     // Build Program
  380.     err = clBuildProgram(program, 1, &device, "", NULL, NULL);
  381.     CHECK_ERROR(err);
  382.  
  383.     // Create Kernel
  384.     kernel_mat_mul = clCreateKernel(program, "mat_mul", &err);
  385.     CHECK_ERROR(err);
  386.  
  387.     // Write Buffer
  388.     bufA = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_A, A, &err);
  389.     CHECK_ERROR(err);
  390.     bufB = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_B * COL_B, B, &err);
  391.     CHECK_ERROR(err);
  392.     bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_B, C, &err);
  393.     CHECK_ERROR(err);
  394.  
  395.     // Set Kernel Arg
  396.     err = clSetKernelArg(kernel_mat_mul, 0, sizeof(cl_mem), &bufA);
  397.     CHECK_ERROR(err);
  398.     err = clSetKernelArg(kernel_mat_mul, 1, sizeof(cl_mem), &bufB);
  399.     CHECK_ERROR(err);
  400.     err = clSetKernelArg(kernel_mat_mul, 2, sizeof(cl_mem), &bufC);
  401.     CHECK_ERROR(err);
  402.     err = clSetKernelArg(kernel_mat_mul, 3, sizeof(int), &ROW_A);
  403.     CHECK_ERROR(err);
  404.     err = clSetKernelArg(kernel_mat_mul, 4, sizeof(int), &COL_B);
  405.     CHECK_ERROR(err);
  406.     err = clSetKernelArg(kernel_mat_mul, 5, sizeof(int), &COL_A);
  407.     CHECK_ERROR(err);
  408.  
  409.     start = clock();
  410.  
  411.     // Execute Kernel
  412.     clEnqueueNDRangeKernel(queue, kernel_mat_mul, 2, NULL, global_size, local_size, 0, NULL, NULL);
  413.  
  414.     err = clFinish(queue);
  415.     CHECK_ERROR(err);
  416.  
  417.     end = clock();
  418.  
  419.     // Read Buffer
  420.     err = clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(int) * ROW_A * COL_B, C, 0, NULL, NULL);
  421.     CHECK_ERROR(err);
  422.  
  423.     err = clReleaseKernel(kernel_mat_mul);
  424.     CHECK_ERROR(err);
  425.     err = clReleaseProgram(program);
  426.     CHECK_ERROR(err);
  427.     err = clReleaseMemObject(bufA);
  428.     CHECK_ERROR(err);
  429.     err = clReleaseMemObject(bufB);
  430.     CHECK_ERROR(err);
  431.     err = clReleaseMemObject(bufC);
  432.     CHECK_ERROR(err);
  433.     err = clReleaseCommandQueue(queue);
  434.     CHECK_ERROR(err);
  435.     err = clReleaseContext(context);
  436.     CHECK_ERROR(err);
  437.  
  438.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  439.     return C;
  440. }
  441.  
  442. int* mat_mul_opencl_3(int* A, int* B) {
  443.     cl_uint num_platforms;
  444.     cl_platform_id *platforms;
  445.     cl_uint num_devices;
  446.     char str[1024];
  447.     cl_device_type device_type;
  448.     cl_device_id device;
  449.     cl_context context;
  450.     cl_command_queue queue;
  451.     cl_program program;
  452.     char *kernel_source;
  453.     size_t kernel_source_size;
  454.     cl_kernel kernel_mat_mul;
  455.     cl_mem bufA, bufB, bufC;
  456.     cl_int err;
  457.  
  458.     size_t global_size[1] = { ROW_A };
  459.     size_t local_size[1] = { 256 };
  460.     int i;
  461.  
  462.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  463.     for (i = 0; i < ROW_A * COL_B; i++) C[i] = 0;
  464.  
  465.     err = clGetPlatformIDs(0, NULL, &num_platforms);
  466.     CHECK_ERROR(err);
  467.  
  468.     platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  469.     err = clGetPlatformIDs(num_platforms, platforms, NULL);
  470.     CHECK_ERROR(err);
  471.  
  472.     err = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  473.     CHECK_ERROR(err);
  474.  
  475.     // Create Context
  476.     context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  477.     CHECK_ERROR(err);
  478.  
  479.     // Create Command Queue
  480.     queue = clCreateCommandQueue(context, device, 0, &err);
  481.     CHECK_ERROR(err);
  482.  
  483.     // Create Program Object
  484.     kernel_source = get_source_code("kernel3.cl", &kernel_source_size);
  485.     program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
  486.     CHECK_ERROR(err);
  487.  
  488.     // Build Program
  489.     err = clBuildProgram(program, 1, &device, "", NULL, NULL);
  490.     CHECK_ERROR(err);
  491.  
  492.     // Create Kernel
  493.     kernel_mat_mul = clCreateKernel(program, "mat_mul", &err);
  494.     CHECK_ERROR(err);
  495.  
  496.     // Write Buffer
  497.     bufA = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_A, A, &err);
  498.     CHECK_ERROR(err);
  499.     bufB = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_B * COL_B, B, &err);
  500.     CHECK_ERROR(err);
  501.     bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_B, C, &err);
  502.     CHECK_ERROR(err);
  503.  
  504.     // Set Kernel Arg
  505.     err = clSetKernelArg(kernel_mat_mul, 0, sizeof(cl_mem), &bufA);
  506.     CHECK_ERROR(err);
  507.     err = clSetKernelArg(kernel_mat_mul, 1, sizeof(cl_mem), &bufB);
  508.     CHECK_ERROR(err);
  509.     err = clSetKernelArg(kernel_mat_mul, 2, sizeof(cl_mem), &bufC);
  510.     CHECK_ERROR(err);
  511.     err = clSetKernelArg(kernel_mat_mul, 3, sizeof(int), &ROW_A);
  512.     CHECK_ERROR(err);
  513.     err = clSetKernelArg(kernel_mat_mul, 4, sizeof(int), &COL_B);
  514.     CHECK_ERROR(err);
  515.     err = clSetKernelArg(kernel_mat_mul, 5, sizeof(int), &COL_A);
  516.     CHECK_ERROR(err);
  517.  
  518.     start = clock();
  519.  
  520.     // Execute Kernel
  521.     clEnqueueNDRangeKernel(queue, kernel_mat_mul, 1, NULL, global_size, local_size, 0, NULL, NULL);
  522.  
  523.     err = clFinish(queue);
  524.     CHECK_ERROR(err);
  525.  
  526.     end = clock();
  527.  
  528.     // Read Buffer
  529.     err = clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(int) * ROW_A * COL_B, C, 0, NULL, NULL);
  530.     CHECK_ERROR(err);
  531.  
  532.     err = clReleaseKernel(kernel_mat_mul);
  533.     CHECK_ERROR(err);
  534.     err = clReleaseProgram(program);
  535.     CHECK_ERROR(err);
  536.     err = clReleaseMemObject(bufA);
  537.     CHECK_ERROR(err);
  538.     err = clReleaseMemObject(bufB);
  539.     CHECK_ERROR(err);
  540.     err = clReleaseMemObject(bufC);
  541.     CHECK_ERROR(err);
  542.     err = clReleaseCommandQueue(queue);
  543.     CHECK_ERROR(err);
  544.     err = clReleaseContext(context);
  545.     CHECK_ERROR(err);
  546.  
  547.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  548.     return C;
  549. }
  550.  
  551. int* mat_mul_opencl_4(int* A, int* B) {
  552.     cl_uint num_platforms;
  553.     cl_platform_id *platforms;
  554.     cl_uint num_devices;
  555.     char str[1024];
  556.     cl_device_type device_type;
  557.     cl_device_id device;
  558.     cl_context context;
  559.     cl_command_queue queue;
  560.     cl_program program;
  561.     char *kernel_source;
  562.     size_t kernel_source_size;
  563.     cl_kernel kernel_mat_mul;
  564.     cl_mem bufA, bufB, bufC;
  565.     cl_int err;
  566.  
  567.     size_t global_size[2] = { ROW_A, COL_B };
  568.     size_t local_size[2] = { 16, 16 };
  569.  
  570.     int i;
  571.  
  572.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  573.  
  574.     err = clGetPlatformIDs(0, NULL, &num_platforms);
  575.     CHECK_ERROR(err);
  576.  
  577.     platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  578.     err = clGetPlatformIDs(num_platforms, platforms, NULL);
  579.     CHECK_ERROR(err);
  580.  
  581.     err = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  582.     CHECK_ERROR(err);
  583.  
  584.     // Create Context
  585.     context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  586.     CHECK_ERROR(err);
  587.  
  588.     // Create Command Queue
  589.     queue = clCreateCommandQueue(context, device, 0, &err);
  590.     CHECK_ERROR(err);
  591.  
  592.     // Create Program Object
  593.     kernel_source = get_source_code("kernel4.cl", &kernel_source_size);
  594.     program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
  595.     CHECK_ERROR(err);
  596.  
  597.     // Build Program
  598.     err = clBuildProgram(program, 1, &device, "", NULL, NULL);
  599.     CHECK_ERROR(err);
  600.  
  601.     // Create Kernel
  602.     kernel_mat_mul = clCreateKernel(program, "mat_mul", &err);
  603.     CHECK_ERROR(err);
  604.  
  605.     // Write Buffer
  606.     bufA = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_A, A, &err);
  607.     CHECK_ERROR(err);
  608.     bufB = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_B * COL_B, B, &err);
  609.     CHECK_ERROR(err);
  610.     bufC = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_B, C, &err);
  611.     CHECK_ERROR(err);
  612.  
  613.     // Set Kernel Arg
  614.     err = clSetKernelArg(kernel_mat_mul, 0, sizeof(cl_mem), &bufA);
  615.     CHECK_ERROR(err);
  616.     err = clSetKernelArg(kernel_mat_mul, 1, sizeof(cl_mem), &bufB);
  617.     CHECK_ERROR(err);
  618.     err = clSetKernelArg(kernel_mat_mul, 2, sizeof(cl_mem), &bufC);
  619.     CHECK_ERROR(err);
  620.     err = clSetKernelArg(kernel_mat_mul, 3, sizeof(int), &ROW_A);
  621.     CHECK_ERROR(err);
  622.     err = clSetKernelArg(kernel_mat_mul, 4, sizeof(int), &COL_B);
  623.     CHECK_ERROR(err);
  624.     err = clSetKernelArg(kernel_mat_mul, 5, sizeof(int), &COL_A);
  625.     CHECK_ERROR(err);
  626.  
  627.     start = clock();
  628.  
  629.     // Execute Kernel
  630.     clEnqueueNDRangeKernel(queue, kernel_mat_mul, 2, NULL, global_size, local_size, 0, NULL, NULL);
  631.  
  632.     err = clFinish(queue);
  633.     CHECK_ERROR(err);
  634.  
  635.     end = clock();
  636.  
  637.     // Read Buffer
  638.     err = clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(int) * ROW_A * COL_B, C, 0, NULL, NULL);
  639.     CHECK_ERROR(err);
  640.  
  641.     err = clReleaseKernel(kernel_mat_mul);
  642.     CHECK_ERROR(err);
  643.     err = clReleaseProgram(program);
  644.     CHECK_ERROR(err);
  645.     err = clReleaseMemObject(bufA);
  646.     CHECK_ERROR(err);
  647.     err = clReleaseMemObject(bufB);
  648.     CHECK_ERROR(err);
  649.     err = clReleaseMemObject(bufC);
  650.     CHECK_ERROR(err);
  651.     err = clReleaseCommandQueue(queue);
  652.     CHECK_ERROR(err);
  653.     err = clReleaseContext(context);
  654.     CHECK_ERROR(err);
  655.  
  656.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  657.     return C;
  658. }
  659.  
  660. int* mat_mul_opencl_5(int* A, int* B) {
  661.     cl_uint num_platforms;
  662.     cl_platform_id *platforms;
  663.     cl_uint num_devices;
  664.     char str[1024];
  665.     cl_device_type device_type;
  666.     cl_device_id device;
  667.     cl_context context;
  668.     cl_command_queue queue;
  669.     cl_program program;
  670.     char *kernel_source;
  671.     size_t kernel_source_size;
  672.     cl_kernel kernel_mat_mul;
  673.     cl_mem bufA, bufB, bufC;
  674.     cl_int err;
  675.  
  676.     size_t global_size[1] = { ROW_A };
  677.     size_t local_size[1] = { 256 };
  678.  
  679.     int i;
  680.  
  681.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  682.     for (i = 0; i < ROW_A * COL_B; i++) C[i] = 0;
  683.  
  684.     err = clGetPlatformIDs(0, NULL, &num_platforms);
  685.     CHECK_ERROR(err);
  686.  
  687.     platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  688.     err = clGetPlatformIDs(num_platforms, platforms, NULL);
  689.     CHECK_ERROR(err);
  690.  
  691.     err = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  692.     CHECK_ERROR(err);
  693.  
  694.     // Create Context
  695.     context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  696.     CHECK_ERROR(err);
  697.  
  698.     // Create Command Queue
  699.     queue = clCreateCommandQueue(context, device, 0, &err);
  700.     CHECK_ERROR(err);
  701.  
  702.     // Create Program Object
  703.     kernel_source = get_source_code("kernel5.cl", &kernel_source_size);
  704.     program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
  705.     CHECK_ERROR(err);
  706.  
  707.     // Build Program
  708.     err = clBuildProgram(program, 1, &device, "", NULL, NULL);
  709.     CHECK_ERROR(err);
  710.  
  711.     // Create Kernel
  712.     kernel_mat_mul = clCreateKernel(program, "mat_mul", &err);
  713.     CHECK_ERROR(err);
  714.  
  715.     // Write Buffer
  716.     bufA = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_A, A, &err);
  717.     CHECK_ERROR(err);
  718.     bufB = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_B * COL_B, B, &err);
  719.     CHECK_ERROR(err);
  720.     bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_B, C, &err);
  721.     CHECK_ERROR(err);
  722.  
  723.     // Set Kernel Arg
  724.     err = clSetKernelArg(kernel_mat_mul, 0, sizeof(cl_mem), &bufA);
  725.     CHECK_ERROR(err);
  726.     err = clSetKernelArg(kernel_mat_mul, 1, sizeof(cl_mem), &bufB);
  727.     CHECK_ERROR(err);
  728.     err = clSetKernelArg(kernel_mat_mul, 2, sizeof(cl_mem), &bufC);
  729.     CHECK_ERROR(err);
  730.     err = clSetKernelArg(kernel_mat_mul, 3, sizeof(int), &ROW_A);
  731.     CHECK_ERROR(err);
  732.     err = clSetKernelArg(kernel_mat_mul, 4, sizeof(int), &COL_B);
  733.     CHECK_ERROR(err);
  734.     err = clSetKernelArg(kernel_mat_mul, 5, sizeof(int), &COL_A);
  735.     CHECK_ERROR(err);
  736.     err = clSetKernelArg(kernel_mat_mul, 6, sizeof(int) * COL_B, NULL);
  737.     CHECK_ERROR(err);
  738.  
  739.     start = clock();
  740.  
  741.     // Execute Kernel
  742.     clEnqueueNDRangeKernel(queue, kernel_mat_mul, 1, NULL, global_size, local_size, 0, NULL, NULL);
  743.  
  744.     err = clFinish(queue);
  745.     CHECK_ERROR(err);
  746.  
  747.     end = clock();
  748.  
  749.     // Read Buffer
  750.     err = clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(int) * ROW_A * COL_B, C, 0, NULL, NULL);
  751.     CHECK_ERROR(err);
  752.  
  753.     err = clReleaseKernel(kernel_mat_mul);
  754.     CHECK_ERROR(err);
  755.     err = clReleaseProgram(program);
  756.     CHECK_ERROR(err);
  757.     err = clReleaseMemObject(bufA);
  758.     CHECK_ERROR(err);
  759.     err = clReleaseMemObject(bufB);
  760.     CHECK_ERROR(err);
  761.     err = clReleaseMemObject(bufC);
  762.     CHECK_ERROR(err);
  763.     err = clReleaseCommandQueue(queue);
  764.     CHECK_ERROR(err);
  765.     err = clReleaseContext(context);
  766.     CHECK_ERROR(err);
  767.  
  768.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  769.     return C;
  770. }
  771.  
  772. int* mat_mul_opencl_6(int* A, int* B) {
  773.     cl_uint num_platforms;
  774.     cl_platform_id *platforms;
  775.     cl_uint num_devices;
  776.     char str[1024];
  777.     cl_device_type device_type;
  778.     cl_device_id device;
  779.     cl_context context;
  780.     cl_command_queue queue;
  781.     cl_program program;
  782.     char *kernel_source;
  783.     size_t kernel_source_size;
  784.     cl_kernel kernel_mat_mul;
  785.     cl_mem bufA, bufB, bufC;
  786.     cl_int err;
  787.  
  788.     size_t global_size[2] = { ROW_A, COL_B };
  789.     size_t local_size[2] = { 256, 1 };
  790.  
  791.     int i;
  792.  
  793.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  794.     for (i = 0; i < ROW_A * COL_B; i++) C[i] = 0;
  795.  
  796.     err = clGetPlatformIDs(0, NULL, &num_platforms);
  797.     CHECK_ERROR(err);
  798.  
  799.     platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  800.     err = clGetPlatformIDs(num_platforms, platforms, NULL);
  801.     CHECK_ERROR(err);
  802.  
  803.     err = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 1, &device, NULL);
  804.     CHECK_ERROR(err);
  805.  
  806.     // Create Context
  807.     context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
  808.     CHECK_ERROR(err);
  809.  
  810.     // Create Command Queue
  811.     queue = clCreateCommandQueue(context, device, 0, &err);
  812.     CHECK_ERROR(err);
  813.  
  814.     // Create Program Object
  815.     kernel_source = get_source_code("kernel6.cl", &kernel_source_size);
  816.     program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, &kernel_source_size, &err);
  817.     CHECK_ERROR(err);
  818.  
  819.     // Build Program
  820.     err = clBuildProgram(program, 1, &device, "", NULL, NULL);
  821.     CHECK_ERROR(err);
  822.  
  823.     // Create Kernel
  824.     kernel_mat_mul = clCreateKernel(program, "mat_mul", &err);
  825.     CHECK_ERROR(err);
  826.  
  827.     // Write Buffer
  828.     bufA = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_A, A, &err);
  829.     CHECK_ERROR(err);
  830.     bufB = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_B * COL_B, B, &err);
  831.     CHECK_ERROR(err);
  832.     bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int) * ROW_A * COL_B, C, &err);
  833.     CHECK_ERROR(err);
  834.  
  835.     // Set Kernel Arg
  836.     err = clSetKernelArg(kernel_mat_mul, 0, sizeof(cl_mem), &bufA);
  837.     CHECK_ERROR(err);
  838.     err = clSetKernelArg(kernel_mat_mul, 1, sizeof(cl_mem), &bufB);
  839.     CHECK_ERROR(err);
  840.     err = clSetKernelArg(kernel_mat_mul, 2, sizeof(cl_mem), &bufC);
  841.     CHECK_ERROR(err);
  842.     err = clSetKernelArg(kernel_mat_mul, 3, sizeof(int), &ROW_A);
  843.     CHECK_ERROR(err);
  844.     err = clSetKernelArg(kernel_mat_mul, 4, sizeof(int), &COL_B);
  845.     CHECK_ERROR(err);
  846.     err = clSetKernelArg(kernel_mat_mul, 5, sizeof(int), &COL_A);
  847.     CHECK_ERROR(err);
  848.     err = clSetKernelArg(kernel_mat_mul, 6, sizeof(int) * COL_B, NULL);
  849.     CHECK_ERROR(err);
  850.  
  851.     start = clock();
  852.  
  853.     // Execute Kernel
  854.     clEnqueueNDRangeKernel(queue, kernel_mat_mul, 2, NULL, global_size, local_size, 0, NULL, NULL);
  855.  
  856.     err = clFinish(queue);
  857.     CHECK_ERROR(err);
  858.  
  859.     end = clock();
  860.  
  861.     // Read Buffer
  862.     err = clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(int) * ROW_A * COL_B, C, 0, NULL, NULL);
  863.     CHECK_ERROR(err);
  864.  
  865.     err = clReleaseKernel(kernel_mat_mul);
  866.     CHECK_ERROR(err);
  867.     err = clReleaseProgram(program);
  868.     CHECK_ERROR(err);
  869.     err = clReleaseMemObject(bufA);
  870.     CHECK_ERROR(err);
  871.     err = clReleaseMemObject(bufB);
  872.     CHECK_ERROR(err);
  873.     err = clReleaseMemObject(bufC);
  874.     CHECK_ERROR(err);
  875.     err = clReleaseCommandQueue(queue);
  876.     CHECK_ERROR(err);
  877.     err = clReleaseContext(context);
  878.     CHECK_ERROR(err);
  879.  
  880.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  881.     return C;
  882. }
  883.  
  884. int* mat_mul_strassen(int* A, int* B) {
  885.     int* C = (int*)malloc(sizeof(int) * ROW_A * COL_B);
  886.     int i, j, k;
  887.  
  888.     start = clock();
  889.  
  890.     C = strassen(C, A, B, N);
  891.  
  892.     end = clock();
  893.  
  894.     printf("Elapsed time: %f sec\n", (double)(end - start) / CLK_TCK);
  895.  
  896.     return C;
  897. }
  898.  
  899. int *strassen(int *C, int *A, int *B, int n)
  900. {
  901.     int nn = n / 2;
  902.     int i, j;
  903.     int size = nn * nn;
  904.     int *buffer = (int *)malloc(sizeof(int) * size * 21);
  905.     int *a11 = buffer,
  906.         *a12 = buffer + size,
  907.         *a21 = buffer + size * 2,
  908.         *a22 = buffer + size * 3,
  909.         *b11 = buffer + size * 4,
  910.         *b12 = buffer + size * 5,
  911.         *b21 = buffer + size * 6,
  912.         *b22 = buffer + size * 7,
  913.         *c11 = buffer + size * 8,
  914.         *c12 = buffer + size * 9,
  915.         *c21 = buffer + size * 10,
  916.         *c22 = buffer + size * 11,
  917.         *m1 = buffer + size * 12,
  918.         *m2 = buffer + size * 13,
  919.         *m3 = buffer + size * 14,
  920.         *m4 = buffer + size * 15,
  921.         *m5 = buffer + size * 16,
  922.         *m6 = buffer + size * 17,
  923.         *m7 = buffer + size * 18,
  924.         *AA = buffer + size * 19,
  925.         *BB = buffer + size * 20;
  926.  
  927.     for (i = 0; i < nn; i++) {
  928.         for (j = 0; j < nn; j++) {
  929.             a11[i * nn + j] = A[i * n + j];
  930.             a12[i * nn + j] = A[i * n + j + nn];
  931.             a21[i * nn + j] = A[(i + nn) * n + j];
  932.             a22[i * nn + j] = A[(i + nn) * n + j + nn];
  933.  
  934.             b11[i * nn + j] = B[(i * n) + j];
  935.             b12[i * nn + j] = B[(i * n) + j + nn];
  936.             b21[i * nn + j] = B[(i + nn) * n + j];
  937.             b22[i * nn + j] = B[(i + nn) * n + j + nn];
  938.         }
  939.     }
  940.  
  941.     strassen_mul(m1, strassen_add(AA, a11, a22, nn), strassen_add(BB, b11, b22, nn), nn);
  942.     strassen_mul(m2, strassen_add(AA, a21, a22, nn), b11, nn);
  943.     strassen_mul(m3, a11, strassen_sub(BB, b12, b22, nn), nn);
  944.     strassen_mul(m4, a22, strassen_sub(BB, b21, b11, nn), nn);
  945.     strassen_mul(m5, strassen_add(AA, a11, a12, nn), b22, nn);
  946.     strassen_mul(m6, strassen_sub(AA, a21, a11, nn), strassen_add(BB, b11, b12, nn), nn);
  947.     strassen_mul(m7, strassen_sub(AA, a12, a22, nn), strassen_add(BB, b21, b22, nn), nn);
  948.  
  949.     strassen_sub(c11, strassen_add(AA, m1, m4, nn), strassen_sub(BB, m5, m7, nn), nn);
  950.     strassen_add(c12, m3, m5, nn);
  951.     strassen_add(c21, m2, m4, nn);
  952.     strassen_add(c22, strassen_sub(AA, m1, m2, nn), strassen_add(BB, m3, m6, nn), nn);
  953.  
  954.     strassen_result(C, c11, c12, c21, c22, nn);
  955.  
  956.     free(buffer);
  957.  
  958.     return C;
  959. }
  960.  
  961. void strassen_result(int *C, int *c11, int *c12, int *c21, int *c22, int n) {
  962.     int i, j;
  963.  
  964.     for (i = 0; i < n; i++)
  965.         for (j = 0; j < n; j++) {
  966.             C[i * 2 * n + j] = c11[i * n + j];
  967.             C[i * 2 * n + j + n] = c12[i * n + j];
  968.             C[(i + n) * n * 2 + j] = c21[i * n + j];
  969.             C[(i + n) * n * 2 + j + n] = c22[i * n + j];
  970.         }
  971. }
  972.  
  973. int *strassen_add(int *res, int *x, int *y, int n) {
  974.     int i, j;
  975.  
  976.     memset(res, 0, sizeof(int) * n * n);
  977.  
  978.     for (i = 0; i < n; i++)
  979.         for (j = 0; j < n; j++)
  980.             res[i * n + j] = x[i * n + j] + y[i * n + j];
  981.  
  982.     return res;
  983. }
  984.  
  985. int *strassen_sub(int *res, int *x, int *y, int n) {
  986.     int i, j;
  987.  
  988.     memset(res, 0, sizeof(int) * n * n);
  989.  
  990.     for (i = 0; i < n; i++)
  991.         for (j = 0; j < n; j++)
  992.             res[i * n + j] = x[i * n + j] - y[i * n + j];
  993.  
  994.     return res;
  995. }
  996.  
  997. int *strassen_mul(int *res, int *a, int *b, int n) {
  998.     int m1, m2, m3, m4, m5, m6, m7;
  999.  
  1000.     if (n == 2) {
  1001.         m1 = (a[0 * n + 0] + a[1 * n + 1]) * (b[0 * n + 0] + b[1 * n + 1]);
  1002.         m2 = (a[1 * n + 0] + a[1 * n + 1]) * b[0 * n + 0];
  1003.         m3 = a[0 * n + 0] * (b[0 * n + 1] - b[1 * n + 1]);
  1004.         m4 = a[1 * n + 1] * (b[1 * n + 0] - b[0 * n + 0]);
  1005.         m5 = (a[0 * n + 0] + a[0 * n + 1]) * b[1 * n + 1];
  1006.         m6 = (a[1 * n + 0] - a[0 * n + 0]) * (b[0 * n + 0] + b[0 * n + 1]);
  1007.         m7 = (a[0 * n + 1] - a[1 * n + 1]) * (b[1 * n + 0] + b[1 * n + 1]);
  1008.  
  1009.         res[0 * n + 0] = m1 + m4 - m5 + m7;
  1010.         res[0 * n + 1] = m3 + m5;
  1011.         res[1 * n + 0] = m2 + m4;
  1012.         res[1 * n + 1] = m1 - m2 + m3 + m6;
  1013.  
  1014.         return res;
  1015.     }
  1016.  
  1017.     return strassen(res, a, b, n);
  1018. }
RAW Paste Data
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy. OK, I Understand
 
Top