Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Slight OpenCL improvements #1634

Merged
merged 2 commits into from
Feb 22, 2019
Merged

Slight OpenCL improvements #1634

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c041db1
Slight OpenCL improvements
SergiySW Jan 7, 2019
3ad265b
Merge remote-tracking branch 'upstream/master' into opencl_improvements
SergiySW Jan 25, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 4 additions & 18 deletions nano/node/openclwork.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,15 +47,15 @@ typedef struct __blake2b_state
uchar last_node;
} blake2b_state;

__constant static ulong blake2b_IV[8] =
__constant static const ulong blake2b_IV[8] =
{
0x6a09e667f3bcc908UL, 0xbb67ae8584caa73bUL,
0x3c6ef372fe94f82bUL, 0xa54ff53a5f1d36f1UL,
0x510e527fade682d1UL, 0x9b05688c2b3e6c1fUL,
0x1f83d9abfb41bd6bUL, 0x5be0cd19137e2179UL
};

__constant static uchar blake2b_sigma[12][16] =
__constant static const uchar blake2b_sigma[12][16] =
{
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } ,
{ 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } ,
Expand Down Expand Up @@ -94,23 +94,9 @@ static inline int blake2b_increment_counter( blake2b_state *S, const ulong inc )
return 0;
}

static inline uint load32( const void *src )
{
#if defined(NATIVE_LITTLE_ENDIAN)
return *( uint * )( src );
#else
const uchar *p = ( uchar * )src;
uint w = *p++;
w |= ( uint )( *p++ ) << 8;
w |= ( uint )( *p++ ) << 16;
w |= ( uint )( *p++ ) << 24;
return w;
#endif
}

static inline ulong load64( const void *src )
{
#if defined(NATIVE_LITTLE_ENDIAN)
#if defined(__ENDIAN_LITTLE__)
return *( ulong * )( src );
#else
const uchar *p = ( uchar * )src;
Expand Down Expand Up @@ -359,7 +345,7 @@ static void ucharcpyglb (uchar * dst, __global uchar const * src, size_t count)
}
}

__kernel void nano_work (__global ulong * attempt, __global ulong * result_a, __global uchar * item_a)
__kernel void nano_work (__global ulong const * attempt, __global ulong * result_a, __global uchar const * item_a)
{
int const thread = get_global_id (0);
uchar item_l [32];
Expand Down