CCminer(SP-MOD) Modded GPU kernels.

[ryen123]

Hey there,

Thanks for having a look at my CUDA kernel. I actually took the keccak and ROTL64 code from your ccminer fork (or was it someone else's fork..not sure) a few months ago, so I'm surprised it can even be faster .

Unfortunately I don't have a working Maxwell development environment to build and test, only Kepler.

The bug causing low hashrate on windows hope it can be solved.

[skunk]

Hey there,

Thanks for having a look at my CUDA kernel. I actually took the keccak and ROTL64 code from your ccminer fork (or was it someone else's fork..not sure) a few months ago, so I'm surprised it can even be faster .

Unfortunately I don't have a working Maxwell development environment to build and test, only Kepler.

i think making ethminer less cpu intensive (full cpu usage on a G3240 processor) would boost performance considerably as ccminer barely uses the cpu...

[sp_]

To improve the cpu usage you can try to add:

cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);

like this:

Code:

void run_ethash_search(
	uint32_t blocks,
	uint32_t threads,
	cudaStream_t stream,
	uint32_t* g_output,
	hash32_t const* g_header,
	hash128_t const* g_dag,
	uint64_t start_nonce,
	uint64_t target
)
{
	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
	ethash_search
#if __CUDA_ARCH__ >= SHUFFLE_MIN_VER
	ethash_search <<<blocks, threads, 0, stream >>>(g_output, g_header, g_dag, start_nonce, target);
#else
	ethash_search <<<blocks, threads, (sizeof(compute_hash_share) * threads) / THREADS_PER_HASH, stream>>>(g_output, g_header, g_dag, start_nonce, target);
#endif
}

[sp_]

Nice work SP , I wish to test it , but can't compile unfortunately...

copy cuda_helper.h from sp's ccminer into libethash-cu directory and add the following line on top of keccak.cuh:

Code:

__constant__ uint64_t keccak_round_constants[24];

unfortunately no improvements on a 6 750ti rig (~52 mhs)

Use cuda 6.5. and remember to compile for compute5.0 (Seems like compute5.0 is included in (CMakeLists.txt) ). Use the version 2nd version I posted with #pragma unroll and bitselect.

And add cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); If your cpu is weak high cpu will slow down alot.

[bensam1231]

Ethereum has no GUI wallet, little windows support, difficult configurations..
I think it's too early to call it a coin, but it seems to have some money support on the exchanges. (pump)
A waste of time.

(Is Ethereum a volunteer botnet?) Join us and we will give you some cash every day.

lol... a lot of promises and hype, little to show for it. It reminds me of Star Citizen. But I go where the money does, so its' what I've been mining.

Also this coin really sucks without pools. Solo mining is for big miners or you get shafted hard. The one pool that is available has closed source code. If they wanted to be fair they should kick off everyone who has greater > 1GH. They can effectively solo mine quite easily. You know, if they actually cared about 51% attacks and being fair.

[sp_]

Hey there,
Thanks for having a look at my CUDA kernel. I actually took the keccak and ROTL64 code from your ccminer fork (or was it someone else's fork..not sure) a few months ago, so I'm surprised it can even be faster .

It must have been another fork. In my latest fork of ccminer(sp-mod 60) (maxcoin) Keccak256 (uint2/bitselect) is running at:

400MHASH gtx 970
155MHASH gtx 750ti

In old cudaminer it was something like 70MHASH on the 750ti and 200 on the 970.

Found a doc with the old keccak hashrates (maxcoin) in cudaminer:

In the old cudaminer a 780ti is doing 218MHASH (250watt)

https://docs.google.com/spreadsheets/d/1BIaD-12rmsoz3t64k3_hU79morafOURkEUjXKSdJVLo/edit?pli=1#gid=0


These algos can not be compared directly

The ether algo is doing a double round of modded keccak512 (24 iterations) + some some other stuff.

[skunk]

To improve the cpu usage you can try to add:

cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);

like this:

Code:

void run_ethash_search(
	uint32_t blocks,
	uint32_t threads,
	cudaStream_t stream,
	uint32_t* g_output,
	hash32_t const* g_header,
	hash128_t const* g_dag,
	uint64_t start_nonce,
	uint64_t target
)
{
	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
	ethash_search
#if __CUDA_ARCH__ >= SHUFFLE_MIN_VER
	ethash_search <<<blocks, threads, 0, stream >>>(g_output, g_header, g_dag, start_nonce, target);
#else
	ethash_search <<<blocks, threads, (sizeof(compute_hash_share) * threads) / THREADS_PER_HASH, stream>>>(g_output, g_header, g_dag, start_nonce, target);
#endif
}

this change causes this:

Code:

  ✘  17:49:23|cudaminer1  Error CUDA mining: cannot set while device is active in this process
  ✘  17:49:23|cudaminer2  Error CUDA mining: cannot set while device is active in this process
  ✘  17:49:23|cudaminer3  Error CUDA mining: cannot set while device is active in this process
  ✘  17:49:23|cudaminer0  Error CUDA mining: cannot set while device is active in this process
  ✘  17:49:23|cudaminer4  Error CUDA mining: cannot set while device is active in this process
  ✘  17:49:23|cudaminer5  Error CUDA mining: cannot set while device is active in this process

[skunk]

Nice work SP , I wish to test it , but can't compile unfortunately...

copy cuda_helper.h from sp's ccminer into libethash-cu directory and add the following line on top of keccak.cuh:

Code:

__constant__ uint64_t keccak_round_constants[24];

unfortunately no improvements on a 6 750ti rig (~52 mhs)

Use cuda 6.5. and remember to compile for compute5.0 (Seems like compute5.0 is included in (CMakeLists.txt) ). Use the version 2nd version I posted with #pragma unroll and bitselect.

And add cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync); If your cpu is weak high cpu will slow down alot.

never moved away from cuda 6.5 and yes, i'm sure it's compiled for compute 5.0...

[sp_]

sorry

Add it after cudadevicereset()

in the ethash_cu_miner.cpp

Code:

	cudaDeviceReset();
//	cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);

for maxwell comment out cudaSharedMemBankSizeEightByte and add cudaFuncCachePreferL1

[ryen123]

sorry

Add it after cudadevicereset()

in the ethash_cu_miner.cpp

Code:

	cudaDeviceReset();
//	cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);

for maxwell comment out cudaSharedMemBankSizeEightByte and add cudaFuncCachePreferL1

Wow cpu utilization dropped to zero.

[sp_]

sorry

Add it after cudadevicereset()

in the ethash_cu_miner.cpp

Code:

	cudaDeviceReset();
//	cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);

for maxwell comment out cudaSharedMemBankSizeEightByte and add cudaFuncCachePreferL1

Wow cpu utilization dropped to zero.

And the speed? Any faster?

[ryen123]

And the speed? Any faster?

I compiled for windows with cuda 6.5. There is a windows bug in it, my 750Ti is only at 2.2+MHs instead of the possible 9+MHs.

[skunk]

sorry

Add it after cudadevicereset()

in the ethash_cu_miner.cpp

Code:

	cudaDeviceReset();
//	cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
	cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
	cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);

for maxwell comment out cudaSharedMemBankSizeEightByte and add cudaFuncCachePreferL1

yess! that did it
cpu is almost idle now but performance wise the hash rate is still the same (even with your keccak changes)

[hashbrown9000]

submitted a pull request for file util.cpp that prevents compiling under linux.  somehow it shows up under tpruvot's fork. can you check, @sp_ ?

[djm34]

There is no projectfile for visual studio 2013. and the cuda 7.0 compiler is slow. Perhaps 30% faster with cuda 6.5.

You can tip the crypto mining blog, they have compiled a windows version before..

Ask them to build for cuda 6.5

doesn't make much of a difference actually, everything is slowed down my mem load/store. (moving to 32bit would probably help a little as well but get some problem with some dll)
actually can't get my 780ti/750ti to work properly because of the high mem allocation (which actually isn't so high... but it doesn't like that)
 

[sp_]

Nvidia fixed a memory allocation bug in one of the latest drivers. 353.30 or newer.you should upgrade

[sp_]

And the speed? Any faster?

I compiled for windows with cuda 6.5. There is a windows bug in it, my 750Ti is only at 2.2+MHs instead of the possible 9+MHs.

Did you build with debug? then build release. I think the kernal need some tuning. Might work on it later. I see work has been done by tvprovut in his fork.

[skunk]

There is no projectfile for visual studio 2013. and the cuda 7.0 compiler is slow. Perhaps 30% faster with cuda 6.5.

You can tip the crypto mining blog, they have compiled a windows version before..

Ask them to build for cuda 6.5

doesn't make much of a difference actually, everything is slowed down my mem load/store. (moving to 32bit would probably help a little as well but get some problem with some dll)
actually can't get my 780ti/750ti to work properly because of the high mem allocation (which actually isn't so high... but it doesn't like that)
 

under linux ethminer allocates 1087mb on each gpu and almost 800mb on the system, i wonder if it's moving hashes between cpu and gpu's memory...

[ryen123]

And the speed? Any faster?

I compiled for windows with cuda 6.5. There is a windows bug in it, my 750Ti is only at 2.2+MHs instead of the possible 9+MHs.

Did you build with debug? then build release. I think the kernal need some tuning. Might work on it later. I see work has been done by tvprovut in his fork.

I always build release. From what I've read on the other forum, the same code works correctly in linux. Others that run in windows get bad hashrates. Whatever the problem is, is beyond my understanding.

My system: Win10, Nvidia driver 355.60, CUDA 6.5

[skunk]

And the speed? Any faster?

I compiled for windows with cuda 6.5. There is a windows bug in it, my 750Ti is only at 2.2+MHs instead of the possible 9+MHs.

Did you build with debug? then build release. I think the kernal need some tuning. Might work on it later. I see work has been done by tvprovut in his fork.

I always build release. From what I've read on the other forum, the same code works correctly in linux. Others that run in windows get bad hashrates. Whatever the problem is, is beyond my understanding.

My system: Win10, Nvidia driver 355.60, CUDA 6.5

windows 10 is the problem, it sends the missing hashes to microsoft's pool