 					break;
 			}
 
-				// // set the valid bitmap - all bits were set to 0 to start
-				int bitmapIdx 	= whichBitmap(rec_id);  	// which bitmap
-				int bitIdx		= whichBit(rec_id);		// which bit - over an 8-bit index
-				setBit(valid[col]+bitmapIdx, bitIdx);		// This is done with atomics
-
+			// set the valid bitmap - all bits were set to 0 to start
+			int bitmapIdx 	= whichBitmap(rec_id);  	// which bitmap
+			int bitIdx		= whichBit(rec_id);		// which bit - over an 8-bit index
+			setBit(valid[col]+bitmapIdx, bitIdx);		// This is done with atomics
+
+			atomicAdd((unsigned long long int*)&num_valid[col],(unsigned long long int)1);
 		}
 		else if(dtype[col]==gdf_dtype::GDF_STRING){
 			str_cols[stringCol][rec_id].first 	= NULL;

	long num_bitmaps = (N + 31) / 8;			// 8 bytes per bitmap

	//--- allocate space for the valid bitmaps
	CUDA_TRY(cudaMallocManaged(&gdf->valid, (sizeof(gdf_valid_type) 	* num_bitmaps)));
	CUDA_TRY(cudaMemset(gdf->valid, 0, (sizeof(gdf_valid_type) 	* num_bitmaps)) );

	//--- Allocate space for the data
	switch(gdf->dtype) {
		case gdf_dtype::GDF_INT8:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(int8_t) * N)));
			break;
		case gdf_dtype::GDF_INT16:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(int16_t) * N)));
			break;
		case gdf_dtype::GDF_INT32:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(int32_t) * N)));
			break;
		case gdf_dtype::GDF_INT64:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(int64_t) * N)));
			break;
		case gdf_dtype::GDF_FLOAT32:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(float) * N)));
			break;
		case gdf_dtype::GDF_FLOAT64:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(double) * N)));
			break;
		case gdf_dtype::GDF_DATE32:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(gdf_date32) * N)));
			break;
		case gdf_dtype::GDF_DATE64:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(gdf_date64) * N)));
			break;
		case gdf_dtype::GDF_TIMESTAMP:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(int64_t) * N)));
			break;
		case gdf_dtype::GDF_CATEGORY:
			CUDA_TRY(cudaMallocManaged(&gdf->data, (sizeof(gdf_category) * N)));
			break;
		case gdf_dtype::GDF_STRING:
			// Memory for gdf->data allocated by string class eventually

	string_pair** str_cols = NULL;

	if (stringColCount > 0 ) {
		CUDA_TRY( cudaMalloc ((void**)&str_cols, 	(sizeof(string_pair *)		* stringColCount)) );

		for (int col = 0; col < stringColCount; col++) {
			CUDA_TRY( cudaMalloc ((void**)(str_cols + col), sizeof(string_pair) * (raw_csv->num_records)) );
		}
	}

