Skip to content

op:switch&where验证完成 #55

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
miaobyte merged 2 commits into from
May 14, 2025
Merged

op:switch&where验证完成 #55

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
2e76f5e
op:switch&where 待验证
miaobyte May 12, 2025
f8ab5b8
op:switch&where验证完成
miaobyte May 14, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion doc/excuter/op-mem-cuda/list.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,7 @@

| Operation | Author | Math Formula | IR Instruction |
|-----------|--------|--------------|----------------|
| switch | miaobyte | C=switch(tensors,cases) | switch(listtensor<any> tensors, tensor<int8> cases)->(tensor<any> result) |
| switch | miaobyte | C=switch(tensors,cases) | switch(listtensor<any> tensors, tensor<int32|bool> cases)->(tensor<any> result) |
| greaterscalar | miaobyte | mask=compare(T1, scalar) | greaterscalar(tensor<any> A, var<any> scalar)->(tensor<bool> mask) |
| notequal | miaobyte | T1!=T2->mask | notequal(tensor<any> A, tensor<any> B, var<float32> epsilon)->(tensor<bool> mask) |
| equalscalar | miaobyte | T1==scalar->mask | equalscalar(tensor<any> A, var<any> scalar, var<float32> epsilon)->(tensor<bool> mask) |
Expand Down
4 changes: 2 additions & 2 deletions doc/excuter/op-mem-ompsimd/list.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@

| Operation | Author | Math Formula | IR Instruction |
|-----------|--------|--------------|----------------|
| switch | miaobyte | C=switch([tensors],case) | switch(listtensor<any> tensors, tensor<int8> cases)->(tensor<any> C) |
| switch | miaobyte | C=switch([tensors],case) | switch(listtensor<any> tensors, tensor<int32|bool> cases)->(tensor<any> C) |
| greaterscalar | miaobyte | mask=greater(T1,scalar) | greaterscalar(tensor<any> A, var<any> scalar)->(tensor<bool> mask) |
| notequal | miaobyte | notequal(T1,T2)->mask | notequal(tensor<any> A, tensor<any> B, var<float32> epsilon)->(tensor<bool> mask) |
| equalscalar | miaobyte | mask=equal(T1,scalar) | equalscalar(tensor<any> A, var<any> scalar, var<float32> eposilon)->(tensor<bool> mask) |
Expand All @@ -63,7 +63,7 @@
| lessscalar | miaobyte | mask=less(T1,scalar) | lessscalar(tensor<any> A, var<any> scalar)->(tensor<bool> mask) |
| notequalscalar | miaobyte | mask=notequal(T1,scalar) | notequalscalar(tensor<any> A, var<any> scalar, var<float32> epsilon)->(tensor<bool> mask) |
| minscalar | miaobyte | T3=min(T1,scalar) | minscalar(tensor<any> A, var<any> scalar)->(tensor<any> C) |
| rpowscalar | miaobyte | T3=scalar^T1 | rpowscalar(var<any> scalar, tensor<any> A)->(tensor<any> C) |
| rpowscalar | miaobyte | T3=scalar^T1 | rpowscalar(var<float32> scalar, tensor<any> A)->(tensor<any> C) |
| rdivscalar | miaobyte | T3=scalar/T1 | rdivscalar(var<any> scalar, tensor<any> A)->(tensor<any> C) |
| less | miaobyte | mask=less(T1,T2) | less(tensor<any> A, tensor<any> B)->(tensor<bool> mask) |
| powscalar | miaobyte | T3=T1^scalar | powscalar(tensor<any> A, var<any> scalar)->(tensor<any> C) |
Expand Down
2 changes: 1 addition & 1 deletion excuter/op-mem-cuda/src/client/tfs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -486,7 +486,7 @@ namespace deepx::tf
tffactory.add_tf(std::make_shared<Switch<miaobyte>>(vector<Param>(
{
Param("tensors", DataCategory::ListTensor, Precision::Any),
Param("cases", DataCategory::Tensor, Precision::Int8),
Param("cases", DataCategory::Tensor, Precision::Int32|Precision::Bool),
}),
vector<Param>(
{
Expand Down
28 changes: 19 additions & 9 deletions excuter/op-mem-cuda/src/deepx/tensorfunc/elementwise_miaobyte_compare.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -524,15 +524,25 @@ namespace deepx::tensorfunc
}
}

template void launch_switch<double,int8_t>(const double **tensorsdata, const int numTensors, const int8_t *cases, double *C, const int size);
template void launch_switch<float,int8_t>(const float **tensorsdata, const int numTensors, const int8_t *cases, float *C, const int size);
template void launch_switch<nv_bfloat16,int8_t>(const nv_bfloat16 **tensorsdata, const int numTensors, const int8_t *cases, nv_bfloat16 *C, const int size);
template void launch_switch<__half,int8_t>(const __half **tensorsdata, const int numTensors, const int8_t *cases, __half *C, const int size);
template void launch_switch<int64_t,int8_t>(const int64_t **tensorsdata, const int numTensors, const int8_t *cases, int64_t *C, const int size);
template void launch_switch<int32_t,int8_t>(const int32_t **tensorsdata, const int numTensors, const int8_t *cases, int32_t *C, const int size);
template void launch_switch<int16_t,int8_t>(const int16_t **tensorsdata, const int numTensors, const int8_t *cases, int16_t *C, const int size);
template void launch_switch<int8_t,int8_t>(const int8_t **tensorsdata, const int numTensors, const int8_t *cases, int8_t *C, const int size);
template void launch_switch<bool,int8_t>(const bool **tensorsdata, const int numTensors, const int8_t *cases, bool *C, const int size);
template void launch_switch<double,int32_t>(const double **tensorsdata, const int numTensors, const int32_t *cases, double *C, const int size);
template void launch_switch<float,int32_t>(const float **tensorsdata, const int numTensors, const int32_t *cases, float *C, const int size);
template void launch_switch<nv_bfloat16,int32_t>(const nv_bfloat16 **tensorsdata, const int numTensors, const int32_t *cases, nv_bfloat16 *C, const int size);
template void launch_switch<__half,int32_t>(const __half **tensorsdata, const int numTensors, const int32_t *cases, __half *C, const int size);
template void launch_switch<int64_t,int32_t>(const int64_t **tensorsdata, const int numTensors, const int32_t *cases, int64_t *C, const int size);
template void launch_switch<int32_t,int32_t>(const int32_t **tensorsdata, const int numTensors, const int32_t *cases, int32_t *C, const int size);
template void launch_switch<int16_t,int32_t>(const int16_t **tensorsdata, const int numTensors, const int32_t *cases, int16_t *C, const int size);
template void launch_switch<int8_t,int32_t>(const int8_t **tensorsdata, const int numTensors, const int32_t *cases, int8_t *C, const int size);
template void launch_switch<bool,int32_t>(const bool **tensorsdata, const int numTensors, const int32_t *cases, bool *C, const int size);

template void launch_switch<double,bool>(const double **tensorsdata, const int numTensors, const bool *cases, double *C, const int size);
template void launch_switch<float,bool>(const float **tensorsdata, const int numTensors, const bool *cases, float *C, const int size);
template void launch_switch<nv_bfloat16,bool>(const nv_bfloat16 **tensorsdata, const int numTensors, const bool *cases, nv_bfloat16 *C, const int size);
template void launch_switch<__half,bool>(const __half **tensorsdata, const int numTensors, const bool *cases, __half *C, const int size);
template void launch_switch<int64_t,bool>(const int64_t **tensorsdata, const int numTensors, const bool *cases, int64_t *C, const int size);
template void launch_switch<int32_t,bool>(const int32_t **tensorsdata, const int numTensors, const bool *cases, int32_t *C, const int size);
template void launch_switch<int16_t,bool>(const int16_t **tensorsdata, const int numTensors, const bool *cases, int16_t *C, const int size);
template void launch_switch<int8_t,bool>(const int8_t **tensorsdata, const int numTensors, const bool *cases, int8_t *C, const int size);
template void launch_switch<bool,bool>(const bool **tensorsdata, const int numTensors, const bool *cases, bool *C, const int size);

}
#endif // DEEPX_TENSORFUNC_ELEMENTWISE_MIAO_BYTE_COMPARE_CU
88 changes: 76 additions & 12 deletions excuter/op-mem-cuda/src/deepx/tf/elementwise_compare.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -835,35 +835,99 @@ namespace deepx::tf
{

Precision C_type = mem->gettensor(this->returns[0].textvalue).get()->shape.dtype;

Precision cases_type = mem->gettensor(this->args[1].textvalue).get()->shape.dtype;

switch (C_type)
{
case Precision::Float64:
tensorfunc::Switch<Author, double>(mem->gettensors<double>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<double>(this->returns[0].textvalue));
break;
case Precision::Float32:
tensorfunc::Switch<Author, float>(mem->gettensors<float>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<float>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, double,bool>(mem->gettensors<double>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<double>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, double,int32_t>(mem->gettensors<double>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<double>(this->returns[0].textvalue));
}
break;
case Precision::Float32:
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, float,bool>(mem->gettensors<float>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<float>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, float,int32_t>(mem->gettensors<float>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<float>(this->returns[0].textvalue));
}
break;
case Precision::Float16:
tensorfunc::Switch<Author, half>(mem->gettensors<half>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<half>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, half,bool>(mem->gettensors<half>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<half>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, half,int32_t>(mem->gettensors<half>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<half>(this->returns[0].textvalue));
}
break;
case Precision::BFloat16:
tensorfunc::Switch<Author, nv_bfloat16>(mem->gettensors<nv_bfloat16>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<nv_bfloat16>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, nv_bfloat16,bool>(mem->gettensors<nv_bfloat16>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<nv_bfloat16>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, nv_bfloat16,int32_t>(mem->gettensors<nv_bfloat16>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<nv_bfloat16>(this->returns[0].textvalue));
}
break;
case Precision::Int64:
tensorfunc::Switch<Author, int64_t>(mem->gettensors<int64_t>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<int64_t>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, int64_t,bool>(mem->gettensors<int64_t>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<int64_t>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, int64_t,int32_t>(mem->gettensors<int64_t>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<int64_t>(this->returns[0].textvalue));
}
break;
case Precision::Int32:
tensorfunc::Switch<Author, int32_t>(mem->gettensors<int32_t>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<int32_t>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, int32_t,bool>(mem->gettensors<int32_t>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<int32_t>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, int32_t,int32_t>(mem->gettensors<int32_t>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<int32_t>(this->returns[0].textvalue));
}
break;
case Precision::Int16:
tensorfunc::Switch<Author, int16_t>(mem->gettensors<int16_t>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<int16_t>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, int16_t,bool>(mem->gettensors<int16_t>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<int16_t>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, int16_t,int32_t>(mem->gettensors<int16_t>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<int16_t>(this->returns[0].textvalue));
}
break;
case Precision::Int8:
tensorfunc::Switch<Author, int8_t>(mem->gettensors<int8_t>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<int8_t>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, int8_t,bool>(mem->gettensors<int8_t>(this->getvector<string>(0)), *mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<int8_t>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, int8_t,int32_t>(mem->gettensors<int8_t>(this->getvector<string>(0)), *mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<int8_t>(this->returns[0].textvalue));
}
break;
case Precision::Bool:
tensorfunc::Switch<Author, bool>(mem->gettensors<bool>(this->getvector<string>(0)), *mem->gettensor<int8_t>(this->args[1].textvalue), *mem->gettensor<bool>(this->returns[0].textvalue));
if (cases_type == Precision::Bool)
{
tensorfunc::Switch<Author, bool,bool>(mem->gettensors<bool>(this->getvector<string>(0)),*mem->gettensor<bool>(this->args[1].textvalue), *mem->gettensor<bool>(this->returns[0].textvalue));
}
else
{
tensorfunc::Switch<Author, bool,int32_t>(mem->gettensors<bool>(this->getvector<string>(0)),*mem->gettensor<int32_t>(this->args[1].textvalue), *mem->gettensor<bool>(this->returns[0].textvalue));
}
break;
default:
error = "Unsupported type: " + precision_str(C_type);
Expand Down
2 changes: 1 addition & 1 deletion excuter/op-mem-ompsimd/src/client/tfs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -492,7 +492,7 @@ namespace deepx::tf
tffactory.add_tf(std::make_shared<Switch<miaobyte>>(vector<Param>(
{
Param("tensors", DataCategory::ListTensor, Precision::Any),
Param("cases", DataCategory::Tensor, Precision::Int8),
Param("cases", DataCategory::Tensor, Precision::Bool|Precision::Int32),
}),
vector<Param>(
{
Expand Down
4 changes: 2 additions & 2 deletions excuter/op-mem-ompsimd/src/deepx/tensorfunc/elementwise_miaobyte.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -995,8 +995,8 @@ namespace deepx::tensorfunc
{
for (int j = 0; j < i_end; j++)
{
int which_tensor=cases.data[i];
C.data[i+j]=tensors[which_tensor]->data[i];
int which_tensor=cases.data[i+j];
C.data[i+j]=tensors[which_tensor]->data[i+j];
} });
}
else
Expand Down
Loading
Loading