Skip to content

Commit 6268734

Browse files
bwyogatamabwyogatama
authored
Merge pull request #38 from sirius-db/issue/11_28
Fix warnings (Issue 28) and unfreed cpu/gpu memory (Issue 11)
2 parents 988b921 + 77e0486 commit 6268734

31 files changed

Lines changed: 221 additions & 197 deletions

‎src/cuda/communication.cu‎

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -126,16 +126,4 @@ void callCudaMemcpyDeviceToDevice(T* dest, T* src, size_t size, int gpu) {
126126
STOP_TIMER();
127127
}
128128

129-
// Define the host function that launches the CUDA kernel
130-
int* sendDataToGPU(int* data, int size) {
131-
SIRIUS_LOG_DEBUG("Send data to GPU");
132-
// use cudamemcpy
133-
int* target;
134-
cudaMalloc((void**) &target, size * sizeof(int));
135-
cudaMemcpy(target, data, size * sizeof(int), cudaMemcpyHostToDevice);
136-
cudaDeviceSynchronize();
137-
SIRIUS_LOG_DEBUG("Done sending data to GPU");
138-
return target;
139-
}
140-
141129
} // namespace duckdb

‎src/cuda/operator/arbitrary_expression.cu‎

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -326,10 +326,8 @@ tableScanExpression(uint8_t **col, uint64_t** offset, uint8_t *constant_compare,
326326
cudaMemcpy(d_data_type, data_type, num_expr * sizeof(int), cudaMemcpyHostToDevice);
327327
CHECK_ERROR();
328328

329-
uint8_t** d_col;
330-
uint64_t** d_offset;
331-
cudaMalloc((void**) &d_col, num_expr * sizeof(uint8_t*));
332-
cudaMalloc((void**) &d_offset, num_expr * sizeof(uint64_t*));
329+
uint8_t** d_col = gpuBufferManager->customCudaMalloc<uint8_t*>(num_expr, 0, 0);
330+
uint64_t** d_offset = gpuBufferManager->customCudaMalloc<uint64_t*>(num_expr, 0, 0);
333331
cudaMemcpy(d_col, col, num_expr * sizeof(uint8_t*), cudaMemcpyHostToDevice);
334332
cudaMemcpy(d_offset, offset, num_expr * sizeof(uint64_t*), cudaMemcpyHostToDevice);
335333
CHECK_ERROR();
@@ -352,8 +350,8 @@ tableScanExpression(uint8_t **col, uint64_t** offset, uint8_t *constant_compare,
352350
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(count), 0);
353351
count = h_count;
354352

355-
cudaFree(d_col);
356-
cudaFree(d_offset);
353+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_col), 0);
354+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_offset), 0);
357355
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_constant_compare), 0);
358356
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_constant_offset), 0);
359357
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_compare_mode), 0);

‎src/cuda/operator/distinct_grouped_aggregate.cu‎

Lines changed: 7 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -179,7 +179,7 @@ void groupedDistinctAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t
179179
//allocate temp memory and copying keys
180180
T* row_keys = gpuBufferManager->customCudaMalloc<T>((num_keys + num_aggregates) * N, 0, 0);
181181
sort_keys_type* materialized_temp = reinterpret_cast<sort_keys_type*> (gpuBufferManager->customCudaMalloc<pointer_and_key>(N, 0, 0));
182-
T** keys_row_id = new T*[num_keys + num_aggregates];
182+
T** keys_row_id = gpuBufferManager->customCudaHostAlloc<T*>(num_keys + num_aggregates);
183183
for (uint64_t i = 0; i < num_keys; i++) {
184184
keys_row_id[i] = reinterpret_cast<T*> (keys[i]);
185185
}
@@ -188,8 +188,7 @@ void groupedDistinctAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t
188188
keys_row_id[num_keys + i] = reinterpret_cast<T*> (aggregate_keys[i]);
189189
}
190190

191-
T** keys_dev;
192-
cudaMalloc((void**) &keys_dev, (num_keys + num_aggregates) * sizeof(T*));
191+
T** keys_dev = gpuBufferManager->customCudaMalloc<T*>(num_keys + num_aggregates, 0, 0);
193192
cudaMemcpy(keys_dev, keys_row_id, (num_keys + num_aggregates) * sizeof(T*), cudaMemcpyHostToDevice);
194193

195194
int tile_items = BLOCK_THREADS * ITEMS_PER_THREAD;
@@ -227,7 +226,7 @@ void groupedDistinctAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t
227226
uint64_t* d_num_runs_out = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
228227
cudaMemset(d_num_runs_out, 0, sizeof(uint64_t));
229228
uint64_t* h_count = gpuBufferManager->customCudaHostAlloc<uint64_t>(1);
230-
uint8_t** output_agg = new uint8_t*[num_aggregates];
229+
uint8_t** output_agg = gpuBufferManager->customCudaHostAlloc<uint8_t*>(num_aggregates);
231230

232231
for (int agg = 0; agg < num_aggregates; agg++) {
233232
// SIRIUS_LOG_DEBUG("Aggregating {}", agg);
@@ -278,9 +277,8 @@ void groupedDistinctAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t
278277
SIRIUS_LOG_DEBUG("Grouped Distinct Aggregate Distinct Count: {}", h_count[0]);
279278
count[0] = h_count[0];
280279

281-
T** keys_dev_result;
282-
T** keys_result = new T*[num_keys];
283-
cudaMalloc((void**) &keys_dev_result, num_keys * sizeof(T*));
280+
T** keys_dev_result = gpuBufferManager->customCudaMalloc<T*>(num_keys, 0, 0);
281+
T** keys_result = gpuBufferManager->customCudaHostAlloc<T*>(num_keys);
284282
for (uint64_t i = 0; i < num_keys; i++) {
285283
keys_result[i] = gpuBufferManager->customCudaMalloc<T>(count[0], 0, 0);
286284
}
@@ -301,8 +299,8 @@ void groupedDistinctAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t
301299
}
302300
}
303301

304-
cudaFree(keys_dev_result);
305-
cudaFree(keys_dev);
302+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
303+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev_result), 0);
306304
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(row_keys), 0);
307305
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(materialized_temp), 0);
308306
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_temp_storage), 0);

‎src/cuda/operator/grouped_aggregate.cu‎

Lines changed: 15 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -256,7 +256,7 @@ void groupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* count,
256256
T* row_keys = gpuBufferManager->customCudaMalloc<T>((num_keys + 1) * N, 0, 0);
257257
sort_keys_type* materialized_temp = reinterpret_cast<sort_keys_type*> (gpuBufferManager->customCudaMalloc<pointer_and_key>(N, 0, 0));
258258
// T* keys_row_id[num_keys + 1];
259-
T** keys_row_id = new T*[num_keys + 1];
259+
T** keys_row_id = gpuBufferManager->customCudaHostAlloc<T*>(num_keys + 1);
260260
for (uint64_t i = 0; i < num_keys; i++) {
261261
keys_row_id[i] = reinterpret_cast<T*> (keys[i]);
262262
}
@@ -267,8 +267,7 @@ void groupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* count,
267267
sequence<BLOCK_THREADS, ITEMS_PER_THREAD><<<(N + tile_items - 1)/tile_items, BLOCK_THREADS>>>(row_sequence, N);
268268
keys_row_id[num_keys] = row_sequence;
269269

270-
T** keys_dev;
271-
cudaMalloc((void**) &keys_dev, (num_keys + 1) * sizeof(T*));
270+
T** keys_dev = gpuBufferManager->customCudaMalloc<T*>(num_keys + 1, 0, 0);
272271
cudaMemcpy(keys_dev, keys_row_id, (num_keys + 1) * sizeof(T*), cudaMemcpyHostToDevice);
273272

274273
columns_to_rows<T, BLOCK_THREADS, ITEMS_PER_THREAD><<<(N + tile_items - 1)/tile_items, BLOCK_THREADS>>>(keys_dev, row_keys, materialized_temp, N, num_keys + 1);
@@ -303,13 +302,13 @@ void groupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* count,
303302

304303
//gather the aggregates based on the row_sequence
305304
SIRIUS_LOG_DEBUG("Gathering Aggregates");
306-
V** aggregate_keys_temp = new V*[num_aggregates];
307-
uint64_t** aggregate_star_temp = new uint64_t*[num_aggregates];
305+
V** aggregate_keys_temp = gpuBufferManager->customCudaHostAlloc<V*>(num_aggregates);
306+
uint64_t** aggregate_star_temp = gpuBufferManager->customCudaHostAlloc<uint64_t*>(num_aggregates);
308307
sort_keys_type* group_by_rows = reinterpret_cast<sort_keys_type*> (gpuBufferManager->customCudaMalloc<pointer_and_key>(N, 0, 0));
309308
uint64_t* d_num_runs_out = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
310309
cudaMemset(d_num_runs_out, 0, sizeof(uint64_t));
311310
uint64_t* h_count = gpuBufferManager->customCudaHostAlloc<uint64_t>(1);
312-
uint8_t** output_agg = new uint8_t*[num_aggregates];
311+
uint8_t** output_agg = gpuBufferManager->customCudaHostAlloc<uint8_t*>(num_aggregates);
313312
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_temp_storage), 0);
314313

315314
for (int agg = 0; agg < num_aggregates; agg++) {
@@ -531,9 +530,8 @@ void groupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* count,
531530
}
532531
}
533532

534-
T** keys_dev_result;
535-
T** keys_result = new T*[num_keys];
536-
cudaMalloc((void**) &keys_dev_result, num_keys * sizeof(T*));
533+
T** keys_dev_result = gpuBufferManager->customCudaMalloc<T*>(num_keys, 0, 0);
534+
T** keys_result = gpuBufferManager->customCudaHostAlloc<T*>(num_keys);
537535
for (uint64_t i = 0; i < num_keys; i++) {
538536
keys_result[i] = gpuBufferManager->customCudaMalloc<T>(count[0], 0, 0);
539537
}
@@ -559,9 +557,9 @@ void groupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* count,
559557
}
560558
}
561559

562-
cudaFree(keys_dev_result);
563-
cudaFree(keys_dev);
564560
//call customCudaFree for all the allocated memory
561+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
562+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev_result), 0);
565563
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(row_keys), 0);
566564
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(materialized_temp), 0);
567565
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(row_sequence), 0);
@@ -587,15 +585,14 @@ void groupedWithoutAggregate(uint8_t **keys, uint64_t* count, uint64_t N, uint64
587585
//allocate temp memory and copying keys
588586
T* row_keys = gpuBufferManager->customCudaMalloc<T>(num_keys * N, 0, 0);
589587
sort_keys_type* materialized_temp = reinterpret_cast<sort_keys_type*> (gpuBufferManager->customCudaMalloc<pointer_and_key>(N, 0, 0));
590-
T** keys_row_id = new T*[num_keys];
588+
T** keys_row_id = gpuBufferManager->customCudaHostAlloc<T*>(num_keys);
591589
for (uint64_t i = 0; i < num_keys; i++) {
592590
keys_row_id[i] = reinterpret_cast<T*> (keys[i]);
593591
}
594592

595593
//generate sequence
596594
int tile_items = BLOCK_THREADS * ITEMS_PER_THREAD;
597-
T** keys_dev;
598-
cudaMalloc((void**) &keys_dev, num_keys * sizeof(T*));
595+
T** keys_dev = gpuBufferManager->customCudaMalloc<T*>(num_keys, 0, 0);
599596
cudaMemcpy(keys_dev, keys_row_id, num_keys * sizeof(T*), cudaMemcpyHostToDevice);
600597

601598
columns_to_rows<T, BLOCK_THREADS, ITEMS_PER_THREAD><<<(N + tile_items - 1)/tile_items, BLOCK_THREADS>>>(keys_dev, row_keys, materialized_temp, N, num_keys);
@@ -669,9 +666,8 @@ void groupedWithoutAggregate(uint8_t **keys, uint64_t* count, uint64_t N, uint64
669666
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_temp_storage), 0);
670667
count[0] = h_count[0];
671668

672-
T** keys_dev_result;
673-
T** keys_result = new T*[num_keys];
674-
cudaMalloc((void**) &keys_dev_result, num_keys * sizeof(T*));
669+
T** keys_dev_result = gpuBufferManager->customCudaMalloc<T*>(num_keys, 0, 0);
670+
T** keys_result = gpuBufferManager->customCudaHostAlloc<T*>(num_keys);
675671
for (uint64_t i = 0; i < num_keys; i++) {
676672
keys_result[i] = gpuBufferManager->customCudaMalloc<T>(count[0], 0, 0);
677673
}
@@ -688,9 +684,9 @@ void groupedWithoutAggregate(uint8_t **keys, uint64_t* count, uint64_t N, uint64
688684
keys[i] = reinterpret_cast<uint8_t*> (keys_result[i]);
689685
}
690686

691-
cudaFree(keys_dev_result);
692-
cudaFree(keys_dev);
693687
//call customCudaFree for all the allocated memory
688+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
689+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev_result), 0);
694690
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(row_keys), 0);
695691
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(materialized_temp), 0);
696692
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(group_by_rows), 0);

‎src/cuda/operator/hash_grouped_aggregate.cu‎

Lines changed: 12 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -405,12 +405,10 @@ void hashGroupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* co
405405
SIRIUS_LOG_DEBUG("Launching Hash Grouped Aggregate Kernel");
406406
GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
407407

408-
T** keys_dev;
409-
cudaMalloc((void**) &keys_dev, num_keys * sizeof(T*));
408+
T** keys_dev = gpuBufferManager->customCudaMalloc<T*>(num_keys, 0, 0);
410409
cudaMemcpy(keys_dev, reinterpret_cast<T*>(keys), num_keys * sizeof(T*), cudaMemcpyHostToDevice);
411410

412-
V** aggregate_keys_dev;
413-
cudaMalloc((void**) &aggregate_keys_dev, num_aggregates * sizeof(V*));
411+
V** aggregate_keys_dev = gpuBufferManager->customCudaMalloc<V*>(num_aggregates, 0, 0);
414412
cudaMemcpy(aggregate_keys_dev, reinterpret_cast<V*>(aggregate_keys), num_aggregates * sizeof(V*), cudaMemcpyHostToDevice);
415413

416414
int* agg_mode_dev = gpuBufferManager->customCudaMalloc<int>(num_aggregates, 0, 0);
@@ -437,8 +435,8 @@ void hashGroupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* co
437435
CHECK_ERROR();
438436
}
439437

440-
T* max_key_host = new T[num_keys];
441-
T* min_key_host = new T[num_keys];
438+
T* max_key_host = gpuBufferManager->customCudaHostAlloc<T>(num_keys);
439+
T* min_key_host = gpuBufferManager->customCudaHostAlloc<T>(num_keys);
442440
cudaMemcpy(max_key_host, max_key, num_keys * sizeof(T), cudaMemcpyDeviceToHost);
443441
cudaMemcpy(min_key_host, min_key, num_keys * sizeof(T), cudaMemcpyDeviceToHost);
444442

@@ -501,8 +499,8 @@ void hashGroupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* co
501499

502500
CHECK_ERROR();
503501

504-
uint8_t** keys_dev_result;
505-
uint8_t** aggregate_keys_dev_result;
502+
uint8_t** keys_dev_result = gpuBufferManager->customCudaMalloc<uint8_t*>(num_keys, 0, 0);
503+
uint8_t** aggregate_keys_dev_result = gpuBufferManager->customCudaMalloc<uint8_t*>(num_aggregates, 0, 0);
506504
uint64_t* d_count = gpuBufferManager->customCudaMalloc<uint64_t>(1, 0, 0);
507505
cudaMemset(d_count, 0, sizeof(uint64_t));
508506
scan_hash_group<T, V, BLOCK_THREADS, ITEMS_PER_THREAD><<<(ht_len + tile_items - 1)/tile_items, BLOCK_THREADS>>>(ht, nullptr, nullptr,
@@ -513,8 +511,8 @@ void hashGroupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* co
513511
cudaMemcpy(h_count, d_count, sizeof(uint64_t), cudaMemcpyDeviceToHost);
514512
assert(h_count[0] > 0);
515513

516-
uint8_t** keys_result = new uint8_t*[num_keys];
517-
uint8_t** aggregate_keys_result = new uint8_t*[num_aggregates];
514+
uint8_t** keys_result = gpuBufferManager->customCudaHostAlloc<uint8_t*>(num_keys);
515+
uint8_t** aggregate_keys_result = gpuBufferManager->customCudaHostAlloc<uint8_t*>(num_aggregates);
518516
for (int i = 0; i < num_keys; i++) {
519517
keys_result[i] = gpuBufferManager->customCudaMalloc<uint8_t>(h_count[0] * sizeof(T), 0, 0);
520518
}
@@ -528,8 +526,6 @@ void hashGroupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* co
528526
}
529527
}
530528

531-
cudaMalloc((void**) &keys_dev_result, num_keys * sizeof(uint8_t*));
532-
cudaMalloc((void**) &aggregate_keys_dev_result, num_aggregates * sizeof(uint8_t*));
533529
cudaMemcpy(keys_dev_result, keys_result, num_keys * sizeof(uint8_t*), cudaMemcpyHostToDevice);
534530
cudaMemcpy(aggregate_keys_dev_result, aggregate_keys_result, num_aggregates * sizeof(uint8_t*), cudaMemcpyHostToDevice);
535531
cudaMemset(d_count, 0, sizeof(uint64_t));
@@ -556,10 +552,10 @@ void hashGroupedAggregate(uint8_t **keys, uint8_t **aggregate_keys, uint64_t* co
556552

557553
count[0] = h_count[0];
558554

559-
cudaFree(keys_dev);
560-
cudaFree(aggregate_keys_dev);
561-
cudaFree(keys_dev_result);
562-
cudaFree(aggregate_keys_dev_result);
555+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
556+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(aggregate_keys_dev), 0);
557+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev_result), 0);
558+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(aggregate_keys_dev_result), 0);
563559
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(d_count), 0);
564560
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(agg_mode_dev), 0);
565561
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(max_key), 0);

‎src/cuda/operator/hash_join_inner.cu‎

Lines changed: 6 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -124,7 +124,6 @@ __global__ void probe_multikey(uint64_t **keys, unsigned long long* ht, uint64_t
124124
else if (equal_keys == 2) slot = hash64_multikey(keys[0][tile_offset + threadIdx.x + ITEM * B], keys[1][tile_offset + threadIdx.x + ITEM * B]) % ht_len;
125125
else cudaAssert(0);
126126

127-
bool found = 0;
128127
while (ht[slot * n_ht_column] != 0xFFFFFFFFFFFFFFFF) {
129128
bool local_found = 1;
130129
for (int n = 0; n < num_keys; n++) {
@@ -211,13 +210,12 @@ void buildHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uin
211210
GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
212211

213212
//reinterpret cast the keys to uint64_t
214-
uint64_t** keys_data = new uint64_t*[num_keys];
213+
uint64_t** keys_data = gpuBufferManager->customCudaHostAlloc<uint64_t*>(num_keys);
215214
for (int idx = 0; idx < num_keys; idx++) {
216215
keys_data[idx] = reinterpret_cast<uint64_t*>(keys[idx]);
217216
}
218217

219-
uint64_t** keys_dev;
220-
cudaMalloc((void**) &keys_dev, num_keys * sizeof(uint64_t*));
218+
uint64_t** keys_dev = gpuBufferManager->customCudaMalloc<uint64_t*>(num_keys, 0, 0);
221219
cudaMemcpy(keys_dev, keys_data, num_keys * sizeof(uint64_t*), cudaMemcpyHostToDevice);
222220

223221
int equal_keys = 0;
@@ -235,7 +233,7 @@ void buildHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uin
235233
cudaDeviceSynchronize();
236234
STOP_TIMER();
237235

238-
cudaFree(keys_dev);
236+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
239237
}
240238

241239
void probeHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uint64_t* &row_ids_left, uint64_t* &row_ids_right, uint64_t* &count, uint64_t N, int* condition_mode, int num_keys, bool is_right) {
@@ -258,13 +256,12 @@ void probeHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uin
258256
uint64_t* offset_each_thread = gpuBufferManager->customCudaMalloc<uint64_t>(((N + tile_items - 1)/tile_items) * BLOCK_THREADS, 0, 0);
259257

260258
//reinterpret cast the keys to uint64_t
261-
uint64_t** keys_data = new uint64_t*[num_keys];
259+
uint64_t** keys_data = gpuBufferManager->customCudaHostAlloc<uint64_t*>(num_keys);
262260
for (int idx = 0; idx < num_keys; idx++) {
263261
keys_data[idx] = reinterpret_cast<uint64_t*>(keys[idx]);
264262
}
265263

266-
uint64_t** keys_dev;
267-
cudaMalloc((void**) &keys_dev, num_keys * sizeof(uint64_t*));
264+
uint64_t** keys_dev = gpuBufferManager->customCudaMalloc<uint64_t*>(num_keys, 0, 0);
268265
cudaMemcpy(keys_dev, keys_data, num_keys * sizeof(uint64_t*), cudaMemcpyHostToDevice);
269266

270267
int equal_keys = 0;
@@ -307,7 +304,7 @@ void probeHashTable(uint8_t **keys, unsigned long long* ht, uint64_t ht_len, uin
307304
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(offset_each_thread), 0);
308305
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(condition_mode_dev), 0);
309306
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(count), 0);
310-
cudaFree(keys_dev);
307+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
311308
count = h_count;
312309
STOP_TIMER();
313310
}

‎src/cuda/operator/hash_join_right.cu‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -191,13 +191,12 @@ void probeHashTableRightSemiAnti(uint8_t **keys, unsigned long long* ht, uint64_
191191
GPUBufferManager* gpuBufferManager = &(GPUBufferManager::GetInstance());
192192

193193
//reinterpret cast the keys to uint64_t
194-
uint64_t** keys_data = new uint64_t*[num_keys];
194+
uint64_t** keys_data = gpuBufferManager->customCudaHostAlloc<uint64_t*>(num_keys);
195195
for (int idx = 0; idx < num_keys; idx++) {
196196
keys_data[idx] = reinterpret_cast<uint64_t*>(keys[idx]);
197197
}
198198

199-
uint64_t** keys_dev;
200-
cudaMalloc((void**) &keys_dev, num_keys * sizeof(uint64_t*));
199+
uint64_t** keys_dev = gpuBufferManager->customCudaMalloc<uint64_t*>(num_keys, 0, 0);
201200
cudaMemcpy(keys_dev, keys_data, num_keys * sizeof(uint64_t*), cudaMemcpyHostToDevice);
202201

203202
int equal_keys = 0;
@@ -215,7 +214,7 @@ void probeHashTableRightSemiAnti(uint8_t **keys, unsigned long long* ht, uint64_
215214

216215
SIRIUS_LOG_DEBUG("Finished probe right");
217216
STOP_TIMER();
218-
cudaFree(keys_dev);
217+
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(keys_dev), 0);
219218
gpuBufferManager->customCudaFree(reinterpret_cast<uint8_t*>(condition_mode_dev), 0);
220219
}
221220

0 commit comments

Comments
 (0)