move opencl files to amd backend

- inlcude opencl source code during compile

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
+

CMakeFiles/cmake.check_cache

@@ -0,0 +1 @@
+# This file is generated by cmake for dependency checking of the CMakeCache.txt file

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;
 }

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); \
 }
-
-
-
-
-
+)==="

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);
 }
+
+)==="

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)
 
+)==="
+		

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
 
+)==="

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
+
+)==="

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
+
+)==="