OmniSciDB  3a86f6ec37
TableFunctions.hpp
Go to the documentation of this file.
1 /*
2  * Copyright 2019 OmniSci, Inc.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
17 #include "../../QueryEngine/OmniSciTypes.h"
18 #include "../../Shared/funcannotations.h"
19 
20 #define EXTENSION_INLINE extern "C" ALWAYS_INLINE DEVICE
21 #define EXTENSION_NOINLINE extern "C" NEVER_INLINE DEVICE
22 
23 EXTENSION_NOINLINE int32_t row_copier(const Column<double>& input_col,
24  int copy_multiplier,
25  Column<double>& output_col) {
26  int32_t output_row_count = copy_multiplier * input_col.getSize();
27  if (output_row_count > 100) {
28  // Test failure propagation.
29  return -1;
30  }
31  if (output_col.getSize() != output_row_count) {
32  return -1;
33  }
34 
35 #ifdef __CUDACC__
36  int32_t start = threadIdx.x + blockDim.x * blockIdx.x;
37  int32_t stop = static_cast<int32_t>(input_col.getSize());
38  int32_t step = blockDim.x * gridDim.x;
39 #else
40  auto start = 0;
41  auto stop = input_col.getSize();
42  auto step = 1;
43 #endif
44 
45  for (auto i = start; i < stop; i += step) {
46  for (int c = 0; c < copy_multiplier; c++) {
47  output_col[i + (c * input_col.getSize())] = input_col[i];
48  }
49  }
50 
51  return output_row_count;
52 }
53 
54 EXTENSION_NOINLINE int32_t row_adder(const int copy_multiplier,
55  const Column<double>& input_col1,
56  const Column<double>& input_col2,
57  Column<double>& output_col) {
58  int32_t output_row_count = copy_multiplier * input_col1.getSize();
59  if (output_row_count > 100) {
60  // Test failure propagation.
61  return -1;
62  }
63  if (output_col.getSize() != output_row_count) {
64  return -1;
65  }
66 
67 #ifdef __CUDACC__
68  int32_t start = threadIdx.x + blockDim.x * blockIdx.x;
69  int32_t stop = static_cast<int32_t>(input_col1.getSize());
70  int32_t step = blockDim.x * gridDim.x;
71 #else
72  auto start = 0;
73  auto stop = input_col1.getSize();
74  auto step = 1;
75 #endif
76  auto stride = input_col1.getSize();
77  for (auto i = start; i < stop; i += step) {
78  for (int c = 0; c < copy_multiplier; c++) {
79  if (input_col1.isNull(i) || input_col2.isNull(i)) {
80  output_col.setNull(i + (c * stride));
81  } else {
82  output_col[i + (c * stride)] = input_col1[i] + input_col2[i];
83  }
84  }
85  }
86 
87  return output_row_count;
88 }
89 
90 EXTENSION_NOINLINE int32_t row_addsub(const int copy_multiplier,
91  const Column<double>& input_col1,
92  const Column<double>& input_col2,
93  Column<double>& output_col1,
94  Column<double>& output_col2) {
95  int32_t output_row_count = copy_multiplier * input_col1.getSize();
96  if (output_row_count > 100) {
97  // Test failure propagation.
98  return -1;
99  }
100  if ((output_col1.getSize() != output_row_count) ||
101  (output_col2.getSize() != output_row_count)) {
102  return -1;
103  }
104 
105 #ifdef __CUDACC__
106  int32_t start = threadIdx.x + blockDim.x * blockIdx.x;
107  int32_t stop = static_cast<int32_t>(input_col1.getSize());
108  int32_t step = blockDim.x * gridDim.x;
109 #else
110  auto start = 0;
111  auto stop = input_col1.getSize();
112  auto step = 1;
113 #endif
114  auto stride = input_col1.getSize();
115  for (auto i = start; i < stop; i += step) {
116  for (int c = 0; c < copy_multiplier; c++) {
117  output_col1[i + (c * stride)] = input_col1[i] + input_col2[i];
118  output_col2[i + (c * stride)] = input_col1[i] - input_col2[i];
119  }
120  }
121  return output_row_count;
122 }
123 
125  Column<int>& output_max_col,
126  Column<int>& output_max_row_col) {
127  if ((output_max_col.getSize() != 1) || output_max_row_col.getSize() != 1) {
128  return -1;
129  }
130 #ifdef __CUDACC__
131  int32_t start = threadIdx.x + blockDim.x * blockIdx.x;
132  int32_t stop = static_cast<int32_t>(input_col.getSize());
133  int32_t step = blockDim.x * gridDim.x;
134 #else
135  auto start = 0;
136  auto stop = input_col.getSize();
137  auto step = 1;
138 #endif
139 
140  int curr_max = -2147483648;
141  int curr_max_row = -1;
142  for (auto i = start; i < stop; i += step) {
143  if (input_col[i] > curr_max) {
144  curr_max = input_col[i];
145  curr_max_row = i;
146  }
147  }
148  output_max_col[0] = curr_max;
149  output_max_row_col[0] = curr_max_row;
150  return 1;
151 }
152 
153 #include "TableFunctionsTesting.hpp"
EXTENSION_NOINLINE int32_t row_adder(const int copy_multiplier, const Column< double > &input_col1, const Column< double > &input_col2, Column< double > &output_col)
EXTENSION_NOINLINE int32_t row_addsub(const int copy_multiplier, const Column< double > &input_col1, const Column< double > &input_col2, Column< double > &output_col1, Column< double > &output_col2)
DEVICE bool isNull(int64_t index) const
Definition: OmniSciTypes.h:157
EXTENSION_NOINLINE int32_t get_max_with_row_offset(const Column< int > &input_col, Column< int > &output_max_col, Column< int > &output_max_row_col)
EXTENSION_NOINLINE int32_t row_copier(const Column< double > &input_col, int copy_multiplier, Column< double > &output_col)
DEVICE void setNull(int64_t index)
Definition: OmniSciTypes.h:158
DEVICE int64_t getSize() const
Definition: OmniSciTypes.h:154
#define EXTENSION_NOINLINE