Quote from: Kangoo* on January 24, 2021, 12:30:12 AM

Quote from: renedx on January 22, 2021, 03:50:32 PM

Quote from: WhyMe on January 22, 2021, 02:38:04 PM

Last summer I was running bitcrack on some 2060 without any problem. I don't understand all the lasts messages.

EDIT
oh, it's about driver version, ok

Yeah, the Ampere cards do not run on these drivers. Sad story.

Not right. It's bitcrack's kernel code that need to be update. Here the solution (GPUMath), if some developer can apply it in the bitcrack's kernel code

https://github.com/JeanLucPons/VanitySearch/commit/d3c1debb12233722f6ccc09ed3317769161a4773

I just see a bunch of __device__ functions in GPUMath.h, the other files are unrelated to this problem since they are C source code.

In that commit we have:

__device__ __forceinline__ uint32_t ctz(uint64_t x) {

__device__ void _DivStep62(uint64_t u[5],uint64_t v[5],
int32_t* pos

__device__ void MatrixVecMulHalf(uint64_t dest[5],uint64_t u[5],uint64_t v[5],int64_t _11,int64_t _12,uint64_t* carry) {

__device__ void MatrixVecMul(uint64_t u[5],uint64_t v[5],int64_t _11,int64_t _12,int64_t _21,int64_t _22) {

__device__ uint64_t AddCh(uint64_t r[5],uint64_t a[5],uint64_t carry) {

The only thing they have in common with Bitcrack code is that these are also using array parameters and are somehow succeeding.

Here's what I did find in bitcrack code.

In cudaDeviceKeys.cu, elliptic curve-related code about x and y points:

Code:

__constant__ unsigned int *_xPtr[1];

__constant__ unsigned int *_yPtr[1];

__device__ unsigned int *ec::getXPtr()
{
return _xPtr[0];
}

__device__ unsigned int *ec::getYPtr()
{
return _yPtr[0];
}

Why are these a single-element array of a pointer in the first place? It looks redundant and the size of 1 is just making more opportunities to get pointer handling wrong.

It's relevant because after xPtr is copied to device memory in x, then this code is called that uses the arrays x and chain (and also inverse but that is initialized to fixed values so we know that's already aligned).

Code:

beginBatchAdd(_INC_X, x, chain, i, i, inverse);

chain is also initialized to _CHAIN which is also a 1-array of a pointer.

I read some things about __constant__ variables in constant memory and learned that it's read-only and can only be initialized during declaration but I could be wrong and may have to check again.

But I think this part could be the real problem:

Code:

// inside for loop of doIteration()
unsigned int x[8];

unsigned int digest[5];
// ... more array decelerations inside loop, such as for y

I think because it's forcing CUDA to discard and allocate memory repeatedly, it might be putting this array on some unaligned address or something.