Advertisement
Guest User

Untitled

a guest
Feb 8th, 2022
53
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C 7.02 KB | None | 0 0
  1. #include <stdint.h>
  2. #include <stddef.h>
  3. #include <stdlib.h>
  4. #include <stdio.h>
  5. #include <stdint.h>
  6.  
  7. #define CL_TARGET_OPENCL_VERSION 200
  8.  
  9. #include <CL/cl.h>
  10.  
  11.  
  12. const char* resizeKernel =
  13. "__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n"
  14. "__kernel void resize_nn(read_only image2d_t source, write_only image2d_t dest, uint src_img_width, uint src_img_height, uint dst_img_width, uint dst_img_height, float ratioX, float ratioY)\n"
  15. "{\n"
  16. "   const int gx = get_global_id(0);\n"
  17. "   const int gy = get_global_id(1);\n"
  18. "   const int2 pos = { gx, gy };\n"
  19. "\n"
  20. "   if (pos.x >= dst_img_width || pos.y >= dst_img_height)\n"
  21. "      return;\n"
  22. "\n"
  23. "   float2 srcpos = {(pos.x + 0.4995f) / ratioX, (pos.y + 0.4995f) / ratioY};\n"
  24. "   int2 SrcSize = (int2)(src_img_width, src_img_height);\n"
  25. "\n"
  26. "   float4 value;\n"
  27. "\n"
  28. "   int2 ipos = convert_int2(srcpos);\n"
  29. "   if (ipos.x < 0 || ipos.x >= SrcSize.x || ipos.y < 0 || ipos.y >= SrcSize.y)\n"
  30. "      value = 0;\n"
  31. "   else\n"
  32. "      value = read_imagef(source, sampler, ipos);\n"
  33. "\n"
  34. "   write_imagef(dest, pos, value);\n"
  35. "}";
  36.  
  37. cl_device_id deviceID;
  38. cl_context deviceCtx;
  39. cl_command_queue queue;
  40. cl_program program;
  41.  
  42. int init() {
  43.     cl_int ret;
  44.     cl_uint len;
  45.     ret = clGetPlatformIDs(0, NULL, &len);
  46.     if (ret != 0) {
  47.         return ret;
  48.     }
  49.     cl_platform_id* platformIDs = malloc(len*sizeof(cl_platform_id));
  50.     ret = clGetPlatformIDs(len, platformIDs, NULL);
  51.     if (ret != 0) {
  52.         return ret;
  53.     }
  54.     for (int i=0;i<len;i++) {
  55.         ret = clGetDeviceIDs(platformIDs[i], CL_DEVICE_TYPE_GPU, 1, &deviceID, NULL);
  56.         if (ret == 0) {
  57.             break;
  58.         }
  59.         deviceID = NULL;
  60.     }
  61.     if (deviceID == NULL) {
  62.         return -100;
  63.     }
  64.     deviceCtx = clCreateContext(NULL, 1, &deviceID, NULL, NULL, &ret);
  65.     if (ret != 0) {
  66.         return ret;
  67.     }
  68.     queue = clCreateCommandQueueWithProperties(deviceCtx, deviceID, NULL, &ret);
  69.     if (ret != 0) {
  70.         return ret;
  71.     }
  72.     program = clCreateProgramWithSource(deviceCtx, 1, &resizeKernel, NULL, &ret);
  73.     if (ret != 0) {
  74.         return ret;
  75.     }
  76.     ret = clBuildProgram(program, 1, &deviceID, NULL, NULL, NULL);
  77.     if (ret != 0) {
  78.         return ret;
  79.     }
  80.  
  81.     return 0;
  82. }
  83.  
  84. size_t max(int a, int b) {
  85.     if (a > b) {
  86.         return a;
  87.     }
  88.     return b;
  89. }
  90.  
  91. void* createImage(
  92.         int width,
  93.         int height,
  94.         int row_pitch,
  95.         unsigned int image_channel_order,
  96.         unsigned int image_channel_data_type,
  97.         void* data,
  98.         cl_int* ret
  99.         ){
  100.     cl_image_format format;
  101. format.image_channel_order = image_channel_order;
  102.     format.image_channel_data_type = image_channel_data_type;
  103.     cl_mem_flags inFlags = CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR; //(1 << 2) | (1 << 5);
  104.     cl_image_desc inDesc;
  105.     inDesc.image_type = CL_MEM_OBJECT_IMAGE2D; //0x10F1;
  106.     inDesc.image_width = width;
  107.     inDesc.image_height = height;
  108.     inDesc.image_depth= 0;// not usedfor 2d
  109.     inDesc.image_array_size = 0;// not used for 2d
  110.     inDesc.image_row_pitch = 0;//(size_t)(row_pitch);
  111.     inDesc.image_slice_pitch = 0;
  112.     inDesc.num_mip_levels = 0;
  113.     inDesc.num_samples = 0;
  114.     inDesc.buffer = NULL;
  115.     cl_mem in = clCreateImage(deviceCtx, inFlags, &format, &inDesc, data, ret);
  116.     return (void*)(in);
  117. }
  118.  
  119. int releaseImage(void* img) {
  120.     return clReleaseMemObject((cl_mem)(img));
  121. }
  122.  
  123. int resize(
  124.         void *in,
  125.         int inWidth,
  126.         int inHeight,
  127.         int outWidth,
  128.         int outHeight,
  129.         int rowPitch,
  130.         unsigned int image_channel_order,
  131.         unsigned int image_channel_data_type,
  132.         void* outData
  133.         ) {
  134.     cl_int ret;
  135.  
  136.     cl_image_format format;
  137.     format.image_channel_order = image_channel_order;
  138.     format.image_channel_data_type = image_channel_data_type;
  139.  
  140.     //output image
  141.     cl_mem_flags outFlags = CL_MEM_WRITE_ONLY;
  142.     cl_image_desc outDesc;
  143.     outDesc.image_type = CL_MEM_OBJECT_IMAGE2D; //0x10F1;
  144.     outDesc.image_width = outWidth;
  145.     outDesc.image_height = outHeight;
  146.     outDesc.image_row_pitch = 0;//rowPitch;
  147.     outDesc.image_slice_pitch = 0;
  148.     outDesc.num_mip_levels = 0;
  149.     outDesc.num_samples = 0;
  150.     outDesc.buffer = NULL;
  151.     cl_mem out = clCreateImage(deviceCtx, outFlags, &format, &outDesc, NULL, &ret);
  152.     if (ret != 0) {
  153.         return ret;
  154.     }
  155.  
  156.     //kernel
  157.     cl_kernel kernel = clCreateKernel(program, "resize_nn", &ret);
  158.     if (ret != 0) {
  159.         return ret;
  160.     }
  161.     ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &in);
  162.     if (ret != 0) {
  163.         return ret;
  164.     }
  165.     ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), &out);
  166.     if (ret != 0) {
  167.         return ret;
  168.     }
  169.     ret = clSetKernelArg(kernel, 2, sizeof(int), &inWidth);
  170.     if (ret != 0) {
  171.         return ret;
  172.     }
  173.     ret = clSetKernelArg(kernel, 3, sizeof(int), &inHeight);
  174.     if (ret != 0) {
  175.         return ret;
  176.     }
  177.     ret = clSetKernelArg(kernel, 4, sizeof(int), &outWidth);
  178.     if (ret != 0) {
  179.         return ret;
  180.     }
  181.     ret = clSetKernelArg(kernel, 5, sizeof(int), &outHeight);
  182.     if (ret != 0) {
  183.         return ret;
  184.     }
  185.     float ratioX = (float)(outWidth) / (float)(inWidth);
  186.     ret = clSetKernelArg(kernel, 6, sizeof(float), &ratioX);
  187.     if (ret != 0) {
  188.         return ret;
  189.     }
  190.     float ratioY = (float)(outHeight) / (float)(inHeight);
  191.     ret = clSetKernelArg(kernel, 7, sizeof(float), &ratioY);
  192.     if (ret != 0) {
  193.         return ret;
  194.     }
  195.  
  196.     //run kernel
  197.     cl_event kernelEvent;
  198.     size_t globalSize[] = {max(inWidth, outWidth), max(inHeight, outHeight)};
  199.     ret = clEnqueueNDRangeKernel(
  200.             queue,
  201.             kernel,
  202.             2, //work_dim
  203.             0, //global_work_offset
  204.             globalSize, //global_work_size
  205.             0, //local_work_size
  206.             0, //num_events_in_wait_list
  207.             NULL, //event_wait_list
  208.             &kernelEvent //event
  209.             );
  210.     if (ret != 0) {
  211.         return ret;
  212.     }
  213.     ret = clWaitForEvents(1, &kernelEvent);
  214.     if (ret != 0) {
  215.         return ret;
  216.     }
  217.     ret = clReleaseEvent(kernelEvent);
  218.     if (ret != 0) {
  219.         return ret;
  220.     }
  221.  
  222.     //get image
  223.     cl_event readImageEvent;
  224.     size_t origin[] = {0,0,0};
  225.     size_t region[] = {outWidth,outHeight,1};
  226.     ret = clEnqueueReadImage(
  227.             queue,
  228.             out,
  229.             CL_TRUE,
  230.             origin,
  231.             region,
  232.             0,
  233.             0,
  234.             outData,
  235.             0,
  236.             NULL,
  237.             &readImageEvent
  238.             );
  239.     if (ret != 0) {
  240.         return ret;
  241.     }
  242.     ret = clWaitForEvents(1, &readImageEvent);
  243.     if (ret != 0) {
  244.         return ret;
  245.     }
  246.     ret = clReleaseEvent(readImageEvent);
  247.     if (ret != 0) {
  248.         return ret;
  249.     }
  250.     ret = clReleaseMemObject(out);
  251.     if (ret != 0) {
  252.         return ret;
  253.     }
  254.     ret = clReleaseKernel(kernel);
  255.     if (ret != 0) {
  256.         return ret;
  257.     }
  258.  
  259.     return 0;
  260. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement