Quote from: jkister on November 16, 2014, 06:28:42 AM

Quote from: ghostlander on November 12, 2014, 05:30:11 AM

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

hmm - i replaced my neoscrypt.cl with yours then deleted my bin file. when sgminer goes to recreate it:

Quote

[01:20:47] Building binary neoscryptCaymangw128l8tc7000.bin
[01:20:47] Error -11: Building Program (clBuildProgram)
[01:20:47] "/tmp/OCL18375T14.cl", line 596: error: identifier "WORKGROUPSIZE" is undefined
__kernel __attribute__((reqd_work_group_size(WORKGROUPSIZE, 1, 1)))
^

1 error detected in the compilation of "/tmp/OCL18375T14.cl".

Frontend phase failed compilation.

any idea what i've done wrong? not sure if 'workgroupsize' has anything to do with 'worksize' but ive tried it at 128, 256, and undefined.

SGminer should define it as WORKSIZE. If even this fails, try a safe way.

Replace the following in the kernel:

Code:

/* 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:

Code:

/* 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's going to reduce performance by about 1% though.

In addition, don't try to set thread concurrency to anything other than power of 2, i.e. 1024, 2048, 4096... It will only produce HW errors otherwise.