[REVIEW] Jitify versions of binaryops for non-homogeneous types #892
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -85,6 +85,17 @@ typedef struct { | |
| // here we can also hold info for decimal datatype or any other datatype that requires additional information | ||
| } gdf_dtype_extra_info; | ||
|
|
||
|
|
||
| /**---------------------------------------------------------------------------* | ||
| * @brief A struct to hold a scalar (single) value and its type information | ||
| * | ||
| *---------------------------------------------------------------------------**/ | ||
| typedef struct { | ||
| void* data; /**< Pointer to the scalar data */ | ||
|
||
| gdf_dtype dtype; /**< The datatype of the scalar's data */ | ||
| } gdf_scalar; | ||
|
|
||
|
|
||
| typedef struct gdf_column_{ | ||
| void *data; /**< Pointer to the columns data */ | ||
| gdf_valid_type *valid; /**< Pointer to the columns validity bit mask where the 'i'th bit indicates if the 'i'th row is NULL */ | ||
|
|
@@ -156,6 +167,29 @@ typedef enum { | |
| GDF_NUM_COLORS, /** Add new colors above this line */ | ||
| } gdf_color; | ||
|
|
||
|
|
||
| /**---------------------------------------------------------------------------* | ||
| * @brief Types of binary operations that can be performed on data. | ||
| * | ||
| *---------------------------------------------------------------------------**/ | ||
| typedef enum { | ||
| GDF_ADD, /**< operator + */ | ||
| GDF_SUB, /**< operator - */ | ||
| GDF_MUL, /**< operator * */ | ||
| GDF_DIV, /**< operator / using common type of lhs and rhs */ | ||
| GDF_TRUE_DIV, /**< operator / after promoting type to floating point*/ | ||
kkraus14 marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| GDF_FLOOR_DIV, /**< operator / after promoting to float and then flooring the result */ | ||
| GDF_MOD, /**< operator % */ | ||
| GDF_POW, /**< lhs ^ rhs */ | ||
| GDF_EQUAL, /**< operator == */ | ||
| GDF_NOT_EQUAL, /**< operator != */ | ||
| GDF_LESS, /**< operator < */ | ||
| GDF_GREATER, /**< operator > */ | ||
| GDF_LESS_EQUAL, /**< operator <= */ | ||
| GDF_GREATER_EQUAL, /**< operator >= */ | ||
| } gdf_binary_operator; | ||
|
|
||
|
|
||
| /* --------------------------------------------------------------------------*/ | ||
| /** | ||
| * @brief This struct holds various information about how an operation should be | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,35 @@ | ||
| /* | ||
| * Copyright 2018-2019 BlazingDB, Inc. | ||
| * Copyright 2018 Christian Noboa Mardini <[email protected]> | ||
| * | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ | ||
|
|
||
| #ifndef GDF_BINARY_OPERATION_JIT_CODE_CODE_H | ||
| #define GDF_BINARY_OPERATION_JIT_CODE_CODE_H | ||
|
|
||
| namespace cudf { | ||
| namespace binops { | ||
| namespace jit { | ||
| namespace code { | ||
|
|
||
| extern const char* kernel; | ||
| extern const char* traits; | ||
| extern const char* operation; | ||
|
|
||
| } | ||
| } | ||
| } | ||
| } | ||
|
|
||
| #endif |
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,67 @@ | ||
| /* | ||
| * Copyright 2018-2019 BlazingDB, Inc. | ||
| * Copyright 2018 Christian Noboa Mardini <[email protected]> | ||
| * Copyright 2018 Rommel Quintanilla <[email protected]> | ||
| * | ||
| * Licensed under the Apache License, Version 2.0 (the "License"); | ||
| * you may not use this file except in compliance with the License. | ||
| * You may obtain a copy of the License at | ||
| * | ||
| * http://www.apache.org/licenses/LICENSE-2.0 | ||
| * | ||
| * Unless required by applicable law or agreed to in writing, software | ||
| * distributed under the License is distributed on an "AS IS" BASIS, | ||
| * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| * See the License for the specific language governing permissions and | ||
| * limitations under the License. | ||
| */ | ||
|
|
||
| namespace cudf { | ||
| namespace binops { | ||
| namespace jit { | ||
| namespace code { | ||
|
|
||
| const char* kernel = | ||
| R"***( | ||
| #include "operation.h" | ||
| #include "cudf/types.h" | ||
|
|
||
| template <typename TypeOut, typename TypeLhs, typename TypeRhs, typename TypeOpe> | ||
| __global__ | ||
| void kernel_v_s(gdf_size_type size, | ||
| TypeOut* out_data, TypeLhs* lhs_data, TypeRhs* rhs_data) { | ||
| int tid = threadIdx.x; | ||
| int blkid = blockIdx.x; | ||
| int blksz = blockDim.x; | ||
| int gridsz = gridDim.x; | ||
|
|
||
| int start = tid + blkid * blksz; | ||
| int step = blksz * gridsz; | ||
|
|
||
| for (gdf_size_type i=start; i<size; i+=step) { | ||
| out_data[i] = TypeOpe::template operate<TypeOut, TypeLhs, TypeRhs>(lhs_data[i], rhs_data[0]); | ||
| } | ||
| } | ||
|
|
||
| template <typename TypeOut, typename TypeLhs, typename TypeRhs, typename TypeOpe> | ||
| __global__ | ||
| void kernel_v_v(gdf_size_type size, | ||
| TypeOut* out_data, TypeLhs* lhs_data, TypeRhs* rhs_data) { | ||
| int tid = threadIdx.x; | ||
| int blkid = blockIdx.x; | ||
| int blksz = blockDim.x; | ||
| int gridsz = gridDim.x; | ||
|
|
||
| int start = tid + blkid * blksz; | ||
| int step = blksz * gridsz; | ||
|
|
||
| for (gdf_size_type i=start; i<size; i+=step) { | ||
| out_data[i] = TypeOpe::template operate<TypeOut, TypeLhs, TypeRhs>(lhs_data[i], rhs_data[i]); | ||
| } | ||
| } | ||
| )***"; | ||
|
|
||
| } // namespace code | ||
| } // namespace jit | ||
| } // namespace binops | ||
| } // namespace cudf |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Should we pin this to a specific branch / tag / commit?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yes, please pin to a commit to avoid build issues in the future.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Will this do? Or do I need to specify it in
.gitmodules.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
👍
Uh oh!
There was an error while loading. Please reload this page.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
We need to specify it in
.gitmodulesotherwise agit submodule update --init --recursive --remotewill update it to the latest commit of master.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I can't seem to find a way to pin this to a commit/tag in the
.gitmodulesfile. I read this andupdateseems to be the only option. But details on update say that the only way to do this is to usenone. Is that what I should do?