From 0e4cbe2643e5069ab60c53cb2fa81e3afa8456ce Mon Sep 17 00:00:00 2001 From: psychocrypt Date: Wed, 27 Sep 2017 21:30:43 +0200 Subject: [PATCH] move opencl files to amd backend - inlcude opencl source code during compile --- CMakeCache.txt | 38 ++++++++++++++++ CMakeFiles/cmake.check_cache | 1 + backend/amd/amd_gpu/gpu.cpp | 45 +++++++++++++------ .../amd/amd_gpu/opencl}/blake256.cl | 7 +-- .../amd/amd_gpu/opencl}/cryptonight.cl | 18 +++++--- .../amd/amd_gpu/opencl}/groestl256.cl | 3 ++ {opencl => backend/amd/amd_gpu/opencl}/jh.cl | 2 + .../amd/amd_gpu/opencl}/wolf-aes.cl | 3 ++ .../amd/amd_gpu/opencl}/wolf-skein.cl | 3 ++ 9 files changed, 97 insertions(+), 23 deletions(-) create mode 100644 CMakeCache.txt create mode 100644 CMakeFiles/cmake.check_cache rename {opencl => backend/amd/amd_gpu/opencl}/blake256.cl (99%) rename {opencl => backend/amd/amd_gpu/opencl}/cryptonight.cl (99%) rename {opencl => backend/amd/amd_gpu/opencl}/groestl256.cl (99%) rename {opencl => backend/amd/amd_gpu/opencl}/jh.cl (99%) rename {opencl => backend/amd/amd_gpu/opencl}/wolf-aes.cl (99%) rename {opencl => backend/amd/amd_gpu/opencl}/wolf-skein.cl (99%) diff --git a/CMakeCache.txt b/CMakeCache.txt new file mode 100644 index 0000000..3c42ff2 --- /dev/null +++ b/CMakeCache.txt @@ -0,0 +1,38 @@ +# This is the CMakeCache file. +# For build in directory: /home/psychocrypt/workspace/xmr-stak-cpu +# It was generated by CMake: /usr/bin/cmake +# You can edit this file to change values found and used by cmake. +# If you do not want to change any of the values, simply exit the editor. +# If you do want to change a value, simply edit, save, and exit the editor. +# The syntax for the file is as follows: +# KEY:TYPE=VALUE +# KEY is the name of a variable in the cache. +# TYPE is a hint to GUIs for the type of VALUE, DO NOT EDIT TYPE!. +# VALUE is the current value for the KEY. + +######################## +# EXTERNAL cache entries +######################## + + +######################## +# INTERNAL cache entries +######################## + +//This is the directory where this CMakeCache.txt was created +CMAKE_CACHEFILE_DIR:INTERNAL=/home/psychocrypt/workspace/xmr-stak-cpu +//Major version of cmake used to create the current loaded cache +CMAKE_CACHE_MAJOR_VERSION:INTERNAL=3 +//Minor version of cmake used to create the current loaded cache +CMAKE_CACHE_MINOR_VERSION:INTERNAL=5 +//Patch version of cmake used to create the current loaded cache +CMAKE_CACHE_PATCH_VERSION:INTERNAL=1 +//Path to CMake executable. +CMAKE_COMMAND:INTERNAL=/usr/bin/cmake +//Path to cpack program executable. +CMAKE_CPACK_COMMAND:INTERNAL=/usr/bin/cpack +//Path to ctest program executable. +CMAKE_CTEST_COMMAND:INTERNAL=/usr/bin/ctest +//Path to CMake installation. +CMAKE_ROOT:INTERNAL=/usr/share/cmake-3.5 + diff --git a/CMakeFiles/cmake.check_cache b/CMakeFiles/cmake.check_cache new file mode 100644 index 0000000..3dccd73 --- /dev/null +++ b/CMakeFiles/cmake.check_cache @@ -0,0 +1 @@ +# This file is generated by cmake for dependency checking of the CMakeCache.txt file diff --git a/backend/amd/amd_gpu/gpu.cpp b/backend/amd/amd_gpu/gpu.cpp index 30aa5a8..04e442a 100644 --- a/backend/amd/amd_gpu/gpu.cpp +++ b/backend/amd/amd_gpu/gpu.cpp @@ -19,6 +19,7 @@ #include #include #include +#include #ifdef _WIN32 #include @@ -212,7 +213,7 @@ char* LoadTextFile(const char* filename) return out; } -size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, char* source_code) +size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_code) { size_t MaximumWorkSize; cl_int ret; @@ -330,9 +331,9 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, char* source_code) printer::inst()->print_msg(L1,"Error %s when calling clGetProgramBuildInfo for build log.", err_to_str(ret)); return ERR_OCL_API; } - + printer::inst()->print_str("Build log:\n"); - printer::inst()->print_str(BuildLog); + std::cerr<print_msg(L1,"WARNING: using non AMD device: %s", platformName.c_str()); - return ERR_STUPID_PARAMS; } free(platforms); @@ -607,22 +607,41 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) return ERR_OCL_API; } - char* source_code = LoadTextFile(sSourcePath); - if(source_code == NULL) - { - printer::inst()->print_msg(L1,"Couldn't locate GPU source code file at %s.", sSourcePath); - return ERR_STUPID_PARAMS; - } + //char* source_code = LoadTextFile(sSourcePath); + + const char *cryptonightCL = + #include "./opencl/cryptonight.cl" + ; + const char *blake256CL = + #include "./opencl/blake256.cl" + ; + const char *groestl256CL = + #include "./opencl/groestl256.cl" + ; + const char *jhCL = + #include "./opencl/jh.cl" + ; + const char *wolfAesCL = + #include "./opencl/wolf-aes.cl" + ; + const char *wolfSkeinCL = + #include "./opencl/wolf-skein.cl" + ; + + std::string source_code(cryptonightCL); + source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_WOLF_AES"), wolfAesCL); + source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_WOLF_SKEIN"), wolfSkeinCL); + source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_JH"), jhCL); + source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_BLAKE256"), blake256CL); + source_code = std::regex_replace(source_code, std::regex("XMRSTAK_INCLUDE_GROESTL256"), groestl256CL); for(int i = 0; i < num_gpus; ++i) { - if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code)) != ERR_SUCCESS) + if((ret = InitOpenCLGpu(opencl_ctx, &ctx[i], source_code.c_str())) != ERR_SUCCESS) { - free(source_code); return ret; } } - free(source_code); return ERR_SUCCESS; } diff --git a/opencl/blake256.cl b/backend/amd/amd_gpu/opencl/blake256.cl similarity index 99% rename from opencl/blake256.cl rename to backend/amd/amd_gpu/opencl/blake256.cl index 6be0eb0..3d5fe3e 100644 --- a/opencl/blake256.cl +++ b/backend/amd/amd_gpu/opencl/blake256.cl @@ -1,3 +1,4 @@ +R"===( /* * blake256 kernel implementation. * @@ -89,8 +90,4 @@ __constant static const sph_u32 c_u256[16] = { v[b] ^= v[c]; \ v[b] = rotate(v[b], 25U); \ } - - - - - +)===" diff --git a/opencl/cryptonight.cl b/backend/amd/amd_gpu/opencl/cryptonight.cl similarity index 99% rename from opencl/cryptonight.cl rename to backend/amd/amd_gpu/opencl/cryptonight.cl index 0478160..4aae939 100644 --- a/opencl/cryptonight.cl +++ b/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -1,3 +1,4 @@ +R"===( /* * This program is free software: you can redistribute it and/or modify * it under the terms of the GNU General Public License as published by @@ -72,11 +73,16 @@ inline int amd_bfe(const uint src0, const uint offset, const uint width) } #endif -#include "opencl/wolf-aes.cl" -#include "opencl/wolf-skein.cl" -#include "opencl/jh.cl" -#include "opencl/blake256.cl" -#include "opencl/groestl256.cl" +//#include "opencl/wolf-aes.cl" +XMRSTAK_INCLUDE_WOLF_AES +//#include "opencl/wolf-skein.cl" +XMRSTAK_INCLUDE_WOLF_SKEIN +//#include "opencl/jh.cl" +XMRSTAK_INCLUDE_JH +//#include "opencl/blake256.cl" +XMRSTAK_INCLUDE_BLAKE256 +//#include "opencl/groestl256.cl" +XMRSTAK_INCLUDE_GROESTL256 static const __constant ulong keccakf_rndc[24] = { @@ -968,3 +974,5 @@ __kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global //for(int i = 0; i < 4; ++i) output[i] = State[i + 4]; if(as_uint2(State[7]).s1 <= Target) output[atomic_inc(output + 0xFF)] = BranchBuf[idx] + get_global_offset(0); } + +)===" \ No newline at end of file diff --git a/opencl/groestl256.cl b/backend/amd/amd_gpu/opencl/groestl256.cl similarity index 99% rename from opencl/groestl256.cl rename to backend/amd/amd_gpu/opencl/groestl256.cl index 1139c4d..403820d 100644 --- a/opencl/groestl256.cl +++ b/backend/amd/amd_gpu/opencl/groestl256.cl @@ -1,3 +1,4 @@ +R"===( /* $Id: groestl.c 260 2011-07-21 01:02:38Z tp $ */ /* * Groestl256 @@ -287,3 +288,5 @@ static const __constant ulong T4_G[] = ROUND_SMALL_Q(a, r); \ } while (0) +)===" + \ No newline at end of file diff --git a/opencl/jh.cl b/backend/amd/amd_gpu/opencl/jh.cl similarity index 99% rename from opencl/jh.cl rename to backend/amd/amd_gpu/opencl/jh.cl index d034abc..fe70ea3 100644 --- a/opencl/jh.cl +++ b/backend/amd/amd_gpu/opencl/jh.cl @@ -1,3 +1,4 @@ +R"===( /* $Id: jh.c 255 2011-06-07 19:50:20Z tp $ */ /* * JH implementation. @@ -270,3 +271,4 @@ static const __constant ulong C[] = #endif +)===" diff --git a/opencl/wolf-aes.cl b/backend/amd/amd_gpu/opencl/wolf-aes.cl similarity index 99% rename from opencl/wolf-aes.cl rename to backend/amd/amd_gpu/opencl/wolf-aes.cl index 23a41a7..996944b 100644 --- a/opencl/wolf-aes.cl +++ b/backend/amd/amd_gpu/opencl/wolf-aes.cl @@ -1,3 +1,4 @@ +R"===( #ifndef WOLF_AES_CL #define WOLF_AES_CL @@ -85,3 +86,5 @@ uint4 AES_Round(const __local uint *AES0, const __local uint *AES1, const __loca } #endif + +)===" diff --git a/opencl/wolf-skein.cl b/backend/amd/amd_gpu/opencl/wolf-skein.cl similarity index 99% rename from opencl/wolf-skein.cl rename to backend/amd/amd_gpu/opencl/wolf-skein.cl index c5c2bd8..868757b 100644 --- a/opencl/wolf-skein.cl +++ b/backend/amd/amd_gpu/opencl/wolf-skein.cl @@ -1,3 +1,4 @@ +R"===( #ifndef WOLF_SKEIN_CL #define WOLF_SKEIN_CL @@ -109,3 +110,5 @@ ulong8 Skein512Block(ulong8 p, ulong8 h, ulong h8, const ulong *t) } #endif + +)==="