diff options
| -rw-r--r-- | CMakeCache.txt | 38 | ||||
| -rw-r--r-- | CMakeFiles/cmake.check_cache | 1 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/gpu.cpp | 45 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/opencl/blake256.cl (renamed from opencl/blake256.cl) | 7 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/opencl/cryptonight.cl (renamed from opencl/cryptonight.cl) | 18 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/opencl/groestl256.cl (renamed from opencl/groestl256.cl) | 3 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/opencl/jh.cl (renamed from opencl/jh.cl) | 2 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/opencl/wolf-aes.cl (renamed from opencl/wolf-aes.cl) | 3 | ||||
| -rw-r--r-- | backend/amd/amd_gpu/opencl/wolf-skein.cl (renamed from opencl/wolf-skein.cl) | 3 |
9 files changed, 97 insertions, 23 deletions
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 <iostream> #include <vector> #include <algorithm> +#include <regex> #ifdef _WIN32 #include <windows.h> @@ -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<<BuildLog<<std::endl; free(BuildLog); return ERR_OCL_API; @@ -544,7 +545,6 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, size_t platform_idx) if( platformName.find("Advanced Micro Devices") == std::string::npos) { printer::inst()->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 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 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 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 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 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 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 + +)===" |
