🚀 go-pugleaf

W dniu piątek, 20 września 2019 15:00:16 UTC+2 użytkownik M.M. napisał:
> On Thursday, September 19, 2019 at 4:55:42 AM UTC+2, krop...@gmail.com wrote:
> > W dniu wtorek, 17 września 2019 21:15:49 UTC+2 użytkownik M.M. napisał:
> > > Ciekawe jak w praktyce wygląda przyspieszenie obliczeń na Tesli 
> > > względem Xenon Phi. I ciekawe czy w ogole warto inwestować w te drogie
> > > rozwiązania, jak za 150usd można kupić: GeForce GTX 1650 - tania,
> > > wydajna, a 10 takich kart z lekkim underclockingiem na niektórych
> > > obliczeniach pobiera tylko 500 wat mocy. No ale na GPU trzeba się
> > > nauczyć jakiegoś OpenCL albo CUDA.
> > Jak ktoś umie napisać sensownie równoległy kod to OpenCL/CUDA nie są
> > żadnym problemem. Owszem, to nie są jakieś przepiękne API, ale 
> > praktycznie każdy ostro sprzętowy kod w C wygląda równie źle;)
> > Do tego można się tego uczyć bezproblemowo na dowolnym komputerze
> > z jakąś sensowną kartą graficzną, odpalenie Hello World na CUDA
> > (domyślny program to bodajże jakieś równoległe dodawanie wektorów
> > czy tam sortowanie) zajmuje jakieś 15 minut, z czego 10 to rejestracja 
> > na stronie nvidii żeby ściągnać toolchain. To jest wręcz przerażająco
> > łatwe w porównaniu do dawnego oprogramowywania "grubych" platform
> > obliczeniowych gdzie potrzeba było komercyjnego kompilatora za grubą
> > kasę na uczelnianym klastrze i wczytywania się w dokumentację żeby
> > coś się w ogóle uruchomiło.
> > 
> > Pozdrawiam, 
> > -- 
> > Karol Piotrowski
> 
> Możesz polecić jakiś praktyczny tutorial od podstaw OpenCLa dla kogoś, kto wie 
> co to C++ i programowanie rownoległe na CPU, ale z GPU nie miał nigdy 
> do czynienia?
> 
> Pozdrawiam

moge wkleic jakies moje stare notatki/proby z opencl (z 2015, nie pamietam czy to ostatnia wersja ale pamietam ze to dzialalo, i wersja opencl renderowania madelbrota byla najszybsza, pisalem o tym na grupie wiec jest o tym gdzies watek w 2015)


ale szczerze mowiac nie wiem czy polecalbym sie tym zajmowac, jesli ktos chce zbierac odznaki programistycznego skauta i wpisac opencl do cv to na pewno 
- w innym wypadku chyba nie


#include "fist.h"
////////////////////////////////////////////////////////////////////////////////

//#include <fcntl.h>
#include <stdio.h>
#include <stdlib.h>
//#include <string.h>
#include <math.h>
//#include <unistd.h>
//#include <sys/types.h>
//#include <sys/stat.h>
//#include <OpenCL/opencl.h>
//#include "clew.c"
#include "clew.h"
int clewInit(const char* path);

////////////////////////////////////////////////////////////////////////////////

// Use a static data size for simplicity
//
//const int DATA_SIZE = 130*1024;

const int OCL_INPUT_DATA_SIZE = 1*1024;
const int OCL_OUTPUT_DATA_SIZE =  256*256;
const int OCL_PROCESSING_RANGE =  256*256;

//int   DATA_SIZE = DATA_SIZE_MAX/10;

////////////////////////////////////////////////////////////////////////////////

// Simple compute kernel which computes the square of an input array
//
// const char *KernelSource = "\n" \
// "__kernel void square(                                                       \n" \
// "   __global float* input,                                              \n" \
// "   __global float* output,                                             \n" \
// "   const unsigned int count)                                           \n" \
// "{                                                                      \n" \
// "   int i = get_global_id(0);                                           \n" \
// "   if(i < count)                                                       \n" \
// "       output[i] = input[i] * input[i];                                \n" \
// "}                                                                      \n" \
// "\n";

// const char *KernelSource = "\n" \
// "__kernel void square(                                                       \n" \
// "   __global int* input,                                              \n" \
// "   __global int* output,                                             \n" \
// "   const unsigned int count)                                           \n" \
// "{                                                                      \n" \
// "   int i = get_global_id(0);                                           \n" \
// "   if(i < count)                                                       \n" \
// "       output[i] = i+input[i] *2   ;                                \n" \
// "}                                                                      \n" \
// "\n";


const char *KernelSource = "\n" \
// "__kernel void square(                                                       \n" \
// "   __global int* input,                                              \n" \
// "   __global int* output,                                             \n" \
// "   const unsigned int count)                                           \n" \
// "{                                                                      \n" \
// "   int i = get_global_id(0);                                           \n" \
// "   if(i < count)       \n" \
// "   {   int x = i%550; int y=i/550;                                                       \n" \
// "       output[i] = x+input[i] +y   ;                                \n" \
// "   }        \n" \
// "}                                                                      \n" \
// "\n";
"__kernel void square(                                                \n" \
"   __global int* input,                                              \n" \
"   __global int* output,                                             \n" \
"   const unsigned int count)                                         \n" \
"{                                                                    \n" \
"   int i = get_global_id(0);                                         \n" \
"   if(i < count)                                 \n" \
"   {                                             \n" \
"   int x = i%256;                  \n" \
" //  if(x>=256) return;                  \n" \
"   int y = i/256;                   \n" \
" //  if(y>=256) return;                  \n" \
"    float cRe = -0.5 + -1.5 + x/256.*3.;                 \n" \
"    float cIm =  0.0 + -1.5 + y/256.*3.;                  \n" \
"   float re = 0;    \n" \
"   float im = 0;     \n" \
"   int n = 0;     \n" \
"   for( n=0; n<=1000; n++)   {   \n" \
"     if( re * re + im * im > 4.0 )  { output[256*y+x] = n + 256*n + 256*256*n; return;}   \n" \
"     float re_n = re * re - im * im + cRe;      \n" \
"     float im_n = 2 * re * im       + cIm;      \n" \
"     re = re_n;          \n" \
"     im = im_n;          \n" \
"     }                   \n" \
"   output[256*y+x] = 250<<8;   \n" \
"   }                     \n" \
"}                         \n" \
"\n";

    int data[OCL_INPUT_DATA_SIZE];              // original data set given to device
    int results[OCL_OUTPUT_DATA_SIZE];           // results returned from device

// void SetupInputData()
// {
//     // Fill our data set with random float values
//     //
//     int i = 0;
//     unsigned int count = DATA_SIZE;
//     for(i = 0; i < count; i++)
//         data[i] = rand() / (float)RAND_MAX;
// }
////////////////////////////////////////////////////////////////////////////////
    int err;                            // error code returned from api calls
    unsigned int correct;               // number of correct results returned

    size_t global;                      // global domain size for our calculation
    size_t local;                       // local domain size for our calculation

    cl_device_id device_id;             // compute device id
    cl_context context;                 // compute context
    cl_command_queue commands;          // compute command queue
    cl_program program;                 // compute program
    cl_kernel kernel;                   // compute kernel

    cl_mem input;                       // device memory used for the input array
    cl_mem output;                      // device memory used for the output array

int SetupCL()
{
  static int initialised =0;
  if(initialised) return 0;
  initialised=1;

    clewInit("OpenCl.dll");

    int writeCLInfo();

  //  writeCLInfo();




  //////// platform
   static cl_platform_id platform_id[10] = {0};
   cl_uint no_of_platforms_found = 0;

   int ret = clGetPlatformIDs(10, platform_id, &no_of_platforms_found);
 //  if(ret == CL_SUCCESS ) ERROR_("clGetPlatformIDs success") ;
  // alert(" %d platforms found\n",no_of_platforms_found ) ;

   // Connect to a compute device
   //
    int gpu = 1;
    err = clGetDeviceIDs(platform_id[0], gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
    if (err != CL_SUCCESS)
    {
        ERROR_("Error: Failed to create a device group!\n");
        return EXIT_FAILURE;
    }

    // Create a compute context
    //
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    if (!context)
    {
        ERROR_("Error: Failed to create a compute context!\n");
        return EXIT_FAILURE;
    }

    // Create a command commands
    //
    commands = clCreateCommandQueue(context, device_id, 0, &err);
    if (!commands)
    {
        ERROR_("Error: Failed to create a command commands!\n");
        return EXIT_FAILURE;
    }

    // Create the compute program from the source buffer
    //
    program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err);
    if (!program)
    {
        ERROR_("Error: Failed to create compute program!\n");
        return EXIT_FAILURE;
    }

    // Build the program executable
    //
    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        size_t len;
        char buffer[2048];

        ERROR_("Error: Failed to build program executable!\n");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        ERROR__("%s\n", buffer);
        exit(1);
    }

    // Create the compute kernel in the program we wish to run
    //
    kernel = clCreateKernel(program, "square", &err);
    if (!kernel || err != CL_SUCCESS)
    {
        ERROR_("Error: Failed to create compute kernel! \n");
        exit(1);
    }

    // Create the input and output arrays in device memory for our calculation
    //
    if(OCL_INPUT_DATA_SIZE>0)
     input  = clCreateBuffer(context,  CL_MEM_READ_ONLY,  sizeof(float) * OCL_INPUT_DATA_SIZE, NULL, NULL);

    output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * OCL_OUTPUT_DATA_SIZE, NULL, NULL);

    if (!input || !output)
    {
        ERROR_("Error: Failed to allocate device memory!\n");
        exit(1);
    }

    ////////////////////////



    // Set the arguments to our compute kernel
    //
    err = 0;
    int count = OCL_PROCESSING_RANGE;

    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
    err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);

    if (err != CL_SUCCESS)
    {
        ERROR__("Error: Failed to set kernel arguments! %d\n", err);
        exit(1);
    }


    // Get the maximum work group size for executing the kernel on the device
    //
    err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
    if (err != CL_SUCCESS)
    {
        ERROR__("Error: Failed to retrieve kernel work group info! %d\n", err);
        exit(1);
    }

    return 0;
}

///////////////////

int RunOpenCLtask()
{
  //    DATA_SIZE = frame_size_x*frame_size_y/4;

//     SetupInputData();

   if(0)
     for(int i = 0; i < OCL_INPUT_DATA_SIZE; i++)
     {
         data[i] = frame_bitmap[i] ;
     }


    // Write our data set into the input array in device memory
    //
   if(OCL_INPUT_DATA_SIZE>0)
   {
    err = clEnqueueWriteBuffer(commands, input, CL_TRUE, 0, sizeof(float) * OCL_INPUT_DATA_SIZE, data, 0, NULL, NULL);

    if (err != CL_SUCCESS)
         ERROR_EXIT("Error: Failed to write to source array!\n");
   }



    // Execute the kernel over the entire range of our 1d input data set
    // using the maximum number of work group items for this device
    //
    global = OCL_PROCESSING_RANGE;

    err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
    if (err)
         ERROR_EXIT("Error: Failed to execute kernel!\n");

    // Wait for the command commands to get serviced before reading back results
    //
    clFinish(commands);

    // Read back the results from the device to verify the output
    //
    err = clEnqueueReadBuffer( commands, output, CL_TRUE, 0, sizeof(float) * OCL_OUTPUT_DATA_SIZE, results, 0, NULL, NULL );
    if (err != CL_SUCCESS)
    {
        ERROR__("Error: Failed to read output array! %d\n", err);
        exit(1);
    }

    if(1)
     for(int i = 0; i < OCL_OUTPUT_DATA_SIZE; i++)
     {
         int x = i%256;
         int y = i/256;
         frame_bitmap[y*550+x] = results[i];
     }

    // Validate our results
    //
//     correct = 0;
//     for(int i = 0; i < count; i++)
//     {
//         if(results[i] - data[i] * data[i]<0.001)
//             correct++;
//     }
//
//     for(int i=0; i<1024; i++)
//     {
//         printf("data %f square %f \n", data[i], data[i]*data[i]);
//         printf("result %f \n", results[i]);
//     }


    // Print a brief summary detailing the results
    //
//    printf("Computed '%d/%d' correct values!\n", correct, count);
 //   alert("Computed '%d/%d' correct values!\n", correct, count);

}

void ShutdownOpenCl()
{
   // Shutdown and cleanup
    //
    clReleaseMemObject(input);
    clReleaseMemObject(output);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(commands);
    clReleaseContext(context);

}

int writeCLInfo()
{



  int i, j;
  char* value;
  size_t valueSize;
  cl_uint platformCount;
  cl_platform_id* platforms;
  cl_uint deviceCount;
  cl_device_id* devices;
  cl_uint maxComputeUnits;



  // get all platforms

  clGetPlatformIDs(0, NULL, &platformCount);

  platforms = (cl_platform_id*) malloc(sizeof(cl_platform_id) * platformCount);

  clGetPlatformIDs(platformCount, platforms, NULL);



  for (i = 0; i < platformCount; i++) {



      // get all devices

      clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &deviceCount);

      devices = (cl_device_id*) malloc(sizeof(cl_device_id) * deviceCount);

      clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, deviceCount, devices, NULL);



      // for each device print critical attributes

      for (j = 0; j < deviceCount; j++) {



          // print device name

          clGetDeviceInfo(devices[j], CL_DEVICE_NAME, 0, NULL, &valueSize);

          value = (char*) malloc(valueSize);

          clGetDeviceInfo(devices[j], CL_DEVICE_NAME, valueSize, value, NULL);

          printf("%d. Device: %s\n", j+1, value);

          free(value);



          // print hardware device version

          clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, 0, NULL, &valueSize);

          value = (char*) malloc(valueSize);

          clGetDeviceInfo(devices[j], CL_DEVICE_VERSION, valueSize, value, NULL);

          printf(" %d.%d Hardware version: %s\n", j+1, 1, value);

          free(value);



          // print software driver version

          clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, 0, NULL, &valueSize);

          value = (char*) malloc(valueSize);

          clGetDeviceInfo(devices[j], CL_DRIVER_VERSION, valueSize, value, NULL);

          printf(" %d.%d Software version: %s\n", j+1, 2, value);

          free(value);



          // print c version supported by compiler for device

          clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, 0, NULL, &valueSize);

          value = (char*) malloc(valueSize);

          clGetDeviceInfo(devices[j], CL_DEVICE_OPENCL_C_VERSION, valueSize, value, NULL);

          printf(" %d.%d OpenCL C version: %s\n", j+1, 3, value);

          free(value);



          // print parallel compute units

          clGetDeviceInfo(devices[j], CL_DEVICE_MAX_COMPUTE_UNITS,

                  sizeof(maxComputeUnits), &maxComputeUnits, NULL);

          printf(" %d.%d Parallel compute units: %d\n", j+1, 4, maxComputeUnits);



      }



      free(devices);



  }



  free(platforms);

    return 0;
}