aboutsummaryrefslogtreecommitdiff
path: root/compute_kernel_writer
diff options
context:
space:
mode:
Diffstat (limited to 'compute_kernel_writer')
-rw-r--r--compute_kernel_writer/prototype/include/ckw/KernelWriter.h2
-rw-r--r--compute_kernel_writer/prototype/include/ckw/types/Functions.h4
-rw-r--r--compute_kernel_writer/prototype/src/KernelWriter.cpp2
-rw-r--r--compute_kernel_writer/prototype/src/Prototype.h6
4 files changed, 11 insertions, 3 deletions
diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h
index 72f85c78aa..fdb5fedc59 100644
--- a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h
+++ b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h
@@ -282,7 +282,7 @@ public:
* @param[out] dst The tile to be written to.
* @param[in] dim The global ID dimension.
*/
- void op_get_global_id(TileOperand &dst, int32_t dim);
+ void op_get_global_id(const TileOperand &dst, int32_t dim);
// =============================================================================================
// Code generation
diff --git a/compute_kernel_writer/prototype/include/ckw/types/Functions.h b/compute_kernel_writer/prototype/include/ckw/types/Functions.h
index 2dd5ed0b3d..bc1f85c188 100644
--- a/compute_kernel_writer/prototype/include/ckw/types/Functions.h
+++ b/compute_kernel_writer/prototype/include/ckw/types/Functions.h
@@ -39,9 +39,10 @@ enum class UnaryFunction : int32_t
Fabs = 0x0004,
Log = 0x0006,
Round = 0x0007,
+ Floor = 0x0008,
// Misc
- SizeOf = 0x0008,
+ SizeOf = 0x0009,
};
enum class BinaryFunction : int32_t
@@ -53,6 +54,7 @@ enum class BinaryFunction : int32_t
enum class TernaryFunction : int32_t
{
Select = 0x0000,
+ Clamp = 0x0001,
};
} // namespace ckw
diff --git a/compute_kernel_writer/prototype/src/KernelWriter.cpp b/compute_kernel_writer/prototype/src/KernelWriter.cpp
index f29cf12802..5c9a16ee33 100644
--- a/compute_kernel_writer/prototype/src/KernelWriter.cpp
+++ b/compute_kernel_writer/prototype/src/KernelWriter.cpp
@@ -341,7 +341,7 @@ void KernelWriter::op_for_loop(const TileOperand &var_name, BinaryOp cond_op, co
// Misc
// =================================================================================================
-void KernelWriter::op_get_global_id(TileOperand &dst, int32_t dim)
+void KernelWriter::op_get_global_id(const TileOperand &dst, int32_t dim)
{
_impl->op_get_global_id(prototype::Operand(dst.name()), dim);
}
diff --git a/compute_kernel_writer/prototype/src/Prototype.h b/compute_kernel_writer/prototype/src/Prototype.h
index 2b519471ac..88d6e898e4 100644
--- a/compute_kernel_writer/prototype/src/Prototype.h
+++ b/compute_kernel_writer/prototype/src/Prototype.h
@@ -3694,6 +3694,9 @@ public:
case UnaryFunction::Round:
_data->code += "round(";
break;
+ case UnaryFunction::Floor:
+ _data->code += "floor(";
+ break;
default:
CKW_ASSERT_MSG(false, "Unexpected UnaryFunction used.");
}
@@ -3772,6 +3775,9 @@ public:
case TernaryFunction::Select:
_data->code += "select(";
break;
+ case TernaryFunction::Clamp:
+ _data->code += "clamp(";
+ break;
default:
CKW_ASSERT_MSG(false, "Unexpected TernaryFunction used.");
}