Duli/clip cuda #1677
Duli/clip cuda #1677
Conversation
| template <typename T> | ||
| __global__ void _Clip(const T* input, T* output, T min, T max, size_t N) { | ||
| CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N); | ||
| output[id] = (input[id] < min ? min : input[id]) > max ? max : input[id]; |
There was a problem hiding this comment.
Looks like it is not correct if input[id]<min. For example, input[id]=-1, min = 0, and max = 10, the result is -1, which is not expected.
There was a problem hiding this comment.
I've fixed this logic error. Please see the latest commit of this branch.
| REGISTER_KERNEL_TYPED(T) \ | ||
| template Status Clip<T>::ComputeInternal(OpKernelContext* ctx) const; | ||
|
|
||
| SPECIALIZED_COMPUTE(float) |
There was a problem hiding this comment.
why float only? version 6 also have double and float16, right?
why not support opset 11 also?
There was a problem hiding this comment.
Let's add float16 and double
There was a problem hiding this comment.
While using GetAttrOrDefault() to get attribute for minimum/maximum values, it seems that ONNX has not yet supported double and float16 type as listed in AttributeProto_AttributeType. Once ONNX supports double and float16, we can add them as well.
* Moreover, properly handle default input situation
Description: Describe your changes.
Clip cuda kernel