[REVIEW] Adding support for is_null and is_not_null

CHANGELOG.md

@@ -4,7 +4,6 @@
 
 - PR #2962 Adding support for isNULL and isNotNULL
 
-
 ## Improvements
 
 - PR #2904 Move gpu decompressors to cudf::io namespace

cpp/src/unary/null_ops.cu

@@ -35,29 +35,29 @@ gdf_column null_op(gdf_column const& input, bool nulls_are_false = true, cudaStr
     if (not cudf::is_nullable(input)) {
 	gdf_scalar value {nulls_are_false, GDF_BOOL8, true}; 
 	cudf::fill(&output, value, 0, output.size);
-    }
-    else {
+    } else {
         const bit_mask_t* __restrict__ typed_input_valid = reinterpret_cast<bit_mask_t*>(input.valid);
         auto exec = rmm::exec_policy(stream)->on(stream);
 
-	thrust::transform(exec,
-			  thrust::make_counting_iterator(static_cast<gdf_size_type>(0)),
+        thrust::transform(exec,
+                          thrust::make_counting_iterator(static_cast<gdf_size_type>(0)),
                           thrust::make_counting_iterator(static_cast<gdf_size_type>(input.size)),
-			  static_cast<bool*>(output.data),
-			  [=]__device__(auto index){return (nulls_are_false == bit_mask::is_valid(typed_input_valid, index));});
+                          static_cast<bool*>(output.data),
+                          [=]__device__(auto index){
+                              return (nulls_are_false ==
+                                      bit_mask::is_valid(typed_input_valid, index));
+                          });
     }
 
     return output;
 }
 }// detail
 
 gdf_column is_null(gdf_column const& input) {
-
     return detail::null_op(input, false, 0);
 }
 
 gdf_column is_not_null(gdf_column const& input) {
-
     return detail::null_op(input, true, 0);
 }
 

python/cudf/cudf/core/series.py

@@ -1278,15 +1278,12 @@ def to_array(self, fillna=None):
     def isnull(self):
         """Identify missing values in a Series.
         """
-        if not self.has_null_mask:
-            return Series(
-                cudautils.zeros(len(self), np.bool_),
-                name=self.name,
-                index=self.index,
-            )
-
-        mask = cudautils.isnull_mask(self.data, self.nullmask.mem)
-        return Series(mask, name=self.name, index=self.index)
+        result = Series(
+            libcudf.unaryops.isNULL(self._column),
+            name=self.name,
+            index=self.index,
+        )
+        return result
 
     def isna(self):
         """Identify missing values in a Series. Alias for isnull.
@@ -1296,15 +1293,12 @@ def isna(self):
     def notna(self):
         """Identify non-missing values in a Series.
         """
-        if not self.has_null_mask:
-            return Series(
-                cudautils.ones(len(self), np.bool_),
-                name=self.name,
-                index=self.index,
-            )
-
-        mask = cudautils.notna_mask(self.data, self.nullmask.mem)
-        return Series(mask, name=self.name, index=self.index)
+        result = Series(
+            libcudf.unaryops.isNotNULL(self._column),
+            name=self.name,
+            index=self.index,
+        )
+        return result
 
     def notnull(self):
         """Identify non-missing values in a Series. Alias for notna.