Skip to content

Added RandomX JIT for AMD Navi GPUs #1551

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

Merged
xmrig merged 1 commit into from
Feb 13, 2020
Merged

Added RandomX JIT for AMD Navi GPUs #1551

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1,716 changes: 910 additions & 806 deletions src/backend/opencl/cl/rx/randomx_cl.h
View file
Open in desktop

Large diffs are not rendered by default.

223 changes: 165 additions & 58 deletions src/backend/opencl/cl/rx/randomx_jit.cl
View file
Open in desktop

Large diffs are not rendered by default.

742 changes: 742 additions & 0 deletions src/backend/opencl/cl/rx/randomx_run_gfx1010.asm
View file
Open in desktop

Large diffs are not rendered by default.

303 changes: 303 additions & 0 deletions src/backend/opencl/cl/rx/randomx_run_gfx1010.h
View file
Open in desktop

Large diffs are not rendered by default.

1 change: 1 addition & 0 deletions src/backend/opencl/generators/ocl_generic_rx_generator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,7 @@ bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorith
case OclDevice::Navi_10:
case OclDevice::Navi_12:
case OclDevice::Navi_14:
gcnAsm = true;
isNavi = true;
break;

Expand Down
8 changes: 3 additions & 5 deletions src/backend/opencl/kernels/rx/RxRunKernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,12 +30,10 @@
#include "crypto/rx/RxAlgo.h"


void xmrig::RxRunKernel::enqueue(cl_command_queue queue, size_t threads)
void xmrig::RxRunKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
{
const size_t gthreads = threads * 64;
static const size_t lthreads = 64;

enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
const size_t gthreads = threads * workgroup_size;
enqueueNDRange(queue, 1, nullptr, &gthreads, &workgroup_size);
}


Expand Down
2 changes: 1 addition & 1 deletion src/backend/opencl/kernels/rx/RxRunKernel.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,7 +40,7 @@ class RxRunKernel : public OclKernel
public:
inline RxRunKernel(cl_program program) : OclKernel(program, "randomx_run") {}

void enqueue(cl_command_queue queue, size_t threads);
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
void setArgs(cl_mem dataset, cl_mem scratchpads, cl_mem registers, cl_mem rounding, cl_mem programs, uint32_t batch_size, const Algorithm &algorithm);
};

Expand Down
4 changes: 4 additions & 0 deletions src/backend/opencl/runners/OclRxBaseRunner.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,10 @@ xmrig::OclRxBaseRunner::OclRxBaseRunner(size_t index, const OclLaunchData &data)
m_gcn_version = 14;
}

if (data.device.type() == OclDevice::Navi_10 || data.device.type() == OclDevice::Navi_12 || data.device.type() == OclDevice::Navi_14) {
m_gcn_version = 15;
}

m_options += " -DALGO=" + std::to_string(RxAlgo::id(m_algorithm));
m_options += " -DWORKERS_PER_HASH=" + std::to_string(m_worksize);
m_options += " -DGCN_VERSION=" + std::to_string(m_gcn_version);
Expand Down
22 changes: 19 additions & 3 deletions src/backend/opencl/runners/OclRxJitRunner.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,6 +26,7 @@

#include "backend/opencl/cl/rx/randomx_run_gfx803.h"
#include "backend/opencl/cl/rx/randomx_run_gfx900.h"
#include "backend/opencl/cl/rx/randomx_run_gfx1010.h"
#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
#include "backend/opencl/kernels/rx/HashAesKernel.h"
#include "backend/opencl/kernels/rx/RxJitKernel.h"
Expand Down Expand Up @@ -84,7 +85,7 @@ void xmrig::OclRxJitRunner::execute(uint32_t iteration)

OclLib::finish(m_queue);

m_randomx_run->enqueue(m_queue, m_intensity);
m_randomx_run->enqueue(m_queue, m_intensity, (m_gcn_version == 15) ? 32 : 64);
}


Expand Down Expand Up @@ -120,8 +121,23 @@ bool xmrig::OclRxJitRunner::loadAsmProgram()
elf_header_flags = *reinterpret_cast<uint32_t*>((binary_data.data() + elf_header_flags_offset));
}

const size_t len = (m_gcn_version == 14) ? randomx_run_gfx900_bin_size : randomx_run_gfx803_bin_size;
unsigned char *binary = (m_gcn_version == 14) ? randomx_run_gfx900_bin : randomx_run_gfx803_bin;
size_t len;
unsigned char *binary;

switch (m_gcn_version) {
case 14:
len = randomx_run_gfx900_bin_size;
binary = randomx_run_gfx900_bin;
break;
case 15:
len = randomx_run_gfx1010_bin_size;
binary = randomx_run_gfx1010_bin;
break;
default:
len = randomx_run_gfx803_bin_size;
binary = randomx_run_gfx803_bin;
break;
}

// Set correct internal device ID in the pre-compiled binary
if (elf_header_flags) {
Expand Down