You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/09/30 09:18:34 UTC
[GitHub] [tvm] srkreddy1238 opened a new pull request, #12951: [CODEGEN][OPENCL] Sampler definition should be at outermost scope
srkreddy1238 opened a new pull request, #12951:
URL: https://github.com/apache/tvm/pull/12951
Not all OpenCL compiers happy with inline sampler definition.
Specification: "The image read functions take a sampler argument. The sampler can be passed as an argument to the kernel using clSetKernelArg, or can be declared in the outermost scope of kernel functions, or it can be a constant variable of type sampler_t declared in the program source."
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] echuraev commented on pull request #12951: [CODEGEN][OPENCL] Sampler definition should be at outermost scope
Posted by GitBox <gi...@apache.org>.
echuraev commented on PR #12951:
URL: https://github.com/apache/tvm/pull/12951#issuecomment-1263415671
cc: @masahi
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] echuraev commented on a diff in pull request #12951: [CODEGEN][OPENCL] Sampler definition should be at outermost scope
Posted by GitBox <gi...@apache.org>.
echuraev commented on code in PR #12951:
URL: https://github.com/apache/tvm/pull/12951#discussion_r984463109
##########
src/target/source/codegen_opencl.cc:
##########
@@ -89,6 +89,17 @@ void CodeGenOpenCL::InitFuncState(const PrimFunc& f) {
void CodeGenOpenCL::PrintFuncPrefix() { stream << "__kernel void"; }
+void CodeGenOpenCL::PreFunctionBody(const PrimFunc& f) {
+ for (Var arg : f->params) {
+ auto ptr_type = arg->type_annotation.as<PointerTypeNode>();
+ if (ptr_type && runtime::IsTextureStorage(std::string(ptr_type->storage_scope))) {
+ this->stream << " const sampler_t image_sampler = "
Review Comment:
Ok, thank you for the clarification! LGTM.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] srkreddy1238 commented on a diff in pull request #12951: [CODEGEN][OPENCL] Sampler definition should be at outermost scope
Posted by GitBox <gi...@apache.org>.
srkreddy1238 commented on code in PR #12951:
URL: https://github.com/apache/tvm/pull/12951#discussion_r984447470
##########
src/target/source/codegen_opencl.cc:
##########
@@ -89,6 +89,17 @@ void CodeGenOpenCL::InitFuncState(const PrimFunc& f) {
void CodeGenOpenCL::PrintFuncPrefix() { stream << "__kernel void"; }
+void CodeGenOpenCL::PreFunctionBody(const PrimFunc& f) {
+ for (Var arg : f->params) {
+ auto ptr_type = arg->type_annotation.as<PointerTypeNode>();
+ if (ptr_type && runtime::IsTextureStorage(std::string(ptr_type->storage_scope))) {
+ this->stream << " const sampler_t image_sampler = "
Review Comment:
Its the first line of function body.
`__kernel void kernel_name(..) {`
` const sampler_t image_sampler = "CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;`
` // ...`
`}`
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] echuraev commented on a diff in pull request #12951: [CODEGEN][OPENCL] Sampler definition should be at outermost scope
Posted by GitBox <gi...@apache.org>.
echuraev commented on code in PR #12951:
URL: https://github.com/apache/tvm/pull/12951#discussion_r984408792
##########
src/target/source/codegen_opencl.cc:
##########
@@ -89,6 +89,17 @@ void CodeGenOpenCL::InitFuncState(const PrimFunc& f) {
void CodeGenOpenCL::PrintFuncPrefix() { stream << "__kernel void"; }
+void CodeGenOpenCL::PreFunctionBody(const PrimFunc& f) {
+ for (Var arg : f->params) {
+ auto ptr_type = arg->type_annotation.as<PointerTypeNode>();
+ if (ptr_type && runtime::IsTextureStorage(std::string(ptr_type->storage_scope))) {
+ this->stream << " const sampler_t image_sampler = "
Review Comment:
nit: Will we create this variable before function body in global scope?
If yes, then why we need this indent before `const sampler_t ...`?
I suppose it should be something like that:
```
const sampler_t image_sampler = "CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n";
__kernel void kernel_name(..) {
// ...
}
```
Am I right?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org