Skip to content

Commit

Permalink
Debugging copyback
Browse files Browse the repository at this point in the history
  • Loading branch information
chrispypatt committed Apr 22, 2019
1 parent f517c36 commit cf0ad99
Show file tree
Hide file tree
Showing 3 changed files with 18 additions and 13 deletions.
18 changes: 9 additions & 9 deletions cpuGroupby.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -367,31 +367,31 @@ void cpuGroupby::printResults() {
}

void cpuGroupby::printGPUResults(int* GPU_output_keys, int* GPU_output_values){
std::cout << "Printing GPU Results..." << std::endl;
std::cout << "Printing GPU Results..." << std::endl;

for (int cRow=0; cRow<numGroups; cRow++) {
//print keys for a row
for (int keyIdx=0; keyIdx<num_key_columns; keyIdx++) {
if (keyIdx == 0) {
std::cout << "{";
std::cout << "{";
}
std::cout << GPU_output_keys[numGroups*keyIdx + cRow];
std::cout << GPU_output_keys[numGroups*keyIdx + cRow];
if(keyIdx != num_key_columns-1) {
std::cout << ":";
std::cout << ":";
} else {
std::cout << "}:";
std::cout << "}:";
}
}
// Print values for a row
for (int valIdx=0; valIdx<num_value_columns; valIdx++) {
if (valIdx == 0) {
std::cout << "{";
std::cout << "{";
}
std::cout << GPU_output_values[numGroups*valIdx + cRow];
std::cout << GPU_output_values[numGroups*valIdx + cRow];
if(valIdx != num_value_columns-1) {
std::cout << ":";
std::cout << ":";
} else {
std::cout << "}";
std::cout << "}";
}
}
std::cout << std::endl;
Expand Down
3 changes: 2 additions & 1 deletion groupby_hash.cu
Original file line number Diff line number Diff line change
Expand Up @@ -62,9 +62,10 @@ void groupby_hash_GPU(const int* key_columns_h, int num_key_columns, int num_key

//copy hash idex of keys, removeing -1's which signify not used
// copy_if(policy, index seq start, index seq end, hash keys for comparison, result containing idx to non -1's, comparator)
auto newEnd = thrust::copy_if(thrust::device, hash_d_seq, hash_d_seq + HASH_TABLE_SIZE, hash_key_idx_d, hashTable_idxs_d, is_not_neg_1());
auto newEnd = thrust::copy_if(thrust::device, hash_d_seq, hash_d_seq + HASH_TABLE_SIZE, hash_key_idx_d, hashTable_idxs_d, is_pos());

num_output_rows = newEnd - hashTable_idxs_d;
printf("%d output rows!\n", num_output_rows);

int* output_key_columns_d = NULL;
cudaMalloc(&output_key_columns_d, sizeof(int)*num_key_columns*num_output_rows);
Expand Down
10 changes: 7 additions & 3 deletions groupby_hash_templates.cu
Original file line number Diff line number Diff line change
Expand Up @@ -150,10 +150,14 @@ void copyUnique(
int num_key_rows)
{
int idx = threadIdx.x + blockIdx.x * blockDim.x;
while (idx < num_output_rows){
while (idx < num_output_rows){//num_output_rows){
// printf("%d | : %d : %d \n ",hash_key_idx_d[hashTable_idxs_d[idx]], key_columns_d[hash_key_idx_d[hashTable_idxs_d[idx]]+num_key_rows*0], key_columns_d[hash_key_idx_d[hashTable_idxs_d[idx]]+num_key_rows*1]);

for (int i = 0; i < num_key_columns; i++){//each column of key matrix
// printf(" : %d",key_columns_d[hash_key_idx_d[hashTable_idxs_d[idx]]+num_key_rows*i]);
output_key_columns_d[idx+num_output_rows*i] = key_columns_d[hash_key_idx_d[hashTable_idxs_d[idx]]+num_key_rows*i];//copy original key entry to output
}
// printf("\n");
idx += gridDim.x*blockDim.x;//increment idx by thread space
}
}
Expand Down Expand Up @@ -199,11 +203,11 @@ void copyValues(
}
}

struct is_not_neg_1
struct is_pos
{
__host__ __device__
bool operator()(const int x)
{
return x != -1;
return x >= 0;
}
};

0 comments on commit cf0ad99

Please sign in to comment.