Skip to content

T
Tensor.LowPrecision
Commit 3e2f47b2 by ltb
Options

merge ConvertDataType(float/float16/int/int8) and modify XName

parent a73f8e42
正在显示 包含 106 行增加 和 9 行删除
source/tensor/XName.cpp
...@@ -108,9 +108,15 @@ const char * GetOPName(int type) ...@@ -108,9 +108,15 @@ const char * GetOPName(int type)
else if (type == REDUCE_REDUCEVARIANCE) else if (type == REDUCE_REDUCEVARIANCE)
return "R_REDUCEVARIANCE"; return "R_REDUCEVARIANCE";
} }
else if ((type & DATA_BASE) != 0){ else if ((type & DATA_BASE) != 0) {
if (type == GETANDSET_SELECT) if (type == GETANDSET_CONVERTDATATYPE)
return "G_SELECT"; return "G_CONVERTDATATYPE";
else if (type == GETANDSET_INDEXTOONEHOT)
return "G_INDEXTOONEHOT";
else if (type == GETANDSET_ONEHOTTOINDEX)
return "G_ONEHOTTOINDEX";
else if (type == GETANDSET_SELECT)
return "G_SELECT";
else if (type == MOVEMENT_COPYINDEXED) else if (type == MOVEMENT_COPYINDEXED)
return "M_COPYINDEXED"; return "M_COPYINDEXED";
else if (type == MOVEMENT_COPYVALUES) else if (type == MOVEMENT_COPYVALUES)
......
source/tensor/XName.h
...@@ -79,9 +79,13 @@ namespace nts { // namespace nts(NiuTrans.Tensor) ...@@ -79,9 +79,13 @@ namespace nts { // namespace nts(NiuTrans.Tensor)
/* data and shape related operations */ /* data and shape related operations */
#define DATA_BASE MATH_BASE * 2 #define DATA_BASE MATH_BASE * 2
#define GETANDSET DATA_BASE + 1 #define GETANDSET DATA_BASE + 1
#define GETANDSET_SELECT GETANDSET + 1 #define GETANDSET_CONVERTDATATYPE GETANDSET + 1
#define GETANDSET_INDEXTOONEHOT GETANDSET_CONVERTDATATYPE + 1
#define GETANDSET_ONEHOTTOINDEX GETANDSET_INDEXTOONEHOT + 1
#define GETANDSET_SELECT GETANDSET_ONEHOTTOINDEX + 1
#define MOVEMENT GETANDSET_SELECT + 1 #define SHAPE_BASE DATA_BASE * 2
#define MOVEMENT SHAPE_BASE + 1
#define MOVEMENT_COPYINDEXED MOVEMENT + 1 #define MOVEMENT_COPYINDEXED MOVEMENT + 1
#define MOVEMENT_COPYVALUES MOVEMENT_COPYINDEXED + 1 #define MOVEMENT_COPYVALUES MOVEMENT_COPYINDEXED + 1
#define MOVEMENT_GATHER MOVEMENT_COPYVALUES + 1 #define MOVEMENT_GATHER MOVEMENT_COPYVALUES + 1
...@@ -104,7 +108,7 @@ namespace nts { // namespace nts(NiuTrans.Tensor) ...@@ -104,7 +108,7 @@ namespace nts { // namespace nts(NiuTrans.Tensor)
#define SORT_TOPK SORT_SORT + 1 #define SORT_TOPK SORT_SORT + 1
/* activation functions */ /* activation functions */
#define FUNCTION_BASE DATA_BASE * 2 #define FUNCTION_BASE SHAPE_BASE * 2
#define FUNC_DROPOUT FUNCTION_BASE + 1 #define FUNC_DROPOUT FUNCTION_BASE + 1
#define FUNC_HARDTANH FUNC_DROPOUT + 1 #define FUNC_HARDTANH FUNC_DROPOUT + 1
#define FUNC_IDENTITY FUNC_HARDTANH + 1 #define FUNC_IDENTITY FUNC_HARDTANH + 1
......
source/tensor/core/getandset/ConvertDataType.cpp
...@@ -20,20 +20,21 @@ ...@@ -20,20 +20,21 @@
*/ */
#include "../../XTensor.h" #include "../../XTensor.h"
#include "../../XName.h"
#include "ConvertDataType.h" #include "ConvertDataType.h"
#include "ConvertDataType.cuh" #include "ConvertDataType.cuh"
#include "../movement/CopyValues.h"
namespace nts { // namespace nts(NiuTrans.Tensor) namespace nts { // namespace nts(NiuTrans.Tensor)
/* /*
convert data type convert data type
>> input - input tensor >> input - input tensor
>> output - output tensor >> output - output tensor
*/ */
void _ConvertDataType(const XTensor * input, XTensor * output) void _ConvertDataType(const XTensor * input, XTensor * output)
{ {
//CheckNTErrors((input->unitSize == output->unitSize), "Input and Output must be same in size!");
if (input->dataType == output->dataType) if (input->dataType == output->dataType)
return; return;
...@@ -61,4 +62,35 @@ void _ConvertDataType(const XTensor * input, XTensor * output) ...@@ -61,4 +62,35 @@ void _ConvertDataType(const XTensor * input, XTensor * output)
ShowNTErrors("Unsupported data types for conversion!"); ShowNTErrors("Unsupported data types for conversion!");
} }
/*
convert data type (return an XTensor structure)
make a new tensor to keep the result and return it
>> input - input tensor
<< return - output tensor with the specified data type
*/
XTensor ConvertDataType(const XTensor & input, TENSOR_DATA_TYPE dataType)
{
if (input.dataType == dataType) {
XTensor output;
output = CopyValues(input);
return output;
}
int order = input.order;
float dr = (!input.isSparse) ? 1.0F : input.denseRatio;
XTensor output(order, input.dimSize, dataType, dr, input.devID, input.mem);
output.SetTMPFlag();
_ConvertDataType(&input, &output);
/* tensor connection */
XLink::MakeLink(&input, NULL, &output, GETANDSET_CONVERTDATATYPE);
return output;
}
} // namespace nts(NiuTrans.Tensor) } // namespace nts(NiuTrans.Tensor)
\ No newline at end of file
source/tensor/core/getandset/ConvertDataType.cu
...@@ -67,7 +67,49 @@ void KernelIntToFloat(int * inputData, float * outputData, int size) ...@@ -67,7 +67,49 @@ void KernelIntToFloat(int * inputData, float * outputData, int size)
if (i < size){ if (i < size){
outputData[i] = (float)(inputData[i]); outputData[i] = (float)(inputData[i]);
}} }
}
__global__
void KernelFloatToInt8(float * inputData, __int8 * outputData, int size)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < size) {
outputData[i] = (__int8)(inputData[i]);
}
}
__global__
void KernelInt8ToFloat(__int8 * inputData, float * outputData, int size)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < size) {
outputData[i] = (float)(inputData[i]);
}
}
__global__
void KernelIntToInt8(int * inputData, __int8 * outputData, int size)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < size) {
outputData[i] = (__int8)(inputData[i]);
}
}
__global__
void KernelInt8ToInt(__int8 * inputData, int * outputData, int size)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < size) {
outputData[i] = (int)(inputData[i]);
}
}
/* /*
data conversion (cuda code) data conversion (cuda code)
...@@ -138,6 +180,14 @@ void _CudaConvertDataType(const XTensor * input, XTensor * output) ...@@ -138,6 +180,14 @@ void _CudaConvertDataType(const XTensor * input, XTensor * output)
KernelFloatToFloat16<<<blocks, threads>>>((float*)input->data, (__half*)output->data, input->unitNum); KernelFloatToFloat16<<<blocks, threads>>>((float*)input->data, (__half*)output->data, input->unitNum);
else if(input->dataType == X_FLOAT16 && output->dataType == X_FLOAT) else if(input->dataType == X_FLOAT16 && output->dataType == X_FLOAT)
KernelFloat16ToFloat<<<blocks, threads>>>((__half*)input->data, (float*)output->data, input->unitNum); KernelFloat16ToFloat<<<blocks, threads>>>((__half*)input->data, (float*)output->data, input->unitNum);
else if (input->dataType == X_FLOAT && output->dataType == X_INT8)
KernelFloatToInt8 << <blocks, threads >> >((float*)input->data, (__int8*)output->data, input->unitNum);
else if (input->dataType == X_INT8 && output->dataType == X_FLOAT)
KernelInt8ToFloat << <blocks, threads >> >((__int8*)input->data, (float*)output->data, input->unitNum);
else if (input->dataType == X_INT && output->dataType == X_INT8)
KernelIntToInt8 << <blocks, threads >> >((int*)input->data, (__int8*)output->data, input->unitNum);
else if (input->dataType == X_INT8 && output->dataType == X_INT)
KernelInt8ToInt << <blocks, threads >> >((__int8*)input->data, (int*)output->data, input->unitNum);
else{ else{
ShowNTErrors("Unsupported data types for conversion!"); ShowNTErrors("Unsupported data types for conversion!");
} }
......
source/tensor/core/getandset/ConvertDataType.h
...@@ -23,12 +23,16 @@ ...@@ -23,12 +23,16 @@
#define __CONVERTDATATYPE_H__ #define __CONVERTDATATYPE_H__
#include "../../XTensor.h" #include "../../XTensor.h"
#include "../../XDataType.h"
namespace nts { // namespace nts(NiuTrans.Tensor) namespace nts { // namespace nts(NiuTrans.Tensor)
/* convert data type */ /* convert data type */
void _ConvertDataType(const XTensor * input, XTensor * output); void _ConvertDataType(const XTensor * input, XTensor * output);
/* convert data type (return an XTensor structure) */
XTensor ConvertDataType(const XTensor & input, TENSOR_DATA_TYPE dataType);
} // namespace nts(NiuTrans.Tensor) } // namespace nts(NiuTrans.Tensor)
#endif // __CONVERTDATATYPE_H__ #endif // __CONVERTDATATYPE_H__
Markdown 格式
0% 或
您添加了 0 人 到此讨论。请谨慎行事。
请先完成此评论的编辑！
请 注册 或者 后发表评论