Skip to content

Commit

Permalink
Removing unnecessary 512MB host-side allocation in favor of a properl…
Browse files Browse the repository at this point in the history
…y-sized result buffer. This was an artifact left over from earlier dev that should not have been there.
  • Loading branch information
dave-andersen committed Jan 11, 2014
1 parent 3acdf01 commit 14da03c
Show file tree
Hide file tree
Showing 3 changed files with 10 additions and 8 deletions.
5 changes: 2 additions & 3 deletions src/gpuhash.cu
Original file line number Diff line number Diff line change
Expand Up @@ -64,9 +64,8 @@ int GPUHasher::Initialize() {
return -1;
}

#define N_RESULTS 32768
/* Results holds any maybe-colliding keys */
error = cudaMalloc((void **)&dev_results, sizeof(uint64_t)*N_RESULTS);
error = cudaMalloc((void **)&dev_results, sizeof(uint64_t)*GPUHasher::N_RESULTS);
if (error != cudaSuccess) {
fprintf(stderr, "Could not malloc dev_data (%d)\n", error);
exit(-1);
Expand Down Expand Up @@ -216,7 +215,7 @@ void filter_and_rewrite_sha512_kernel(__restrict__ uint64_t *dev_hashes, const _

if (myword && is_in_filter_twice(dev_countbits, (myword>>18))) {
myword = ((myword & (~(((1ULL<<26) - 1)))) | (spot*8+i));
uint32_t result_slot = atomicInc((uint32_t *)dev_results, N_RESULTS);
uint32_t result_slot = atomicInc((uint32_t *)dev_results, GPUHasher::N_RESULTS);
dev_results[result_slot+1] = myword;
}
}
Expand Down
3 changes: 3 additions & 0 deletions src/gpuhash.h
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,9 @@ class GPUHasher {
int Initialize();
int ComputeHashes(uint64_t data[16], uint64_t H[8]);
~GPUHasher();

static const int N_RESULTS = 32768;

private:
int device_id;
uint64_t *dev_data;
Expand Down
10 changes: 5 additions & 5 deletions src/main_poolminer.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -22,8 +22,8 @@
#endif

#define VERSION_MAJOR 0
#define VERSION_MINOR 7
#define VERSION_EXT "RC2 <experimental>"
#define VERSION_MINOR 8
#define VERSION_EXT "GPU0.2 <experimental>"

#define MAX_THREADS 64

Expand Down Expand Up @@ -171,13 +171,13 @@ class CWorkerThread { // worker=miner
}

void run() {
std::cout << "[WORKER" << _id << "] Hello, World!" << std::endl;
std::cout << "[WORKER" << _id << "] starting" << std::endl;

/* Ensure that thread is pinned to its allocation */
_hashblock = (uint64_t *)malloc(sizeof(uint64_t) * (1<<26));
_hashblock = (uint64_t *)malloc(sizeof(uint64_t) * GPUHasher::N_RESULTS);
_gpu = new GPUHasher(gpu_device_id);
_gpu->Initialize();

#ifndef MAP_HUGETLB
_collisionMap = (uint32_t*)malloc(sizeof(uint32_t)*(1 << COLLISION_TABLE_BITS));
#else
Expand Down

0 comments on commit 14da03c

Please sign in to comment.