success with aparapi

Speed test: SHA-256
typedef struct This_s{
   __global int *currentVal;
   __global int *W;
   __global char  *inputBuf;
   __constant int *K;
   int passid;
}This;
int get_pass_id(This *this){
   return this->passid;
}
int fr_cryptohash_SHA2Core$MyKernel__r3(This *this, int i1, int i2, int i3){
   return(((i1 ^ i2) ^ i3));
}
__kernel void run(
   __global int *currentVal,
   __global int *W,
   __global char  *inputBuf,
   __constant int *K,
   int passid
){
   This thisStruct;
   This* this=&thisStruct;
   this->currentVal = currentVal;
   this->W = W;
   this->inputBuf = inputBuf;
   this->K = K;
   this->passid = passid;
   {
      int H = this->currentVal[7];
      int G = this->currentVal[6];
      int F = this->currentVal[5];
      int E = this->currentVal[4];
      int D = this->currentVal[3];
      int C = this->currentVal[2];
      int B = this->currentVal[1];
      int A = this->currentVal[0];
      for (int i = 0; i<16; i++){
         {
            int off = i << 2;
            this->W[i]  = (((this->inputBuf[off++] << 24) | (this->inputBuf[off++] << 16)) | (this->inputBuf[off++] << 8)) | this->inputBuf[off];
         }
      }
      int T2 = 0;
      int T1 = 0;
      for (int i = 16; i<64; i++){
         T1 = this->W[(i - 2)];
         T2 = this->W[(i - 15)];
         this->W[i]  = ((fr_cryptohash_SHA2Core$MyKernel__r3(this, ((T1 << 15) | (((unsigned int)T1) >> 17)), ((((unsigned int)T1) >> 19) | (T1 << 13)), (((unsigned int)T1) >> 10)) + this->W[(i - 7)]) + fr_cryptohash_SHA2Core$MyKernel__r3(this, ((T2 << 25) | (((unsigned int)T2) >> 7)), ((((unsigned int)T2) >> 18) | (T2 << 14)), (((unsigned int)T2) >> 3))) + this->W[(i - 16)];
      }
      for (int i = 0; i<64; i++){
         T1 = (((H + fr_cryptohash_SHA2Core$MyKernel__r3(this, ((E << 26) | (((unsigned int)E) >> 6)), ((((unsigned int)E) >> 11) | (E << 21)), ((E << 7) | (((unsigned int)E) >> 25)))) + ((F & E) ^ (G & (E ^ -1)))) + this->K[i]) + this->W[i];
         T2 = fr_cryptohash_SHA2Core$MyKernel__r3(this, ((((unsigned int)A) >> 2) | (A << 30)), ((A << 19) | (((unsigned int)A) >> 13)), ((((unsigned int)A) >> 22) | (A << 10))) + fr_cryptohash_SHA2Core$MyKernel__r3(this, (A & B), (A & C), (B & C));
         H = G;
         G = F;
         F = E;
         E = D + T1;
         D = C;
         C = B;
         B = A;
         A = T1 + T2;
      }
      this->currentVal[7]  = this->currentVal[7] + H;
      this->currentVal[6]  = this->currentVal[6] + G;
      this->currentVal[5]  = this->currentVal[5] + F;
      this->currentVal[4]  = this->currentVal[4] + E;
      this->currentVal[3]  = this->currentVal[3] + D;
      this->currentVal[2]  = this->currentVal[2] + C;
      this->currentVal[1]  = this->currentVal[1] + B;
      this->currentVal[0]  = this->currentVal[0] + A;
      return;
   }
}

message length =    16 ->    0.02 MBytes/s
message length =    64 ->    0.04 MBytes/s
message length =   256 ->    0.05 MBytes/s
message length =  1024 ->    0.07 MBytes/s
message length =  8192 ->    0.08 MBytes/s
long messages          ->    0.08 MBytes/s

Process finished with exit code 0