From 1f7a349395fbdee37c32cdd2f3c8c1e2d6bc53f1 Mon Sep 17 00:00:00 2001 From: Keith Kraus Date: Wed, 19 Sep 2018 21:52:12 -0400 Subject: [PATCH 1/6] Allocate validity masks on output arrays of groupbys --- pygdf/libgdf_groupby.py | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/pygdf/libgdf_groupby.py b/pygdf/libgdf_groupby.py index 2aa0a03a37e3..cec28c9d2e96 100644 --- a/pygdf/libgdf_groupby.py +++ b/pygdf/libgdf_groupby.py @@ -9,6 +9,8 @@ from libgdf_cffi import ffi, libgdf +from ..utils import make_mask + class LibGdfGroupby(object): """Groupby object returned by pygdf.DataFrame.groupby(). @@ -104,11 +106,18 @@ def _apply_agg(self, agg_type, result, add_col_values, out_col_values = ffi.NULL if agg_type == "count": - out_col_agg_series = Series( - Buffer(cuda.device_array(col_agg.size, dtype=np.int64))) + out_col_agg_series = Series.from_masked_array( + data=cuda.device_array(col_agg.size, dtype=np.int64), + mask=make_mask(col_agg.size) + ) else: - out_col_agg_series = Series(Buffer(cuda.device_array( - col_agg.size, dtype=self._df[val_col]._column.data.dtype))) + out_col_agg_series = Series.from_masked_array( + data=cuda.device_array( + col_agg.size, + dtype=self._df[val_col]._column.data.dtype + ), + mask=make_mask(col_agg.size) + ) out_col_agg = out_col_agg_series._column.cffi_view From 0309270526dfd598de31ab34654977a4209eb5ce Mon Sep 17 00:00:00 2001 From: Keith Kraus Date: Thu, 20 Sep 2018 09:49:33 -0400 Subject: [PATCH 2/6] Fix relative import --- pygdf/libgdf_groupby.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pygdf/libgdf_groupby.py b/pygdf/libgdf_groupby.py index cec28c9d2e96..c880ccf74943 100644 --- a/pygdf/libgdf_groupby.py +++ b/pygdf/libgdf_groupby.py @@ -9,7 +9,7 @@ from libgdf_cffi import ffi, libgdf -from ..utils import make_mask +from .utils import make_mask class LibGdfGroupby(object): From aeca909a7186fc57d5ee85f29cb5a61479c1c10b Mon Sep 17 00:00:00 2001 From: Keith Kraus Date: Thu, 20 Sep 2018 13:44:09 -0400 Subject: [PATCH 3/6] Don't run kernel of trying to create mask if input is zero sized --- pygdf/cudautils.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/pygdf/cudautils.py b/pygdf/cudautils.py index a66b3d9764f4..edf9f41caa34 100755 --- a/pygdf/cudautils.py +++ b/pygdf/cudautils.py @@ -300,8 +300,9 @@ def gpu_mask_from_devary(ary, bits): def mask_from_devary(ary): bits = make_mask(len(ary)) - gpu_fill_value.forall(bits.size)(bits, 0) - gpu_mask_from_devary.forall(bits.size)(ary, bits) + if bits.size > 0: + gpu_fill_value.forall(bits.size)(bits, 0) + gpu_mask_from_devary.forall(bits.size)(ary, bits) return bits # From 0c2de3c536bbc1df88f2e567cf0cb2683201a8fc Mon Sep 17 00:00:00 2001 From: Keith Kraus Date: Thu, 20 Sep 2018 13:48:10 -0400 Subject: [PATCH 4/6] Another place to prevent a zero size kernel --- pygdf/cudautils.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/pygdf/cudautils.py b/pygdf/cudautils.py index edf9f41caa34..3b21745e6036 100755 --- a/pygdf/cudautils.py +++ b/pygdf/cudautils.py @@ -174,7 +174,8 @@ def gpu_fill_value(data, value): def fill_value(arr, value): """Fill *arr* with value """ - gpu_fill_value.forall(arr.size)(arr, value) + if arr.size > 0: + gpu_fill_value.forall(arr.size)(arr, value) @cuda.jit From a6678f6a65b5c6edb25ec0abcacea3661b448151 Mon Sep 17 00:00:00 2001 From: Keith Kraus Date: Thu, 20 Sep 2018 17:52:22 -0400 Subject: [PATCH 5/6] Don't remove valid mask if the null_count is 0 --- pygdf/column.py | 5 ----- 1 file changed, 5 deletions(-) diff --git a/pygdf/column.py b/pygdf/column.py index d1305ae29011..ce05a8138c88 100644 --- a/pygdf/column.py +++ b/pygdf/column.py @@ -98,15 +98,10 @@ def __init__(self, data, mask=None, null_count=None): nnz = cudautils.count_nonzero_mask(self._mask.mem, size=len(self)) null_count = len(self) - nnz - if null_count == 0: - self._mask = None else: null_count = 0 assert 0 <= null_count <= len(self) - if null_count == 0: - # Remove mask if null_count is zero - self._mask = None self._null_count = null_count From 13cdda1e58ec3b0aa52e77b61bf281014b91f3af Mon Sep 17 00:00:00 2001 From: Keith Kraus Date: Thu, 20 Sep 2018 18:06:04 -0400 Subject: [PATCH 6/6] Fix zero sized kernel on fillna --- pygdf/cudautils.py | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/pygdf/cudautils.py b/pygdf/cudautils.py index 3b21745e6036..a3478cf0aa6a 100755 --- a/pygdf/cudautils.py +++ b/pygdf/cudautils.py @@ -374,9 +374,10 @@ def gpu_fill_masked(value, validity, out): def fillna(data, mask, value): out = cuda.device_array_like(data) - out.copy_to_device(data) - configured = gpu_fill_masked.forall(data.size) - configured(value, mask, out) + if out.size > 0: + out.copy_to_device(data) + configured = gpu_fill_masked.forall(data.size) + configured(value, mask, out) return out