#include <vector> 
#include <thread>  
#include <functional>  

int main() {
     const int size_x = 1920;
     const int size_y = 1080;
     uint8_t frame[size_x * size_y], buffer[size_x * size_y];
     load_grayscale_image(frame, size_x, size_y);
     memset(buffer, 0, size_x * size_y * sizeof(uint8_t));

     auto f = [](uint8_t *in, uint8_t *out,
                 int start, int end, int shift) {
         out[start] = in[start];
         for (int offset = start + shift; offset < end; offset++)
             out[offset] = in[offset - 1] / 4 + in[offset] / 2 + in[offset - 1] / 4;
         out[end] = in[end];
     };

   std::vector<std::thread>threads;
   threads.resize(size_y);
   for (int y = 0; y < threads.size(); y++)
       threads[y] = std::thread(f, frame, buffer,
                                y * size_x, (y - 1) * size_x - 1, 1);

   for (auto &t : threads) t.join();

   threads.resize(size_x); 
   for (int x = 0; x < threads.size(); x++) 
      threads[x] = std::thread(f, buffer, frame x, x + (size_y - 1) * size_x, size_x); 
   for (auto &t : threads) t.join(); 
   return 0; }

// Searching passwords hashed with MD5 
#include "maltcc.h"
#include "md5.c"

atomic_t found;
void cpu_check_md5(const char* s) {
  Hash h = HashFn(MD5_Id,s);
  for (int i = 0; i < N_HASHES; i++)        // Result check
    if (!memcmp(h.bytes(),hashes[i],16)) {
      printf("\tFound! %s\n",s); 
      found++;
    }
}

__simd __kernel check_words(u32 w[W_LEN]) { // Executed on accelerators 
  u32 i, j, r;
  u32 hs[N_HASHES], abc[ABC_SZ];                  
  __slave on_init() {
    for (int i=0; i < N_HASHES; i++)        // First words of hashes
      hs[i] = cpu_to_le(hashes[i][0]);      // for express check
    for (unsigned i = 0; i < ABC_SZ; i++)   
      abc[i] = alphabet[i];
  }
  j = 0; REPEAT(ABC_SZ) {            
    w[0] = (w[0] & 0xFFFF) | abc[j] << 16;  j++;
    MD5(w,r); 
    i = 0; REPEAT(N_HASHES) {
        __success |= r==hs[i]; i++;
    }
  }
  __slave on_success() {                    // Callback on success
    u32 ws = le_to_cpu(w[0]);
    char *s = (char*)&ws; s[W_SZ]=0;
    printf("Found? [w=%s]\n",s);
    for (int i = 0; i < ABC_SZ; i++) {      // Looking for all endings
      s[W_SZ-1] = alphabet[i];
      cpu_check_md5(s);
    }
  }
} 

__slave dict_fragment(int d) {              // Execution on scalar cores
  for (u32 i=0; i<ABC_SZ; i++) {            
    u32 w = (u32)(alphabet[d]) | (u32)(alphabet[i]) << 8;
    check_words(&w);                        // Checking series of words on vector accelerators
  }
}

int main(int argc, char *argv[]) {
  for (int i=0; i < ABC_SZ; i++)
    dict_fragment(i);                       // Execution on all scalar cores
  slave_sync();
  return found;  
}

NVIDIA® and CUDA® are the registered trademarks of NVIDIA Corporation

// OpenCL-core example code for marking specific pixels of an image

#define DIM_X   128
#define DIM_Y   128

#define M       9
#define MASK(n) (masks+(n)*M*M)
#define HESS_L  100

int D(__global uchar* img, int x, int y, __constant uchar *mask) {
  x -= M/2; y -= M/2; int v = 0;
  for (int dx=0; dx<M; dx++)
    for (int dy=0; dy<M; dy++) 
      v += mask[dx+M*dy] * img[(x+dx)+DIM_X*(y+dy)];
  return v;    
}

__kernel void mark_ipoints(__global uchar *img, __constant uchar* masks) {
  for (int x=M/2; x<DIM_X-M/2; x++) { 
    for (int y=M/2; y<DIM_Y-M/2; y++) {
      int Dxx = D(img,x,y,MASK(0)), 
          Dyy = D(img,x,y,MASK(1)), 
          Dxy = D(img,x,y,MASK(2));
      img[x+DIM_X*y] |= (abs(5*Dxx*Dyy - 4*Dxy*Dxy) > HESS_L) << 7; 
    } 
  }
}