The problem I'm seeing, though, is that select causes the most inefficient code. I'll give this other code a try and see what it comes up with. But the if statements appear to be best since it checks against what's already in the buffer. What it looks like is happening here with your new code is that it's causing writes to nonce regardless of whether or not the otherwise if statement is true. I think that's what is causing more cycles.
-_- So far, it looks like the if statements are most effective due to the lack of writes to and from vectors and nonce.
I just tested your new method. It DROPPED 10 ALUs! Congrats! The code is now faster.
Topic: Phatk2 Mod (Already seeing improvement!) - page 3. (Read 8054 times)
Yes, this is without any branching (similar to alternative 2) from my previous post except that I had W[3] wrong.
Basically the best would be to profile both and choose the faster one. With branching and without divergence, you have an additional clause (with divergence the penalty is worse as both "paths" would be serialized). However, without branching you introduce 7 dependent additions (can't pack them in two VLIW bundles as the result of the next addition depends on the previous one). I am not sure which would be faster.
BTW for the scalar case, you don't need that:
as direct comparison might be faster, especially with predication. E.g:
Unfortunately, this is not useful in the vector case. Of course you could try:
But that would generate much more inefficient code than that generated by using select().
The only problem is when you have more than one matching component pairs (v.sX and g.sX). For example v.s0==g.s0 and v.s3==g.s3. The version with branches would eventually have one of the two nonces written correctly in the output buffer (namely W[3].s3), the version with select() would have the wrong nonce written in the output buffer (W[3].s0+W[3].s3).
Basically the best would be to profile both and choose the faster one. With branching and without divergence, you have an additional clause (with divergence the penalty is worse as both "paths" would be serialized). However, without branching you introduce 7 dependent additions (can't pack them in two VLIW bundles as the result of the next addition depends on the previous one). I am not sure which would be faster.
BTW for the scalar case, you don't need that:
Code:
#else
v = select(W[3],(u)0,(v==g));
uint nonce = (v);
v = select(W[3],(u)0,(v==g));
uint nonce = (v);
as direct comparison might be faster, especially with predication. E.g:
Code:
nonce = (v==g) ? W[3] : 0;
Unfortunately, this is not useful in the vector case. Of course you could try:
Code:
nonce = (v.s0==g.s0) ? W[3].s0 : nonce;
nonce = (v.s1==g.s1) ? W[3].s1 : nonce;
...
nonce = (v.s1==g.s1) ? W[3].s1 : nonce;
...
But that would generate much more inefficient code than that generated by using select().
Quote
So, will having partial matches in a vector cause for any problems?
The only problem is when you have more than one matching component pairs (v.sX and g.sX). For example v.s0==g.s0 and v.s3==g.s3. The version with branches would eventually have one of the two nonces written correctly in the output buffer (namely W[3].s3), the version with select() would have the wrong nonce written in the output buffer (W[3].s0+W[3].s3).
I am compiling that with clcc and it builds successfully. It should work when you have VECTORS8 defined because you have:
and eq is defined to be of type u.
Yes, that's a valid point, but having nonce as a vector means you should be also increasing the output buffer (vector width) times. This in turn means you'd need (vector width) times larger device-host transfers. People with underclocked GPU memory and PCIe extenders won't be very happy about that
Yep, that's the purpose of replacing branches with select()
No, you don't need to do that. The result of (v==g) is a vector where each component is 0 if the corresponding v and g components are equal. E.g you have:
v = (uint8)(5,5,5,5,5,5,5,5);
g = (uint8)(1,2,3,4,5,6,7,8);
(v==g) would be (0,0,0,0,1,0,0,0)
This is still not useful as nonce is a scalar value. Then also (I noticed that later and corrected it) nonce should equal the matching vector element from W[3], not v or g.
Thus, this is the most straightforward solution:
eq = select(W[3],(u)0,(v==g))
What's the idea? eq is a vector, same width as W[3], v and g.
Let W[3] contain (0x10,0x20,0x30,0x40,0x50,0x60,0x70,0x80)
eq would contain (0,0,0,0,0x50,0,0,0)
since we need a scalar nonce, we just sum all the elements of eq and get 0+0+0+0+0x50+0+0+0 = 0x50.
Of course this would break if we have more than one match between v and g components and in that case the nonce would be wrong. The probability for this is low but it could happen. This is the worst case. Overall though, I think performance improvement due to branches elimination is worth the increased percentage of wrong shares. Also, a quick check on host could prevent the miner from submitting the wrong share. And as I said, this should occur rarely.
How about this instead?Code:
#ifdef VECTORS8
typedef uint8 u;
typedef uint8 u;
and eq is defined to be of type u.
Quote
So it takes that match and makes it nonce. However, it continues through the if statements as though it were looking for another match. So, either A) nonce needs to be increased in size to hold multiple equivalent vectors or B) the if statements needs to be stopped once a suitable nonce is found otherwise it will only serve to overwrite the first.
Yes, that's a valid point, but having nonce as a vector means you should be also increasing the output buffer (vector width) times. This in turn means you'd need (vector width) times larger device-host transfers. People with underclocked GPU memory and PCIe extenders won't be very happy about that
Quote
B) the if statements needs to be stopped once a suitable nonce is found otherwise it will only serve to overwrite the first.
Yep, that's the purpose of replacing branches with select()
Quote
And I think you may be misusing select there. You see, we need to pull apart v and g into it's separate parts before figuring out which parts are equal. That is, unless we can xor v and g, pull apart v and then write any vector = 0 from the equivalent g vector.
No, you don't need to do that. The result of (v==g) is a vector where each component is 0 if the corresponding v and g components are equal. E.g you have:
v = (uint8)(5,5,5,5,5,5,5,5);
g = (uint8)(1,2,3,4,5,6,7,8);
(v==g) would be (0,0,0,0,1,0,0,0)
This is still not useful as nonce is a scalar value. Then also (I noticed that later and corrected it) nonce should equal the matching vector element from W[3], not v or g.
Thus, this is the most straightforward solution:
eq = select(W[3],(u)0,(v==g))
What's the idea? eq is a vector, same width as W[3], v and g.
Let W[3] contain (0x10,0x20,0x30,0x40,0x50,0x60,0x70,0x80)
eq would contain (0,0,0,0,0x50,0,0,0)
since we need a scalar nonce, we just sum all the elements of eq and get 0+0+0+0+0x50+0+0+0 = 0x50.
Of course this would break if we have more than one match between v and g components and in that case the nonce would be wrong. The probability for this is low but it could happen. This is the worst case. Overall though, I think performance improvement due to branches elimination is worth the increased percentage of wrong shares. Also, a quick check on host could prevent the miner from submitting the wrong share. And as I said, this should occur rarely.
#ifdef VECTORS8
v = select(W[3],(u)0,(v==g));
uint nonce = (v.s0+v.s1+v.s2+v.s3+v.s4+v.s5+v.s6+v.s7);
#elif defined VECTORS4
v = select(W[3],(u)0,(v==g));
uint nonce = (v.s0+v.s1+v.s2+v.s3);
#elif defined VECTORS
v = select(W[3],(u)0,(v==g));
uint nonce = (v.s0+v.s1);
#else
v = select(W[3],(u)0,(v==g));
uint nonce = (v);
#endif
So, will having partial matches in a vector cause for any problems?
Also, using select is less efficient than using if statements. Any other method? And, as I suggested earlier, you might want the kernel analyzer as it will tell you the expected amount of cycles and the like. I'm not telling you that you have to, but it really does come in handy.
I am compiling that with clcc and it builds successfully. It should work when you have VECTORS8 defined because you have:
and eq is defined to be of type u.
Yes, that's a valid point, but having nonce as a vector means you should be also increasing the output buffer (vector width) times. This in turn means you'd need (vector width) times larger device-host transfers. People with underclocked GPU memory and PCIe extenders won't be very happy about that
Yep, that's the purpose of replacing branches with select()
No, you don't need to do that. The result of (v==g) is a vector where each component is 0 if the corresponding v and g components are equal. E.g you have:
v = (uint8)(5,5,5,5,5,5,5,5);
g = (uint8)(1,2,3,4,5,6,7,8);
(v==g) would be (0,0,0,0,1,0,0,0)
This is still not useful as nonce is a scalar value. Then also (I noticed that later and corrected it) nonce should equal the matching vector element from W[3], not v or g.
Thus, this is the most straightforward solution:
eq = select(W[3],(u)0,(v==g))
What's the idea? eq is a vector, same width as W[3], v and g.
Let W[3] contain (0x10,0x20,0x30,0x40,0x50,0x60,0x70,0x80)
eq would contain (0,0,0,0,0x50,0,0,0)
since we need a scalar nonce, we just sum all the elements of eq and get 0+0+0+0+0x50+0+0+0 = 0x50.
Of course this would break if we have more than one match between v and g components and in that case the nonce would be wrong. The probability for this is low but it could happen. This is the worst case. Overall though, I think performance improvement due to branches elimination is worth the increased percentage of wrong shares. Also, a quick check on host could prevent the miner from submitting the wrong share. And as I said, this should occur rarely.
Code:
#ifdef VECTORS8
typedef uint8 u;
typedef uint8 u;
and eq is defined to be of type u.
Quote
So it takes that match and makes it nonce. However, it continues through the if statements as though it were looking for another match. So, either A) nonce needs to be increased in size to hold multiple equivalent vectors or B) the if statements needs to be stopped once a suitable nonce is found otherwise it will only serve to overwrite the first.
Yes, that's a valid point, but having nonce as a vector means you should be also increasing the output buffer (vector width) times. This in turn means you'd need (vector width) times larger device-host transfers. People with underclocked GPU memory and PCIe extenders won't be very happy about that
Quote
B) the if statements needs to be stopped once a suitable nonce is found otherwise it will only serve to overwrite the first.
Yep, that's the purpose of replacing branches with select()
Quote
And I think you may be misusing select there. You see, we need to pull apart v and g into it's separate parts before figuring out which parts are equal. That is, unless we can xor v and g, pull apart v and then write any vector = 0 from the equivalent g vector.
No, you don't need to do that. The result of (v==g) is a vector where each component is 0 if the corresponding v and g components are equal. E.g you have:
v = (uint8)(5,5,5,5,5,5,5,5);
g = (uint8)(1,2,3,4,5,6,7,8);
(v==g) would be (0,0,0,0,1,0,0,0)
This is still not useful as nonce is a scalar value. Then also (I noticed that later and corrected it) nonce should equal the matching vector element from W[3], not v or g.
Thus, this is the most straightforward solution:
eq = select(W[3],(u)0,(v==g))
What's the idea? eq is a vector, same width as W[3], v and g.
Let W[3] contain (0x10,0x20,0x30,0x40,0x50,0x60,0x70,0x80)
eq would contain (0,0,0,0,0x50,0,0,0)
since we need a scalar nonce, we just sum all the elements of eq and get 0+0+0+0+0x50+0+0+0 = 0x50.
Of course this would break if we have more than one match between v and g components and in that case the nonce would be wrong. The probability for this is low but it could happen. This is the worst case. Overall though, I think performance improvement due to branches elimination is worth the increased percentage of wrong shares. Also, a quick check on host could prevent the miner from submitting the wrong share. And as I said, this should occur rarely.
OK, try this way:
That still won't work because eq must be of type uint8 to hold the data from v==g. Why not download the AMD APP KernelAnalyzer and copy-paste this code into it with the compile options of: -DOUTPUT_SIZE=0x100 -DOUTPUT_MASK=0xFF -DBFI_INT -DBITALIGN -DWORKSIZE=64 -DVECTORS8Code:
#ifdef VECTORS8
if (any(v==g))
{
u eq = select(W[3],(u)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
if (any(v==g))
{
u eq = select(W[3],(u)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
That way, you can see where the errors are firsthand. Now, the problem that I've found with the original code in general is that it makes the assumption that only one pair of vectors are going to match no matter what. So it takes that match and makes it nonce. However, it continues through the if statements as though it were looking for another match. So, either A) nonce needs to be increased in size to hold multiple equivalent vectors or B) the if statements needs to be stopped once a suitable nonce is found otherwise it will only serve to overwrite the first.
I could really use Phateus here to help me rewrite this portion to hold to these constraints.
And I think you may be misusing select there. You see, we need to pull apart v and g into it's separate parts before figuring out which parts are equal. That is, unless we can xor v and g, pull apart v and then write any vector = 0 from the equivalent g vector. However, that doesn't answer what to do in case of multiple nonce.
OK, try this way:
Code:
#ifdef VECTORS8
if (any(v==g))
{
u eq = select(W[3],(u)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
if (any(v==g))
{
u eq = select(W[3],(u)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
I'd advise you to change this:
To this:
A branchless version would be:
It incurs a penalty of several more ALU ops which might be acceptable or might not as compared to the one-branch version and this needs to be profiled.
Anyway, having 8 branches is a bad idea, even without divergence, this introduces at least 8 clauses and clause latency is ~40 cycles on VLIW hardware. Should be better on GCN though.
There's something wrong with the alternatives. It won't allow uint8 to be used with eq.Code:
#ifdef VECTORS8
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
if (v.s4 == g.s4)
{
nonce = W[3].s4;
}
if (v.s5 == g.s5)
{
nonce = W[3].s5;
}
if (v.s6 == g.s6)
{
nonce = W[3].s6;
}
if (v.s7 == g.s7)
{
nonce = W[3].s7;
}
#elif defined VECTORS4
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
if (v.s4 == g.s4)
{
nonce = W[3].s4;
}
if (v.s5 == g.s5)
{
nonce = W[3].s5;
}
if (v.s6 == g.s6)
{
nonce = W[3].s6;
}
if (v.s7 == g.s7)
{
nonce = W[3].s7;
}
#elif defined VECTORS4
To this:
Code:
#ifdef VECTORS8
uint8 eq=(v==g);
if (any(eq))
{
eq = select(g,(uint8)0,eq);
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
uint8 eq=(v==g);
if (any(eq))
{
eq = select(g,(uint8)0,eq);
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
A branchless version would be:
Code:
#ifdef VECTORS8
uint8 eq;
eq = select(g,(uint8)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
#elif defined VECTORS4
uint8 eq;
eq = select(g,(uint8)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
#elif defined VECTORS4
It incurs a penalty of several more ALU ops which might be acceptable or might not as compared to the one-branch version and this needs to be profiled.
Anyway, having 8 branches is a bad idea, even without divergence, this introduces at least 8 clauses and clause latency is ~40 cycles on VLIW hardware. Should be better on GCN though.
Another thing I would like to add is that there is a problem with efficiency of the nonce code. If nonce is a uint, there's only enough room for a single vector from W[3]. So, if we're running through the entire v==g array and happen upon v.s0==g.s0 right off the bat, it's still going through the rest of the if statements to see if it needs to replace that w[3] value in nonce yet again. My take on it is to ether stop there or make nonce large enough to hold any multiple nonce values and then introduce them into the next round as the full-sized vectors. -_- Too much editing of the kernel to add in efficiency. I want to leave this up to someone more experienced with programming it if someone would be so willing.
List of things to do:
Increase efficiency of nonce at end.
Find register spill in VECTORS8 processing and put a stop to it.
Verify GOFFSET is working properly.
Figure out why the heck moving around the P1 + P2 + P4 etc. increases ALUs when the preprocessor should be able to optimize the simple addition problems for speed.
Achieve at least a 1/15th speed increase by code optimization.
It's going to take a few days in the least to get the nonce working properly with the GOFFSET option. Since we've avoided using certain variables, it makes it difficult to weave it in. The same applies since we're skipping straight to W[3]. I don't want the new code to take up more ALUs though. It's a tough situation.
oh, and I was just going to test this kernel too to see what the hell goffset did for hashes
Yeah interested as well. Can we expect like a massive performance increase from 440 mhash/s for a 5870 to 700 mhash/s
? Guess not but let us know !
oh, and I was just going to test this kernel too to see what the hell goffset did for hashes
Yeah interested as well. Can we expect like a massive performance increase from 440 mhash/s for a 5870 to 700 mhash/s
? Guess not but let us know !
I'll try changing this for VECTORS8, 4 and 2 as there should be similar effects if it does as you say. I'm also working on getting that darn GOFFSET put into the init file. Not fun.
I'd advise you to change this:
To this:
A branchless version would be:
It incurs a penalty of several more ALU ops which might be acceptable or might not as compared to the one-branch version and this needs to be profiled.
Anyway, having 8 branches is a bad idea, even without divergence, this introduces at least 8 clauses and clause latency is ~40 cycles on VLIW hardware. Should be better on GCN though.
Code:
#ifdef VECTORS8
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
if (v.s4 == g.s4)
{
nonce = W[3].s4;
}
if (v.s5 == g.s5)
{
nonce = W[3].s5;
}
if (v.s6 == g.s6)
{
nonce = W[3].s6;
}
if (v.s7 == g.s7)
{
nonce = W[3].s7;
}
#elif defined VECTORS4
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
if (v.s4 == g.s4)
{
nonce = W[3].s4;
}
if (v.s5 == g.s5)
{
nonce = W[3].s5;
}
if (v.s6 == g.s6)
{
nonce = W[3].s6;
}
if (v.s7 == g.s7)
{
nonce = W[3].s7;
}
#elif defined VECTORS4
To this:
Code:
#ifdef VECTORS8
uint8 eq=(v==g);
if (any(eq))
{
eq = select(g,(uint8)0,eq);
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
uint8 eq=(v==g);
if (any(eq))
{
eq = select(g,(uint8)0,eq);
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
}
#elif defined VECTORS4
A branchless version would be:
Code:
#ifdef VECTORS8
uint8 eq;
eq = select(g,(uint8)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
#elif defined VECTORS4
uint8 eq;
eq = select(g,(uint8)0,(v==g));
nonce = (eq.s0+eq.s1+eq.s2+eq.s3+eq.s4+eq.s5+eq.s6+eq.s7);
#elif defined VECTORS4
It incurs a penalty of several more ALU ops which might be acceptable or might not as compared to the one-branch version and this needs to be profiled.
Anyway, having 8 branches is a bad idea, even without divergence, this introduces at least 8 clauses and clause latency is ~40 cycles on VLIW hardware. Should be better on GCN though.
oh, and I was just going to test this kernel too to see what the hell goffset did for hashes
Gah...I need to toss in an init file...crud. I'll mess with it later.
Hey everyone! I've decided to add Diapolo's GOFFSET option to Phatk2. In some cases, it works better, in others, it doesn't. Enable it to find out via GOFFSET!
This kernel is for Phoenix 2.0.0.
Currently, I test kernels on an ATI HD5450 graphics card. If anyone wants to help further my OpenCL expertise, I happily welcome donations to 12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq. They'll be put towards a new GPU or two (VLIW and GCN) and programming guides so I can compile optimum code. I already have the ideas, now I just need to get them out in the open. I'll also be helping Diapolo test his kernel and donate to help him out. (A day's worth of hashing each week should be enough)
This kernel is for Phoenix 2.0.0.
Currently, I test kernels on an ATI HD5450 graphics card. If anyone wants to help further my OpenCL expertise, I happily welcome donations to 12HWUSguWXRCQKfkPeJygVR1ex5wbg3hAq. They'll be put towards a new GPU or two (VLIW and GCN) and programming guides so I can compile optimum code. I already have the ideas, now I just need to get them out in the open. I'll also be helping Diapolo test his kernel and donate to help him out. (A day's worth of hashing each week should be enough)
Code:
// This file is in the public domain
#ifdef VECTORS8
typedef uint8 u;
#elif defined VECTORS4
typedef uint4 u;
#elif defined VECTORS
typedef uint2 u;
#else
typedef uint u;
#endif
__constant uint K[64] = {
0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2
};
__constant uint ConstW[128] = {
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x80000000U, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000280U,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x80000000U, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000100U,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000
};
__constant uint H[8] = {
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
};
#ifdef BITALIGN
#pragma OPENCL EXTENSION cl_amd_media_ops : enable
#define rot(x, y) amd_bitalign(x, x, (uint)(32 - y))
#else
#define rot(x, y) rotate(x, (uint)y)
#endif
// Some AMD devices have the BFI_INT opcode, which behaves exactly like the
// SHA-256 Ch function, but provides it in exactly one instruction. If
// detected, use it for Ch. Otherwise, use bitselect() for Ch.
#ifdef BFI_INT
// Well, slight problem... It turns out BFI_INT isn't actually exposed to
// OpenCL (or CAL IL for that matter) in any way. However, there is
// a similar instruction, BYTE_ALIGN_INT, which is exposed to OpenCL via
// amd_bytealign, takes the same inputs, and provides the same output.
// We can use that as a placeholder for BFI_INT and have the application
// patch it after compilation.
// This is the BFI_INT function
#define Ch(x, y, z) amd_bytealign(x,y,z)
// Ma can also be implemented in terms of BFI_INT...
#define Ma(z, x, y) amd_bytealign(z^x,y,x)
#else
#define Ch(x, y, z) bitselect(z,y,x)
#define Ma(x, y, z) bitselect(x,y,(z^x))
#endif
//Various intermediate calculations for each SHA round
#define s0(n) (S0(Vals[(0 + 128 - (n)) % 8]))
#define S0(n) (rot(n, 30u)^rot(n, 19u)^rot(n,10u))
#define s1(n) (S1(Vals[(4 + 128 - (n)) % 8]))
#define S1(n) (rot(n, 26u)^rot(n, 21u)^rot(n, 7u))
#define ch(n) Ch(Vals[(4 + 128 - (n)) % 8],Vals[(5 + 128 - (n)) % 8],Vals[(6 + 128 - (n)) % 8])
#define maj(n) Ma(Vals[(1 + 128 - (n)) % 8],Vals[(2 + 128 - (n)) % 8],Vals[(0 + 128 - (n)) % 8])
//t1 calc when W is already calculated
#define t1(n) K[(n) % 64] + Vals[(7 + 128 - (n)) % 8] + W[(n)] + s1(n) + ch(n)
//t1 calc which calculates W
#define t1W(n) K[(n) % 64] + Vals[(7 + 128 - (n)) % 8] + W(n) + s1(n) + ch(n)
//Used for constant W Values (the compiler optimizes out zeros)
#define t1C(n) (K[(n) % 64]+ ConstW[(n)]) + Vals[(7 + 128 - (n)) % 8] + s1(n) + ch(n)
//t2 Calc
#define t2(n) maj(n) + s0(n)
#define rotC(x,n) (x<> (32-n))
//W calculation used for SHA round
#define W(n) (W[n] = P4(n) + P3(n) + P2(n) + P1(n))
//Partial W calculations (used for the begining where only some values are nonzero)
#define P1(n) ((rot(W[(n)-2],15u)^rot(W[(n)-2],13u)^((W[(n)-2])>>10U)))
#define P2(n) ((rot(W[(n)-15],25u)^rot(W[(n)-15],14u)^((W[(n)-15])>>3U)))
#define p1(x) ((rot(x,15u)^rot(x,13u)^((x)>>10U)))
#define p2(x) ((rot(x,25u)^rot(x,14u)^((x)>>3U)))
#define P3(n) W[n-7]
#define P4(n) W[n-16]
//Partial Calcs for constant W values
#define P1C(n) ((rotC(ConstW[(n)-2],15)^rotC(ConstW[(n)-2],13)^((ConstW[(n)-2])>>10U)))
#define P2C(n) ((rotC(ConstW[(n)-15],25)^rotC(ConstW[(n)-15],14)^((ConstW[(n)-15])>>3U)))
#define P3C(x) ConstW[x-7]
#define P4C(x) ConstW[x-16]
//SHA round with built in W calc
#define sharoundW(n) Barrier1(n); Vals[(3 + 128 - (n)) % 8] += t1W(n); Vals[(7 + 128 - (n)) % 8] = t1W(n) + t2(n);
//SHA round without W calc
#define sharound(n) Barrier2(n); Vals[(3 + 128 - (n)) % 8] += t1(n); Vals[(7 + 128 - (n)) % 8] = t1(n) + t2(n);
//SHA round for constant W values
#define sharoundC(n) Barrier2(n); Vals[(3 + 128 - (n)) % 8] += t1C(n); Vals[(7 + 128 - (n)) % 8] = t1C(n) + t2(n);
//The compiler is stupid... I put this in there only to stop the compiler from (de)optimizing the order
#define Barrier1(n) t1 = t1C((n+1))
#define Barrier2(n) t1 = t1C((n))
__kernel
//removed this to allow detection of invalid work size
//__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
void search( const uint state0, const uint state1, const uint state2, const uint state3,
const uint state4, const uint state5, const uint state6, const uint state7,
const uint B1, const uint C1, const uint D1,
const uint F1, const uint G1, const uint H1,
#ifndef GOFFSET
const u base,
#endif
const uint W16, const uint W17,
const uint PreVal4, const uint PreVal0,
const uint PreW31, const uint PreW32,
const uint PreW19, const uint PreW20,
__global uint * output)
{
u W[124];
u Vals[8];
//Dummy Variable to prevent compiler from reordering between rounds
u t1;
//Vals[0]=state0;
Vals[1]=B1;
Vals[2]=C1;
Vals[3]=D1;
//Vals[4]=PreVal4;
Vals[5]=F1;
Vals[6]=G1;
Vals[7]=H1;
W[16] = W16;
W[17] = W17;
#ifdef VECTORS8
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0) << 3) + (u)(0, 1, 2, 3, 4, 5, 6, 7);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
W[3] = base + (uint)(get_local_id(0)) * 8u + (uint)(get_group_id(0)) * (WORKSIZE * 8u);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
//Since only the 2 LSB is opposite between the nonces, we can save an instruction by flipping the 4 bits in W18 rather than the 1 bit in W3
W[18] = PreW20 + (u){r, r ^ 0x2004000U, r ^ 0x4008000U, r ^ 0x600C000U, r ^ 0x8010000U, r ^ 0xA050000U, r ^ 0xC090000U, r ^ 0xE0D0000U};
#endif
#elif defined VECTORS4
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0) << 2) + (u)(0, 1, 2, 3);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
//Less dependencies to get both the local id and group id and then add them
W[3] = base + (uint)(get_local_id(0)) * 4u + (uint)(get_group_id(0)) * (WORKSIZE * 4u);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
//Since only the 2 LSB is opposite between the nonces, we can save an instruction by flipping the 4 bits in W18 rather than the 1 bit in W3
W[18] = PreW20 + (u){r, r ^ 0x2004000U, r ^ 0x4008000U, r ^ 0x600C000U};
#endif
#elif defined VECTORS
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0) << 1) + (u)(0, 1,);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
W[3] = base + (uint)(get_local_id(0)) * 2u + (uint)(get_group_id(0)) * (WORKSIZE * 2u);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
W[18] = PreW20 + (u){r, r ^ 0x2004000U};
#endif
#else
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0));
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
W[3] = base + get_local_id(0) + get_group_id(0) * (WORKSIZE);
u r = rot(W[3],25u)^rot(W[3],14u)^((W[3])>>3U);
W[18] = PreW20 + r;
#endif
#endif
//the order of the W calcs and Rounds is like this because the compiler needs help finding how to order the instructions
Vals[4] = PreVal4 + W[3];
Vals[0] = PreVal0 + W[3];
sharoundC(4);
W[19] = PreW19 + W[3];
sharoundC(5);
W[20] = P4C(20) + P1(20);
sharoundC(6);
W[21] = P1(21);
sharoundC(7);
W[22] = P3C(22) + P1(22);
sharoundC(8);
W[23] = W[16] + P1(23);
sharoundC(9);
W[24] = W[17] + P1(24);
sharoundC(10);
W[25] = P1(25) + P3(25);
W[26] = P1(26) + P3(26);
sharoundC(11);
W[27] = P1(27) + P3(27);
W[28] = P1(28) + P3(28);
sharoundC(12);
W[29] = P1(29) + P3(29);
sharoundC(13);
W[30] = P1(30) + P2C(30) + P3(30);
W[31] = PreW31 + (P1(31) + P3(31));
sharoundC(14);
W[32] = PreW32 + (P1(32) + P3(32));
sharoundC(15);
sharound(16);
sharound(17);
sharound(18);
sharound(19);
sharound(20);
sharound(21);
sharound(22);
sharound(23);
sharound(24);
sharound(25);
sharound(26);
sharound(27);
sharound(28);
sharound(29);
sharound(30);
sharound(31);
sharound(32);
sharoundW(33);
sharoundW(34);
sharoundW(35);
sharoundW(36);
sharoundW(37);
sharoundW(38);
sharoundW(39);
sharoundW(40);
sharoundW(41);
sharoundW(42);
sharoundW(43);
sharoundW(44);
sharoundW(45);
sharoundW(46);
sharoundW(47);
sharoundW(48);
sharoundW(49);
sharoundW(50);
sharoundW(51);
sharoundW(52);
sharoundW(53);
sharoundW(54);
sharoundW(55);
sharoundW(56);
sharoundW(57);
sharoundW(58);
sharoundW(59);
sharoundW(60);
sharoundW(61);
sharoundW(62);
sharoundW(63);
W[64]=state0+Vals[0];
W[65]=state1+Vals[1];
W[66]=state2+Vals[2];
W[67]=state3+Vals[3];
W[68]=state4+Vals[4];
W[69]=state5+Vals[5];
W[70]=state6+Vals[6];
W[71]=state7+Vals[7];
const u Temp = (0xb0edbdd0U + K[0]) + W[64];
Vals[0]=H[0];
Vals[1]=H[1];
Vals[2]=H[2];
Vals[3]=0xa54ff53aU + Temp;
Vals[4]=H[4];
Vals[5]=H[5];
Vals[6]=H[6];
Vals[7]=0x08909ae5U + Temp;
#define P124(n) P2(n) + P1(n) + P4(n)
W[80] = P2(80) + P4(80);
sharound(65);
W[81] = P1C(81) + P2(81) + P4(81);
sharound(66);
W[82] = P124(82);
sharound(67);
W[83] = P124(83);
sharound(68);
W[84] = P124(84);
sharound(69);
W[85] = P124(85);
sharound(70);
W[86] = P4(86) + P3C(86) + P2(86) + P1(86);
sharound(71);
W[87] = P4(87) + P3(87) + P2C(87) + P1(87);
sharoundC(72);
W[88] = P1(88) + P4C(88) + P3(88);
sharoundC(73);
W[89] = P3(89) + P1(89);
sharoundC(74);
W[90] = P3(90) + P1(90);
sharoundC(75);
W[91] = P3(91) + P1(91);
sharoundC(76);
W[92] = P3(92) + P1(92);
sharoundC(77);
W[93] = P1(93) + P3(93);
W[94] = P3(94) + P2C(94) + P1(94);
sharoundC(78);
W[95] = P4C(95) + P3(95) + P2(95) + P1(95);
sharoundC(79);
sharound(80);
sharound(81);
sharound(82);
sharound(83);
sharound(84);
sharound(85);
sharound(86);
sharound(87);
sharound(88);
sharound(89);
sharound(90);
sharound(91);
sharound(92);
sharound(93);
sharound(94);
sharound(95);
sharoundW(96);
sharoundW(97);
sharoundW(98);
sharoundW(99);
sharoundW(100);
sharoundW(101);
sharoundW(102);
sharoundW(103);
sharoundW(104);
sharoundW(105);
sharoundW(106);
sharoundW(107);
sharoundW(108);
sharoundW(109);
sharoundW(110);
sharoundW(111);
sharoundW(112);
sharoundW(113);
sharoundW(114);
sharoundW(115);
sharoundW(116);
sharoundW(117);
sharoundW(118);
sharoundW(119);
sharoundW(120);
sharoundW(121);
sharoundW(122);
u v = W[117] + W[108] + Vals[3] + Vals[7] + P2(124) + P1(124) + Ch((Vals[0] + Vals[4]) + (K[59] + W(123)) + s1(123)+ ch(123),Vals[1],Vals[2]);
u g = -(K[60] + H[7]) - S1((Vals[0] + Vals[4]) + (K[59] + W(123)) + s1(123)+ ch(123));
uint nonce = 0;
#ifdef VECTORS8
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
if (v.s4 == g.s4)
{
nonce = W[3].s4;
}
if (v.s5 == g.s5)
{
nonce = W[3].s5;
}
if (v.s6 == g.s6)
{
nonce = W[3].s6;
}
if (v.s7 == g.s7)
{
nonce = W[3].s7;
}
#elif defined VECTORS4
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
#elif defined VECTORS
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
#else
if (v == g)
{
nonce = W[3];
}
#endif
if(nonce)
{
//Faster to shift the nonce by 2 due to 4-DWORD addressing and does not add more collisions
output[WORKSIZE] = nonce;
output[get_local_id(0)] = nonce;
}
}
As you may be able to tell, I've started adding VECTORS8 to the code as well, but I am having difficulty keeping it from spilling over into memory. If someone could assist me with this, I would appreciate it. #ifdef VECTORS8
typedef uint8 u;
#elif defined VECTORS4
typedef uint4 u;
#elif defined VECTORS
typedef uint2 u;
#else
typedef uint u;
#endif
__constant uint K[64] = {
0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2
};
__constant uint ConstW[128] = {
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x80000000U, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000280U,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x80000000U, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000100U,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000,
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000
};
__constant uint H[8] = {
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
};
#ifdef BITALIGN
#pragma OPENCL EXTENSION cl_amd_media_ops : enable
#define rot(x, y) amd_bitalign(x, x, (uint)(32 - y))
#else
#define rot(x, y) rotate(x, (uint)y)
#endif
// Some AMD devices have the BFI_INT opcode, which behaves exactly like the
// SHA-256 Ch function, but provides it in exactly one instruction. If
// detected, use it for Ch. Otherwise, use bitselect() for Ch.
#ifdef BFI_INT
// Well, slight problem... It turns out BFI_INT isn't actually exposed to
// OpenCL (or CAL IL for that matter) in any way. However, there is
// a similar instruction, BYTE_ALIGN_INT, which is exposed to OpenCL via
// amd_bytealign, takes the same inputs, and provides the same output.
// We can use that as a placeholder for BFI_INT and have the application
// patch it after compilation.
// This is the BFI_INT function
#define Ch(x, y, z) amd_bytealign(x,y,z)
// Ma can also be implemented in terms of BFI_INT...
#define Ma(z, x, y) amd_bytealign(z^x,y,x)
#else
#define Ch(x, y, z) bitselect(z,y,x)
#define Ma(x, y, z) bitselect(x,y,(z^x))
#endif
//Various intermediate calculations for each SHA round
#define s0(n) (S0(Vals[(0 + 128 - (n)) % 8]))
#define S0(n) (rot(n, 30u)^rot(n, 19u)^rot(n,10u))
#define s1(n) (S1(Vals[(4 + 128 - (n)) % 8]))
#define S1(n) (rot(n, 26u)^rot(n, 21u)^rot(n, 7u))
#define ch(n) Ch(Vals[(4 + 128 - (n)) % 8],Vals[(5 + 128 - (n)) % 8],Vals[(6 + 128 - (n)) % 8])
#define maj(n) Ma(Vals[(1 + 128 - (n)) % 8],Vals[(2 + 128 - (n)) % 8],Vals[(0 + 128 - (n)) % 8])
//t1 calc when W is already calculated
#define t1(n) K[(n) % 64] + Vals[(7 + 128 - (n)) % 8] + W[(n)] + s1(n) + ch(n)
//t1 calc which calculates W
#define t1W(n) K[(n) % 64] + Vals[(7 + 128 - (n)) % 8] + W(n) + s1(n) + ch(n)
//Used for constant W Values (the compiler optimizes out zeros)
#define t1C(n) (K[(n) % 64]+ ConstW[(n)]) + Vals[(7 + 128 - (n)) % 8] + s1(n) + ch(n)
//t2 Calc
#define t2(n) maj(n) + s0(n)
#define rotC(x,n) (x<
//W calculation used for SHA round
#define W(n) (W[n] = P4(n) + P3(n) + P2(n) + P1(n))
//Partial W calculations (used for the begining where only some values are nonzero)
#define P1(n) ((rot(W[(n)-2],15u)^rot(W[(n)-2],13u)^((W[(n)-2])>>10U)))
#define P2(n) ((rot(W[(n)-15],25u)^rot(W[(n)-15],14u)^((W[(n)-15])>>3U)))
#define p1(x) ((rot(x,15u)^rot(x,13u)^((x)>>10U)))
#define p2(x) ((rot(x,25u)^rot(x,14u)^((x)>>3U)))
#define P3(n) W[n-7]
#define P4(n) W[n-16]
//Partial Calcs for constant W values
#define P1C(n) ((rotC(ConstW[(n)-2],15)^rotC(ConstW[(n)-2],13)^((ConstW[(n)-2])>>10U)))
#define P2C(n) ((rotC(ConstW[(n)-15],25)^rotC(ConstW[(n)-15],14)^((ConstW[(n)-15])>>3U)))
#define P3C(x) ConstW[x-7]
#define P4C(x) ConstW[x-16]
//SHA round with built in W calc
#define sharoundW(n) Barrier1(n); Vals[(3 + 128 - (n)) % 8] += t1W(n); Vals[(7 + 128 - (n)) % 8] = t1W(n) + t2(n);
//SHA round without W calc
#define sharound(n) Barrier2(n); Vals[(3 + 128 - (n)) % 8] += t1(n); Vals[(7 + 128 - (n)) % 8] = t1(n) + t2(n);
//SHA round for constant W values
#define sharoundC(n) Barrier2(n); Vals[(3 + 128 - (n)) % 8] += t1C(n); Vals[(7 + 128 - (n)) % 8] = t1C(n) + t2(n);
//The compiler is stupid... I put this in there only to stop the compiler from (de)optimizing the order
#define Barrier1(n) t1 = t1C((n+1))
#define Barrier2(n) t1 = t1C((n))
__kernel
//removed this to allow detection of invalid work size
//__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
void search( const uint state0, const uint state1, const uint state2, const uint state3,
const uint state4, const uint state5, const uint state6, const uint state7,
const uint B1, const uint C1, const uint D1,
const uint F1, const uint G1, const uint H1,
#ifndef GOFFSET
const u base,
#endif
const uint W16, const uint W17,
const uint PreVal4, const uint PreVal0,
const uint PreW31, const uint PreW32,
const uint PreW19, const uint PreW20,
__global uint * output)
{
u W[124];
u Vals[8];
//Dummy Variable to prevent compiler from reordering between rounds
u t1;
//Vals[0]=state0;
Vals[1]=B1;
Vals[2]=C1;
Vals[3]=D1;
//Vals[4]=PreVal4;
Vals[5]=F1;
Vals[6]=G1;
Vals[7]=H1;
W[16] = W16;
W[17] = W17;
#ifdef VECTORS8
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0) << 3) + (u)(0, 1, 2, 3, 4, 5, 6, 7);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
W[3] = base + (uint)(get_local_id(0)) * 8u + (uint)(get_group_id(0)) * (WORKSIZE * 8u);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
//Since only the 2 LSB is opposite between the nonces, we can save an instruction by flipping the 4 bits in W18 rather than the 1 bit in W3
W[18] = PreW20 + (u){r, r ^ 0x2004000U, r ^ 0x4008000U, r ^ 0x600C000U, r ^ 0x8010000U, r ^ 0xA050000U, r ^ 0xC090000U, r ^ 0xE0D0000U};
#endif
#elif defined VECTORS4
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0) << 2) + (u)(0, 1, 2, 3);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
//Less dependencies to get both the local id and group id and then add them
W[3] = base + (uint)(get_local_id(0)) * 4u + (uint)(get_group_id(0)) * (WORKSIZE * 4u);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
//Since only the 2 LSB is opposite between the nonces, we can save an instruction by flipping the 4 bits in W18 rather than the 1 bit in W3
W[18] = PreW20 + (u){r, r ^ 0x2004000U, r ^ 0x4008000U, r ^ 0x600C000U};
#endif
#elif defined VECTORS
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0) << 1) + (u)(0, 1,);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
W[3] = base + (uint)(get_local_id(0)) * 2u + (uint)(get_group_id(0)) * (WORKSIZE * 2u);
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
W[18] = PreW20 + (u){r, r ^ 0x2004000U};
#endif
#else
#ifdef GOFFSET
W[3] = ((uint)get_global_id(0));
uint r = rot(W[3].s0,25u)^rot(W[3].s0,14u)^((W[3].s0)>>3U);
#else
W[3] = base + get_local_id(0) + get_group_id(0) * (WORKSIZE);
u r = rot(W[3],25u)^rot(W[3],14u)^((W[3])>>3U);
W[18] = PreW20 + r;
#endif
#endif
//the order of the W calcs and Rounds is like this because the compiler needs help finding how to order the instructions
Vals[4] = PreVal4 + W[3];
Vals[0] = PreVal0 + W[3];
sharoundC(4);
W[19] = PreW19 + W[3];
sharoundC(5);
W[20] = P4C(20) + P1(20);
sharoundC(6);
W[21] = P1(21);
sharoundC(7);
W[22] = P3C(22) + P1(22);
sharoundC(8);
W[23] = W[16] + P1(23);
sharoundC(9);
W[24] = W[17] + P1(24);
sharoundC(10);
W[25] = P1(25) + P3(25);
W[26] = P1(26) + P3(26);
sharoundC(11);
W[27] = P1(27) + P3(27);
W[28] = P1(28) + P3(28);
sharoundC(12);
W[29] = P1(29) + P3(29);
sharoundC(13);
W[30] = P1(30) + P2C(30) + P3(30);
W[31] = PreW31 + (P1(31) + P3(31));
sharoundC(14);
W[32] = PreW32 + (P1(32) + P3(32));
sharoundC(15);
sharound(16);
sharound(17);
sharound(18);
sharound(19);
sharound(20);
sharound(21);
sharound(22);
sharound(23);
sharound(24);
sharound(25);
sharound(26);
sharound(27);
sharound(28);
sharound(29);
sharound(30);
sharound(31);
sharound(32);
sharoundW(33);
sharoundW(34);
sharoundW(35);
sharoundW(36);
sharoundW(37);
sharoundW(38);
sharoundW(39);
sharoundW(40);
sharoundW(41);
sharoundW(42);
sharoundW(43);
sharoundW(44);
sharoundW(45);
sharoundW(46);
sharoundW(47);
sharoundW(48);
sharoundW(49);
sharoundW(50);
sharoundW(51);
sharoundW(52);
sharoundW(53);
sharoundW(54);
sharoundW(55);
sharoundW(56);
sharoundW(57);
sharoundW(58);
sharoundW(59);
sharoundW(60);
sharoundW(61);
sharoundW(62);
sharoundW(63);
W[64]=state0+Vals[0];
W[65]=state1+Vals[1];
W[66]=state2+Vals[2];
W[67]=state3+Vals[3];
W[68]=state4+Vals[4];
W[69]=state5+Vals[5];
W[70]=state6+Vals[6];
W[71]=state7+Vals[7];
const u Temp = (0xb0edbdd0U + K[0]) + W[64];
Vals[0]=H[0];
Vals[1]=H[1];
Vals[2]=H[2];
Vals[3]=0xa54ff53aU + Temp;
Vals[4]=H[4];
Vals[5]=H[5];
Vals[6]=H[6];
Vals[7]=0x08909ae5U + Temp;
#define P124(n) P2(n) + P1(n) + P4(n)
W[80] = P2(80) + P4(80);
sharound(65);
W[81] = P1C(81) + P2(81) + P4(81);
sharound(66);
W[82] = P124(82);
sharound(67);
W[83] = P124(83);
sharound(68);
W[84] = P124(84);
sharound(69);
W[85] = P124(85);
sharound(70);
W[86] = P4(86) + P3C(86) + P2(86) + P1(86);
sharound(71);
W[87] = P4(87) + P3(87) + P2C(87) + P1(87);
sharoundC(72);
W[88] = P1(88) + P4C(88) + P3(88);
sharoundC(73);
W[89] = P3(89) + P1(89);
sharoundC(74);
W[90] = P3(90) + P1(90);
sharoundC(75);
W[91] = P3(91) + P1(91);
sharoundC(76);
W[92] = P3(92) + P1(92);
sharoundC(77);
W[93] = P1(93) + P3(93);
W[94] = P3(94) + P2C(94) + P1(94);
sharoundC(78);
W[95] = P4C(95) + P3(95) + P2(95) + P1(95);
sharoundC(79);
sharound(80);
sharound(81);
sharound(82);
sharound(83);
sharound(84);
sharound(85);
sharound(86);
sharound(87);
sharound(88);
sharound(89);
sharound(90);
sharound(91);
sharound(92);
sharound(93);
sharound(94);
sharound(95);
sharoundW(96);
sharoundW(97);
sharoundW(98);
sharoundW(99);
sharoundW(100);
sharoundW(101);
sharoundW(102);
sharoundW(103);
sharoundW(104);
sharoundW(105);
sharoundW(106);
sharoundW(107);
sharoundW(108);
sharoundW(109);
sharoundW(110);
sharoundW(111);
sharoundW(112);
sharoundW(113);
sharoundW(114);
sharoundW(115);
sharoundW(116);
sharoundW(117);
sharoundW(118);
sharoundW(119);
sharoundW(120);
sharoundW(121);
sharoundW(122);
u v = W[117] + W[108] + Vals[3] + Vals[7] + P2(124) + P1(124) + Ch((Vals[0] + Vals[4]) + (K[59] + W(123)) + s1(123)+ ch(123),Vals[1],Vals[2]);
u g = -(K[60] + H[7]) - S1((Vals[0] + Vals[4]) + (K[59] + W(123)) + s1(123)+ ch(123));
uint nonce = 0;
#ifdef VECTORS8
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
if (v.s4 == g.s4)
{
nonce = W[3].s4;
}
if (v.s5 == g.s5)
{
nonce = W[3].s5;
}
if (v.s6 == g.s6)
{
nonce = W[3].s6;
}
if (v.s7 == g.s7)
{
nonce = W[3].s7;
}
#elif defined VECTORS4
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
if (v.s2 == g.s2)
{
nonce = W[3].s2;
}
if (v.s3 == g.s3)
{
nonce = W[3].s3;
}
#elif defined VECTORS
if (v.s0 == g.s0)
{
nonce = W[3].s0;
}
if (v.s1 == g.s1)
{
nonce = W[3].s1;
}
#else
if (v == g)
{
nonce = W[3];
}
#endif
if(nonce)
{
//Faster to shift the nonce by 2 due to 4-DWORD addressing and does not add more collisions
output[WORKSIZE] = nonce;
output[get_local_id(0)] = nonce;
}
}
Jump to: