Skip to content

Commit

Permalink
64-bit Visual Studio build
Browse files Browse the repository at this point in the history 
-Added x64 build to Visual Studio project. Interestingly, 64-bit runs a little slower than the 32-bit on my machine.
-Small optimization to RIPEMD160 by skipping the MD strengthening step and instead subtracting the IV from the target hash
-Small optimization to SHA256 by hard-coding values known to be 0.
-Removed unused code
  • Loading branch information
@brichard19
brichard19 committed Feb 27, 2018
1 parent eff70cb commit 1317f09
Show file tree
Hide file tree
Showing 18 changed files with 353 additions and 135 deletions.
4 changes: 4 additions & 0 deletions AddrGen/AddrGen.vcxproj
View file
Original file line number Diff line number Diff line change
Expand Up @@ -103,10 +103,12 @@
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<PreprocessorDefinitions>_DEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\secp256k1lib;$(SolutionDir)\util;$(SolutionDir)\AddressUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Console</SubSystem>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalDependencies>BCrypt.lib;%(AdditionalDependencies)</AdditionalDependencies>
</Link>
</ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
Expand Down Expand Up @@ -137,12 +139,14 @@
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<PreprocessorDefinitions>NDEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\secp256k1lib;$(SolutionDir)\util;$(SolutionDir)\AddressUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Console</SubSystem>
<EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalDependencies>BCrypt.lib;%(AdditionalDependencies)</AdditionalDependencies>
</Link>
</ItemDefinitionGroup>
<ItemGroup>
Expand Down
2 changes: 2 additions & 0 deletions AddressUtil/AddressUtil.vcxproj
View file
Original file line number Diff line number Diff line change
Expand Up @@ -105,6 +105,7 @@
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<PreprocessorDefinitions>_DEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\secp256k1lib;$(SolutionDir)\CryptoUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Windows</SubSystem>
Expand Down Expand Up @@ -136,6 +137,7 @@
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<PreprocessorDefinitions>NDEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\secp256k1lib;$(SolutionDir)\CryptoUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Windows</SubSystem>
Expand Down
2 changes: 1 addition & 1 deletion AddressUtil/Base58.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -56,7 +56,7 @@ void Base58::toHash160(const std::string &s, unsigned int hash[5])

bool Base58::isBase58(std::string s)
{
for(int i = 0; i < s.length(); i++) {
for(unsigned int i = 0; i < s.length(); i++) {
if(BASE58_STRING.find(s[i]) < 0) {
return false;
}
Expand Down
10 changes: 8 additions & 2 deletions KeyFinder/KeyFinder.vcxproj
View file
Original file line number Diff line number Diff line change
Expand Up @@ -42,14 +42,14 @@
<ConfigurationType>Application</ConfigurationType>
<UseDebugLibraries>true</UseDebugLibraries>
<PlatformToolset>v140</PlatformToolset>
<CharacterSet>Unicode</CharacterSet>
<CharacterSet>NotSet</CharacterSet>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<UseDebugLibraries>false</UseDebugLibraries>
<PlatformToolset>v140</PlatformToolset>
<WholeProgramOptimization>true</WholeProgramOptimization>
<CharacterSet>Unicode</CharacterSet>
<CharacterSet>NotSet</CharacterSet>
</PropertyGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
<ImportGroup Label="ExtensionSettings">
Expand Down Expand Up @@ -104,10 +104,13 @@
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<PreprocessorDefinitions>_DEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\KeyFinderLib;$(SolutionDir)\util;$(SolutionDir)\AddressUtil;$(SolutionDir)\secp256k1lib;$(SolutionDir)\CmdParse;$(SolutionDir)\cudaDeviceContext;C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include;$(SolutionDir)\cudaUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Console</SubSystem>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalLibraryDirectories>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\lib\x64;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
<AdditionalDependencies>BCrypt.lib;%(AdditionalDependencies)</AdditionalDependencies>
</Link>
</ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
Expand Down Expand Up @@ -139,12 +142,15 @@
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<PreprocessorDefinitions>NDEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\KeyFinderLib;$(SolutionDir)\util;$(SolutionDir)\AddressUtil;$(SolutionDir)\secp256k1lib;$(SolutionDir)\CmdParse;$(SolutionDir)\cudaDeviceContext;C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include;$(SolutionDir)\cudaUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Console</SubSystem>
<EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalLibraryDirectories>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\lib\x64;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
<AdditionalDependencies>BCrypt.lib;%(AdditionalDependencies)</AdditionalDependencies>
</Link>
</ItemDefinitionGroup>
<ItemGroup>
Expand Down
2 changes: 1 addition & 1 deletion KeyFinder/main.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ int main(int argc, char **argv)
parser.parse(argc, argv);
std::vector<OptArg> args = parser.getArgs();

for(int i = 0; i < args.size(); i++) {
for(unsigned int i = 0; i < args.size(); i++) {
OptArg optArg = args[i];
std::string opt = args[i].option;

Expand Down
8 changes: 4 additions & 4 deletions KeyFinderLib/KeyFinder.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ KeyFinder::KeyFinder(const secp256k1::uint256 &start, unsigned long long range,
throw KeyFinderException("Requires at least 1 target");
}

for(int i = 0; i < targetHashes.size(); i++) {
for(unsigned int i = 0; i < targetHashes.size(); i++) {
KeyFinderTarget t;

if(!Address::verifyAddress(targetHashes[i])) {
Expand Down Expand Up @@ -126,13 +126,13 @@ void KeyFinder::generateStartingPoints()
// Generate key pairs for k, k+1, k+2 ... k + <total points in parallel - 1>
secp256k1::uint256 privKey = _startExponent;

for(int i = 0; i < totalPoints; i++) {
for(unsigned int i = 0; i < totalPoints; i++) {
_exponents.push_back(privKey.add(i));
}

secp256k1::generateKeyPairsBulk(secp256k1::G(), _exponents, _startingPoints);

for(int i = 0; i < _startingPoints.size(); i++) {
for(unsigned int i = 0; i < _startingPoints.size(); i++) {
if(!secp256k1::pointExists(_startingPoints[i])) {
throw KeyFinderException("Point generation error");
}
Expand Down Expand Up @@ -227,7 +227,7 @@ void KeyFinder::run()

_devCtx->getKeyFinderResults(results);

for(int i = 0; i < results.size(); i++) {
for(unsigned int i = 0; i < results.size(); i++) {
unsigned int index = _devCtx->getIndex(results[i].block, results[i].thread, results[i].index);

secp256k1::uint256 exp = _exponents[index];
Expand Down
37 changes: 25 additions & 12 deletions KeyFinderLib/KeyFinder.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,10 +18,31 @@ __constant__ unsigned int _TARGET_HASH[5];
__constant__ unsigned int _INC_X[8];
__constant__ unsigned int _INC_Y[8];

static const unsigned int _RIPEMD160_IV_HOST[5] = {
0x67452301,
0xefcdab89,
0x98badcfe,
0x10325476,
0xc3d2e1f0
};

static unsigned int swp(unsigned int x)
{
return (x << 24) | ((x << 8) & 0x00ff0000) | ((x >> 8) & 0x0000ff00) | (x >> 24);
}

cudaError_t setTargetHash(const unsigned int hash[5])
{
return cudaMemcpyToSymbol(_TARGET_HASH, hash, sizeof(unsigned int) * 5);
unsigned int h[5];


// Subtract the RIPEmd160 IV from the hash and swp endian, to avoid performing
// the addition and endian swap at the end of the RIPEMD160 call
for(int i = 0; i < 5; i++) {
h[i] = swp(hash[i]) - _RIPEMD160_IV_HOST[(i + 1) % 5];
}

return cudaMemcpyToSymbol(_TARGET_HASH, h, sizeof(unsigned int) * 5);
}

cudaError_t setIncrementorPoint(const secp256k1::uint256 &x, const secp256k1::uint256 &y)
Expand Down Expand Up @@ -51,11 +72,7 @@ __device__ void hashPublicKey(const unsigned int *x, const unsigned int *y, unsi
hash[i] = endian(hash[i]);
}

ripemd160sha256(hash, hash);

for(int i = 0; i < 5; i++) {
digestOut[i] = endian(hash[i]);
}
ripemd160sha256NoFinal(hash, digestOut);
}

__device__ void hashPublicKeyCompressed(const unsigned int *x, const unsigned int *y, unsigned int *digestOut)
Expand All @@ -69,11 +86,7 @@ __device__ void hashPublicKeyCompressed(const unsigned int *x, const unsigned in
hash[i] = endian(hash[i]);
}

ripemd160sha256(hash, hash);

for(int i = 0; i < 5; i++) {
digestOut[i] = endian(hash[i]);
}
ripemd160sha256NoFinal(hash, digestOut);
}

__device__ void setHashFoundFlag(unsigned int *flagsAra, int idx, int value)
Expand Down Expand Up @@ -105,7 +118,7 @@ __device__ void setResultFound(unsigned int *numResultsPtr, void *results, int i
}

for(int i = 0; i < 5; i++) {
r.digest[i] = digest[i];
r.digest[i] = endian(digest[i] + _RIPEMD160_IV[(i + 1) % 5]);
}

struct KeyFinderDeviceResult *resultsPtr = (struct KeyFinderDeviceResult *)results;
Expand Down
10 changes: 6 additions & 4 deletions KeyFinderLib/KeyFinderLib.vcxproj
View file
Original file line number Diff line number Diff line change
Expand Up @@ -57,9 +57,9 @@
<PlatformToolset>v140</PlatformToolset>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<ConfigurationType>StaticLibrary</ConfigurationType>
<UseDebugLibraries>true</UseDebugLibraries>
<CharacterSet>MultiByte</CharacterSet>
<CharacterSet>NotSet</CharacterSet>
<PlatformToolset>v140</PlatformToolset>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
Expand All @@ -70,10 +70,10 @@
<PlatformToolset>v140</PlatformToolset>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<ConfigurationType>StaticLibrary</ConfigurationType>
<UseDebugLibraries>false</UseDebugLibraries>
<WholeProgramOptimization>true</WholeProgramOptimization>
<CharacterSet>MultiByte</CharacterSet>
<CharacterSet>NotSet</CharacterSet>
<PlatformToolset>v140</PlatformToolset>
</PropertyGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
Expand Down Expand Up @@ -121,6 +121,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)"</Command>
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<PreprocessorDefinitions>WIN32;WIN64;_DEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)secp256k1lib;$(SolutionDir)\cudaDeviceContext;$(SolutionDir)\cudaMath;$(SolutionDir)\util;$(SolutionDir)\AddressUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<GenerateDebugInformation>true</GenerateDebugInformation>
Expand Down Expand Up @@ -163,6 +164,7 @@ copy "$(CudaToolkitBinDir)\cudart*.dll" "$(OutDir)"</Command>
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<PreprocessorDefinitions>WIN32;WIN64;NDEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)secp256k1lib;$(SolutionDir)\cudaDeviceContext;$(SolutionDir)\cudaMath;$(SolutionDir)\util;$(SolutionDir)\AddressUtil;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<GenerateDebugInformation>true</GenerateDebugInformation>
Expand Down
2 changes: 2 additions & 0 deletions cudaDeviceContext/cudaDeviceContext.vcxproj
View file
Original file line number Diff line number Diff line change
Expand Up @@ -97,6 +97,7 @@
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<PreprocessorDefinitions>_DEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include;$(SolutionDir)\secp256k1lib;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Windows</SubSystem>
Expand Down Expand Up @@ -128,6 +129,7 @@
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<PreprocessorDefinitions>NDEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include;$(SolutionDir)\secp256k1lib;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Windows</SubSystem>
Expand Down
4 changes: 4 additions & 0 deletions cudaInfo/cudaInfo.vcxproj
View file
Original file line number Diff line number Diff line change
Expand Up @@ -104,10 +104,12 @@
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<PreprocessorDefinitions>_DEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\cudautil;C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Console</SubSystem>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalLibraryDirectories>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\lib\x64;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
</Link>
</ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
Expand Down Expand Up @@ -139,12 +141,14 @@
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<PreprocessorDefinitions>NDEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<AdditionalIncludeDirectories>$(SolutionDir)\cudautil;C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\include;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
</ClCompile>
<Link>
<SubSystem>Console</SubSystem>
<EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences>
<GenerateDebugInformation>true</GenerateDebugInformation>
<AdditionalLibraryDirectories>C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v8.0\lib\x64;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
</Link>
</ItemDefinitionGroup>
<ItemGroup>
Expand Down
5 changes: 5 additions & 0 deletions cudaMath/ptx.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -33,4 +33,9 @@ __device__ void releaseLock()
atomicExch(&_lock, 0);
}

__device__ __forceinline__ unsigned int endian(unsigned int x)
{
return (x << 24) | ((x << 8) & 0x00ff0000) | ((x >> 8) & 0x0000ff00) | (x >> 24);
}

#endif
Loading

0 comments on commit 1317f09

Please sign in to comment.