Embed OpenCL code into executable to external files aren't necessary

Added 'embedcl' tool which takes a text file and encodes it into a C++ file and can be accessed with provided symbol.
brichard19 · Nov 4, 2018 · e70f260 · e70f260
1 parent 34add47
commit e70f260
Show file tree

Hide file tree

Showing 17 changed files with 287 additions and 27 deletions.
diff --git a/BitCrack.sln b/BitCrack.sln
@@ -11,7 +11,7 @@ Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "util", "util\util.vcxproj",
 EndProject
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "AddressUtil", "AddressUtil\AddressUtil.vcxproj", "{34042455-D274-432D-9134-C9EA41FD1B54}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "KeyFinder", "KeyFinder\KeyFinder.vcxproj", "{D77642A9-365C-420C-A726-469649D2927E}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "cuKeyFinder", "KeyFinder\KeyFinder.vcxproj", "{D77642A9-365C-420C-A726-469649D2927E}"
 	ProjectSection(ProjectDependencies) = postProject
 		{53EE0C03-4419-4767-A91B-7FC7D4B3D2AA} = {53EE0C03-4419-4767-A91B-7FC7D4B3D2AA}
 		{CCA3D02C-5E5A-4A24-B34B-5961DFA93946} = {CCA3D02C-5E5A-4A24-B34B-5961DFA93946}
@@ -65,6 +65,8 @@ Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "clKeyFinder", "clKeyFinder\
 		{BFF4B5FE-C2C5-4384-8941-CD6CB29E78C6} = {BFF4B5FE-C2C5-4384-8941-CD6CB29E78C6}
 	EndProjectSection
 EndProject
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "embedcl", "embedcl\embedcl.vcxproj", "{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}"
+EndProject
 Global
 	GlobalSection(SolutionConfigurationPlatforms) = preSolution
 		Debug|x64 = Debug|x64
@@ -193,6 +195,14 @@ Global
 		{36400E8D-3D04-430C-90A4-FC989E460B3C}.Release|x64.Build.0 = Release|x64
 		{36400E8D-3D04-430C-90A4-FC989E460B3C}.Release|x86.ActiveCfg = Release|Win32
 		{36400E8D-3D04-430C-90A4-FC989E460B3C}.Release|x86.Build.0 = Release|Win32
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Debug|x64.ActiveCfg = Debug|x64
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Debug|x64.Build.0 = Debug|x64
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Debug|x86.ActiveCfg = Debug|Win32
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Debug|x86.Build.0 = Debug|Win32
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Release|x64.ActiveCfg = Release|x64
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Release|x64.Build.0 = Release|x64
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Release|x86.ActiveCfg = Release|Win32
+		{8DA841AA-42FF-40AA-8F12-BC654DF39FEF}.Release|x86.Build.0 = Release|Win32
 	EndGlobalSection
 	GlobalSection(SolutionProperties) = preSolution
 		HideSolutionNode = FALSE

diff --git a/CLKeySearchDevice/CLKeySearchDevice.cpp b/CLKeySearchDevice/CLKeySearchDevice.cpp
@@ -3,6 +3,9 @@
 #include "util.h"
 #include "CLKeySearchDevice.h"
 
+// Defined in bitcrack_cl.cpp which gets build in the pre-build event
+extern char _bitcrack_cl[];
+
 typedef struct {
     int thread;
     int block;
@@ -48,8 +51,9 @@ CLKeySearchDevice::CLKeySearchDevice(uint64_t device, int threads, int pointsPer
     try {
         // Create the context
         _clContext = new cl::CLContext(_device);
-        Logger::log(LogLevel::Info, "Compiling 'KeySearch.cl'...");
-        _clProgram = new cl::CLProgram(*_clContext, util::getExeDirectory() + "KeySearch.cl");
+        Logger::log(LogLevel::Info, "Compiling OpenCL kernels...");
+        //_clProgram = new cl::CLProgram(*_clContext, util::getExeDirectory() + "KeySearch.cl");
+        _clProgram = new cl::CLProgram(*_clContext, _bitcrack_cl);
 
         // Load the kernels
         _initKeysKernel = new cl::CLKernel(*_clProgram, "multiplyStepKernel");

diff --git a/CLKeySearchDevice/CLKeySearchDevice.vcxproj b/CLKeySearchDevice/CLKeySearchDevice.vcxproj
@@ -89,6 +89,15 @@
       <EnableCOMDATFolding>true</EnableCOMDATFolding>
       <OptimizeReferences>true</OptimizeReferences>
     </Link>
+    <PostBuildEvent>
+      <Command>
+      </Command>
+    </PostBuildEvent>
+    <PreBuildEvent>
+      <Command>type ripemd160.cl secp256k1.cl sha256.cl keysearch.cl &gt; bitcrack.cl
+$(SolutionDir)\tools\embedcl.exe bitcrack.cl bitcrack_cl.cpp _bitcrack_cl</Command>
+      <Message>Embed bitcrack.cl into bitcrack_cl.cpp</Message>
+    </PreBuildEvent>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
     <ClCompile>
@@ -99,6 +108,11 @@
       <AdditionalIncludeDirectories>$(SolutionDir)\KeyFinderLib;$(SolutionDir)\clUtil;C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.2\include;$(SolutionDir)\secp256k1lib;$(SolutionDir)\Logger;$(SolutionDir)\util;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
       <PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
     </ClCompile>
+    <PreBuildEvent>
+      <Command>type ripemd160.cl secp256k1.cl sha256.cl keysearch.cl &gt; bitcrack.cl
+$(SolutionDir)\tools\embedcl.exe bitcrack.cl bitcrack_cl.cpp _bitcrack_cl</Command>
+      <Message>Embed bitcrack.cl into bitcrack_cl.cpp</Message>
+    </PreBuildEvent>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
     <ClCompile>
@@ -109,6 +123,15 @@
       <AdditionalIncludeDirectories>$(SolutionDir)\KeyFinderLib;$(SolutionDir)\clUtil;$(OPENCL_INCLUDE);$(SolutionDir)\secp256k1lib;$(SolutionDir)\Logger;$(SolutionDir)\util;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
       <PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
     </ClCompile>
+    <PostBuildEvent>
+      <Command>
+      </Command>
+    </PostBuildEvent>
+    <PreBuildEvent>
+      <Command>type ripemd160.cl secp256k1.cl sha256.cl keysearch.cl &gt; bitcrack.cl
+$(SolutionDir)\tools\embedcl.exe bitcrack.cl bitcrack_cl.cpp _bitcrack_cl</Command>
+      <Message>Embed bitcrack.cl into bitcrack_cl.cpp</Message>
+    </PreBuildEvent>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
     <ClCompile>
@@ -125,11 +148,17 @@
       <EnableCOMDATFolding>true</EnableCOMDATFolding>
       <OptimizeReferences>true</OptimizeReferences>
     </Link>
+    <PreBuildEvent>
+      <Command>type ripemd160.cl secp256k1.cl sha256.cl keysearch.cl &gt; bitcrack.cl
+$(SolutionDir)\tools\embedcl.exe bitcrack.cl bitcrack_cl.cpp _bitcrack_cl</Command>
+      <Message>Embed bitcrack.cl into bitcrack_cl.cpp</Message>
+    </PreBuildEvent>
   </ItemDefinitionGroup>
   <ItemGroup>
     <ClInclude Include="CLKeySearchDevice.h" />
   </ItemGroup>
   <ItemGroup>
+    <ClCompile Include="bitcrack_cl.cpp" />
     <ClCompile Include="CLKeySearchDevice.cpp" />
   </ItemGroup>
   <ItemGroup>

diff --git a/CLKeySearchDevice/keysearch.cl b/CLKeySearchDevice/keysearch.cl
@@ -1,19 +1,7 @@
-#include "sha256.cl"
-#include "ripemd160.cl"
-#include "secp256k1.cl"
-
 #define COMPRESSED 0
 #define UNCOMPRESSED 1
 #define BOTH 2
 
-/*
-typedef struct {
-    ulong mask;
-    ulong size;
-    unsigned int *ptr;
-}CLTargetList;
-*/
-
 unsigned int endian(unsigned int x)
 {
     return (x << 24) | ((x << 8) & 0x00ff0000) | ((x >> 8) & 0x0000ff00) | (x >> 24);
@@ -426,4 +414,4 @@ __kernel void keyFinderKernelWithDouble(
     __global unsigned int *numResults)
 {
     doIterationWithDouble(pointsPerThread, compression, chain, xPtr, yPtr, incXPtr, incYPtr, targetList, numTargets, mask, results, numResults);
-}
+}
diff --git a/CLKeySearchDevice/ripemd160.cl b/CLKeySearchDevice/ripemd160.cl
@@ -491,4 +491,4 @@ void ripemd160sha256NoFinal(const unsigned int x[8], unsigned int digest[5])
     digest[3] = a1 + b2;
     digest[4] = b1 + c2;
 }
-#endif
+#endif
diff --git a/CLKeySearchDevice/secp256k1.cl b/CLKeySearchDevice/secp256k1.cl
@@ -637,4 +637,4 @@ void doBatchInverse(unsigned int inverse[8])
     invModP(inverse);
 }
 
-#endif
+#endif
diff --git a/CLKeySearchDevice/sha256.cl b/CLKeySearchDevice/sha256.cl
@@ -511,4 +511,4 @@ void sha256PublicKeyCompressed(const unsigned int x[8], unsigned int yParity, un
     digest[6] = g;
     digest[7] = h;
 }
-#endif
+#endif
diff --git a/CudaKeySearchDevice/CudaKeySearchDevice.vcxproj b/CudaKeySearchDevice/CudaKeySearchDevice.vcxproj
@@ -84,7 +84,7 @@
     <CudaCompile>
       <TargetMachinePlatform>64</TargetMachinePlatform>
       <GenerateRelocatableDeviceCode>true</GenerateRelocatableDeviceCode>
-      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_50,sm_50;compute_52,sm_52;compute_61,sm_61;compute_70,sm_70;%(CodeGeneration)</CodeGeneration>
+      <CodeGeneration>%(CodeGeneration)</CodeGeneration>
     </CudaCompile>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'">
@@ -107,7 +107,7 @@
     <CudaCompile>
       <TargetMachinePlatform>64</TargetMachinePlatform>
       <GenerateRelocatableDeviceCode>true</GenerateRelocatableDeviceCode>
-      <CodeGeneration>compute_30,sm_30;compute_35,sm_35;compute_50,sm_50;compute_52,sm_52;compute_61,sm_61;compute_70,sm_70;%(CodeGeneration)</CodeGeneration>
+      <CodeGeneration>%(CodeGeneration)</CodeGeneration>
     </CudaCompile>
   </ItemDefinitionGroup>
   <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />

diff --git a/KeyFinder/DeviceManager.cpp b/KeyFinder/DeviceManager.cpp
@@ -25,6 +25,7 @@ std::vector<DeviceManager::DeviceInfo> DeviceManager::getDevices()
             device.type = DeviceType::CUDA;
             device.id = deviceId;
             device.physicalId = cudaDevices[i].id;
+            device.memory = cudaDevices[i].mem;
 
             devices.push_back(device);
 
@@ -46,6 +47,7 @@ std::vector<DeviceManager::DeviceInfo> DeviceManager::getDevices()
             device.type = DeviceType::OpenCL;
             device.id = deviceId;
             device.physicalId = clDevices[i].id;
+            device.memory = clDevices[i].mem;
 
             devices.push_back(device);
             deviceId++;

diff --git a/KeyFinder/DeviceManager.h b/KeyFinder/DeviceManager.h
@@ -32,6 +32,12 @@ typedef struct {
     int id;
     uint64_t physicalId;
     std::string name;
+
+    uint64_t memory;
+
+    // CUDA device info
+    int cudaMajor;
+    int cudaMinor;
 }DeviceInfo;
 
 std::vector<DeviceInfo> getDevices();

diff --git a/KeyFinder/main.cpp b/KeyFinder/main.cpp
@@ -108,6 +108,7 @@ void usage()
 	printf("-r, --range             Number of keys to search\n");
 	printf("-i, --in                Specify file containing addresses, one per line\n");
 	printf("-o, --out               Specify file where results are written\n");
+    printf("-l, --list-devices      List available devices\n");
 }
 
 
@@ -173,9 +174,11 @@ static KeySearchDevice *getDeviceContext(DeviceManager::DeviceInfo &device, int
 
 static void printDeviceList(const std::vector<DeviceManager::DeviceInfo> &devices)
 {
-    printf("ID  Name\n");
     for(int i = 0; i < devices.size(); i++) {
-        printf("%2d  %s\n", devices[i].id, devices[i].name.c_str());
+        printf("ID:     %d\n", devices[i].id);
+        printf("Name:   %s\n", devices[i].name.c_str());
+        printf("Memory: %lldMB\n", devices[i].memory / (1024 * 1024));
+        printf("\n");
     }
 }
 

diff --git a/clKeyFinder/clKeyFinder.vcxproj b/clKeyFinder/clKeyFinder.vcxproj
@@ -76,10 +76,12 @@
   <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
     <TargetName>clBitCrack</TargetName>
     <CustomBuildAfterTargets>Build</CustomBuildAfterTargets>
+    <PostBuildEventUseInBuild>false</PostBuildEventUseInBuild>
   </PropertyGroup>
   <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'">
     <TargetName>clBitCrack</TargetName>
     <CustomBuildAfterTargets>Build</CustomBuildAfterTargets>
+    <PostBuildEventUseInBuild>false</PostBuildEventUseInBuild>
   </PropertyGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
     <ClCompile>
@@ -106,10 +108,12 @@
       <TreatOutputAsContent>true</TreatOutputAsContent>
     </CustomBuildStep>
     <PostBuildEvent>
-      <Command>xcopy /Y "$(SolutionDir)\CLKeySearchDevice\*.cl" .\ &amp;&amp; xcopy /Y "$(SolutionDir)\CLKeySearchDevice\*.cl" $(OutDir)</Command>
+      <Command>
+      </Command>
     </PostBuildEvent>
     <PostBuildEvent>
-      <Message>Copy .cl files</Message>
+      <Message>
+      </Message>
     </PostBuildEvent>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
@@ -163,10 +167,12 @@
       <TreatOutputAsContent>true</TreatOutputAsContent>
     </CustomBuildStep>
     <PostBuildEvent>
-      <Command>xcopy /Y "$(SolutionDir)\CLKeySearchDevice\*.cl" .\ &amp;&amp; xcopy /Y "$(SolutionDir)\CLKeySearchDevice\*.cl" $(OutDir)</Command>
+      <Command>
+      </Command>
     </PostBuildEvent>
     <PostBuildEvent>
-      <Message>Copy .cl files</Message>
+      <Message>
+      </Message>
     </PostBuildEvent>
   </ItemDefinitionGroup>
   <ItemGroup>

diff --git a/clUtil/clContext.cpp b/clUtil/clContext.cpp
@@ -101,6 +101,31 @@ cl::CLProgram::CLProgram(cl::CLContext &ctx, std::string srcFile) : _ctx(ctx)
     clCall(err);
 }
 
+cl::CLProgram::CLProgram(cl::CLContext &ctx, const char *src) : _ctx(ctx)
+{
+    size_t len = strlen(src);
+    cl_int err;
+
+    _prog = clCreateProgramWithSource(ctx.getContext(), 1, &src, &len, &err);
+    clCall(err);
+
+    err = clBuildProgram(_prog, 0, NULL, NULL, NULL, NULL);
+
+    if(err == CL_BUILD_PROGRAM_FAILURE) {
+        size_t logSize;
+        clGetProgramBuildInfo(_prog, ctx.getDevice(), CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
+
+        char *log = new char[logSize];
+        clGetProgramBuildInfo(_prog, ctx.getDevice(), CL_PROGRAM_BUILD_LOG, logSize, log, NULL);
+
+        _buildLog = std::string(log, logSize);
+        delete[] log;
+
+        throw CLException(err, _buildLog);
+    }
+    clCall(err);
+}
+
 std::string cl::CLProgram::loadSource(std::string srcFile)
 {
     std::ifstream f(srcFile);

diff --git a/clUtil/clContext.h b/clUtil/clContext.h
@@ -46,6 +46,8 @@ class CLProgram {
 
 public:
     CLProgram(CLContext &ctx, std::string src);
+    CLProgram(CLContext &ctx, const char *src);
+
     ~CLProgram();
 
     cl_program getProgram();
-Original file line number
+Diff line change
@@ Expand Up @@
         digest[3] = a1 + b2;
         digest[4] = b1 + c2;
     }
-    #endif
+    #endif