[Dev] NeoScrypt GPU Miner - Public Beta Test

slowhash

Well, I got woke up for an outage at work that is now solved, so I’m back at it for a few minutes while I get drowsy…

Tried out my existing settings from the 14.4 OCL on the 14.6 and 14.7, and my hash rate jumped from 255/245 to 299/294 for the 290/290x.

As mentioned earlier, indecipherable difference between the 14.6 and the 14.7.

Thanks insanid for posting those!!!

Alpha Wolf

Scratch what I said about sgminer and nvidia as its not producing any shares and some hardware errors. :-

insanid

Well, I got woke up for an outage at work that is now solved, so I’m back at it for a few minutes while I get drowsy…

Tried out my existing settings from the 14.4 OCL on the 14.6 and 14.7, and my hash rate jumped from 255/245 to 299/294 for the 290/290x.

As mentioned earlier, indecipherable difference between the 14.6 and the 14.7.

Thanks insanid for posting those!!!

No problem! I am glad to hear you have your 290s running near optimum hash rate with Wolf0’s kernel. Try these settings on the 290x to see if you can get higher hashrates closer to mine:

AMD r9 290x (Hynix):

w64 g2 (g2 is REALLY the key here and ONLY for 290x in this case) I14

Elpida cards may or may not see an increase in speed with these settings.

slowhash

That was one of the 3 settings that gave me the highest speeds, but it was number two, and gives me 297/284.

The best I get is from I13 g2 w96.

GPU 1 & 2 are 290x’s, GPU 3 & 4 are 290’s.

insanid

That was one of the 3 settings that gave me the highest speeds, but it was number two, and gives me 297/284.

The best I get is from I13 g2 w96.

GPU 1 & 2 are 290x’s, GPU 3 & 4 are 290’s.

I guess my 290x is special. It’s a Sapphire Vapor-X 290x with Hynix GDDR5, so that could be it also.

slowhash

Speaking of which, is that bitalign thing normal? I’ve never seen it before doing neoscrypt with SGminer…

ghostlander

Nice will try it very soon and report back, Thank you!
Edit: What miner versions have been tested? Be nice to know what all this works with and drivers.

I see most people use our customised CGminer, so I test my kernel with it. If you are going to put it into SGminer, replace WORKGROUPSIZE with WORKSIZE. Everything else should be the same.

Alpha Wolf

I see most people use our customised CGminer, so I test my kernel with it. If you are going to put it into SGminer, replace WORKGROUPSIZE with WORKSIZE. Everything else should be the same.

My tests showed 47kh/s with 3.7.7c and 3.7.8 with my unlocked 6950 -I 12 -w 64. Thanks

ghostlander

My tests showed 47kh/s with 3.7.7c and 3.7.8 with my unlocked 6950 -I 12 -w 64. Thanks

You may increase work size to 256. It helps these VLIW cards a little.

insanid

Speaking of which, is that bitalign thing normal? I’ve never seen it before doing neoscrypt with SGminer…

Yes it’s perfectly normal to see that in sgminer.

daimyo

Folks, might be a silly question but would neoscrypt work with onboard intel graphics ?

ghostlander

Folks, might be a silly question but would neoscrypt work with onboard intel graphics ?

Intel HD graphics of Sandy Briidge and newer supports OpenCL, so it will. Performance isn’t going to be impressive though.

daimyo

ghostlander, It sure isn`t. Just wanted to make sure its in the arsenal :) thanks

slowhash

Intel HD graphics of Sandy Briidge and newer supports OpenCL, so it will. Performance isn’t going to be impressive though.

Do the APU’s support OpenCL? My wife’s PC is due to get upgraded sometime soon, and that might effect what gets built.

cisahasa

Do the APU’s support OpenCL? My wife’s PC is due to get upgraded sometime soon, and that might effect what gets built.

found something like this:

so i guess it will work?
AMD Embedded GPU and APU Driver 14.201.1009 for Windows 7/8

Thia is the latest graphics driver for AMD’s Embedded GPU and APUs. Windows 7/8/8.1 32-bit and 64-bit supported.

ghostlander

Do the APU’s support OpenCL? My wife’s PC is due to get upgraded sometime soon, and that might effect what gets built.

Of course they do. Even old Llano with 400 pre-GCN units is a good thing to have. Kaveri with 512 GCN units is even better, though I still prefer 4 real cores of Llano to 4 CMT cores of Kaveri.

xIIImaL

neoscrypt_vliw.cl v3

HD6970: 19.5KH/s to 50.5KH/s and no HW errors!

Thanks to a much better global memory management with burst writes. Also tried a trick with copying a workgroup buffer to local memory and writing back asynchronously to global memory through async_work_group_copy(), but it didn’t work out (36KH/s and 15% of HW errors). A 2.5x speed increase is something to celebrate anyway. Anyone cares to donate towards drinks, whores and blackjack? ;)

How it use? Don’t work for me with cgminer 3.7.8. Cards in rig: 6950,6950,6870,5870,5870. Drivers 13.12, Windows 7 x64 sp1

Wellenreiter

What is the content of your .bat file? I suggest to rename your feathercoin.conf and start with a simple 'cgminer’command.

Then enter the url, username and password of your pool when prompted and try that. If the miner works write the config file to disk and start tuning from there.

If it doesn’t work, I leave it to the Wolf0 and Avehre to find a solution

ghostlander

How it use? Don’t work for me with cgminer 3.7.8. Cards in rig: 6950,6950,6870,5870,5870. Drivers 13.12, Windows 7 x64 sp1

Replace this in the kernel:

/* NeoScrypt core engine:
 * N = 128, r = 2, p = 1, salt = password */
__kernel __attribute__((vec_type_hint(uint4)))
__kernel __attribute__((reqd_work_group_size(WORKGROUPSIZE, 1, 1)))
__kernel void search(__global const uint4* restrict password,
  volatile __global uint* restrict output, __global uint16* globalcache,
  const uint target) {
    uint glbid = get_global_id(0);
    uint grpid = get_group_id(0);
    uint lsize = WORKGROUPSIZE;
    uint lclid = glbid & (WORKGROUPSIZE - 1);

    __global uint16 *G = &globalcache[(grpid * WORKGROUPSIZE) << 9];

with this:

/* NeoScrypt core engine:
 * N = 128, r = 2, p = 1, salt = password */
__kernel __attribute__((vec_type_hint(uint4)))
__kernel void search(__global const uint4* restrict password,
  volatile __global uint* restrict output, __global uint16* globalcache,
  const uint target) {
    uint glbid = get_global_id(0);
    uint grpid = get_group_id(0);
    uint lsize = get_local_size(0);
    uint lclid = glbid & (lsize - 1);

    __global uint16 *G = &globalcache[(grpid * lsize) << 9];

It makes the kernel very compatible at -1% of performance.

cisahasa

Replace this in the kernel:

/* NeoScrypt core engine:
 * N = 128, r = 2, p = 1, salt = password */
__kernel __attribute__((vec_type_hint(uint4)))
__kernel __attribute__((reqd_work_group_size(WORKGROUPSIZE, 1, 1)))
__kernel void search(__global const uint4* restrict password,
  volatile __global uint* restrict output, __global uint16* globalcache,
  const uint target) {
    uint glbid = get_global_id(0);
    uint grpid = get_group_id(0);
    uint lsize = WORKGROUPSIZE;
    uint lclid = glbid & (WORKGROUPSIZE - 1);

    __global uint16 *G = &globalcache[(grpid * WORKGROUPSIZE) << 9];

with this:

/* NeoScrypt core engine:
 * N = 128, r = 2, p = 1, salt = password */
__kernel __attribute__((vec_type_hint(uint4)))
__kernel void search(__global const uint4* restrict password,
  volatile __global uint* restrict output, __global uint16* globalcache,
  const uint target) {
    uint glbid = get_global_id(0);
    uint grpid = get_group_id(0);
    uint lsize = get_local_size(0);
    uint lclid = glbid & (lsize - 1);

    __global uint16 *G = &globalcache[(grpid * lsize) << 9];

It makes the kernel very compatible at -1% of performance.

does sgminer with wolf latest kernel work for 6xxx cards as it is? +using dll fix for drivers??

playground with kernels and miners is a little complicated now…

hard to keep on track