Commit 8a00d83f authored by Ondrej Mosnáček's avatar Ondrej Mosnáček 🌽

C++-ify the kernel code a bit

parent 2bf470a4
......@@ -9,6 +9,10 @@ set(CMAKE_CXX_STANDARD_REQUIRED ON)
find_package(CUDA REQUIRED)
set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};
-std=c++11 -O3 --ptxas-options=-v -arch sm_30 -lineinfo
)
add_library(argon2-gpu-common SHARED
lib/argon2-gpu-common/argon2params.cpp
lib/argon2-gpu-common/blake2b.cpp
......
#include "kernels.h"
#include <cuda_runtime.h>
#define ARGON2_D 0
#define ARGON2_I 1
......@@ -15,6 +13,11 @@
#define THREADS_PER_LANE 32
#define QWORDS_PER_THREAD (ARGON2_QWORDS_IN_BLOCK / 32)
namespace argon2 {
namespace cuda {
using namespace std;
__device__ uint64_t rotr64(uint64_t x, uint32_t n)
{
return (x >> n) | (x << (64 - n));
......@@ -592,3 +595,6 @@ void argon2_run_kernel_oneshot(
}
}
}
} // cuda
} // argon2
......@@ -2,16 +2,28 @@
#define ARGON2_CUDA_KERNELS_H
#include <cuda_runtime.h>
#include <stdint.h>
#include <cstdint>
/* workaround weird CMake/CUDA bug: */
#ifdef argon2
#undef argon2
#endif
namespace argon2 {
namespace cuda {
void argon2_run_kernel_segment(
uint32_t type, uint32_t version, uint32_t batchSize,
cudaStream_t stream, void *memory, uint32_t passes, uint32_t lanes,
uint32_t segment_blocks, uint32_t pass, uint32_t slice);
std::uint32_t type, std::uint32_t version, std::uint32_t batchSize,
cudaStream_t stream, void *memory, std::uint32_t passes,
std::uint32_t lanes, std::uint32_t segment_blocks, std::uint32_t pass,
std::uint32_t slice);
void argon2_run_kernel_oneshot(
uint32_t type, uint32_t version, uint32_t batchSize,
cudaStream_t stream, void *memory, uint32_t passes, uint32_t lanes,
uint32_t segment_blocks);
std::uint32_t type, std::uint32_t version, std::uint32_t batchSize,
cudaStream_t stream, void *memory, std::uint32_t passes,
std::uint32_t lanes, std::uint32_t segment_blocks);
} // cuda
} // argon2
#endif // ARGON2_CUDA_KERNELS_H
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment