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 2021/06/18 21:07:11 UTC

[GitHub] [tvm] Lunderberg opened a new pull request #8281: [Vulkan][Codegen] Fixed SPIR-V scoping bug with threadIdx

Lunderberg opened a new pull request #8281:
URL: https://github.com/apache/tvm/pull/8281


   Change is split into two commits, one which implements the fix, and one which refactors the result into a more general use.
   
   * [Vulkan][Codegen] Fixed SPIR-V scoping bug with threadIdx
   
     Unlike in Cuda, where the threadIdx.x/y/z are separate built-in variables, in vulkan the built-in thread index consists of an array
   that must be dereferenced.  If `te.thread_axis("threadIdx.x")` is first declared inside a scope, then that same python object is passed as the IterVar of a separate scope, this breaks the spirv scoping rules, and may result in an undefined variable.  This only applies to threadIdx/blockIdx variables, as all other variable declarations obey tir's scoping rules.
   
     To resolve, the threadIdx and blockIdx variables are declared at the top of the function declaration.  This makes the generated spirv code follow the same semantics as tir/cuda.
   
   * [Vulkan][Codegen] Refactor of GetLocalID and GetWorkgroupID
   
     Shared behavior separated out into GetBuiltInValue.


-- 
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.

For queries about this service, please contact Infrastructure at:
users@infra.apache.org

[GitHub] [tvm] Lunderberg commented on pull request #8281: [Vulkan][Codegen] Fixed SPIR-V scoping bug with threadIdx

Posted by GitBox <gi...@apache.org>.

Lunderberg commented on pull request #8281:
URL: https://github.com/apache/tvm/pull/8281#issuecomment-864187138


   Potential reviewer: @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.

For queries about this service, please contact Infrastructure at:
users@infra.apache.org

[GitHub] [tvm] masahi merged pull request #8281: [Vulkan][Codegen] Fixed SPIR-V scoping bug with threadIdx

Posted by GitBox <gi...@apache.org>.

masahi merged pull request #8281:
URL: https://github.com/apache/tvm/pull/8281


   


-- 
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.

For queries about this service, please contact Infrastructure at:
users@infra.apache.org

[GitHub] [tvm] masahi commented on pull request #8281: [Vulkan][Codegen] Fixed SPIR-V scoping bug with threadIdx

Posted by GitBox <gi...@apache.org>.

masahi commented on pull request #8281:
URL: https://github.com/apache/tvm/pull/8281#issuecomment-864330009


   thanks @Lunderberg 


-- 
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.

For queries about this service, please contact Infrastructure at:
users@infra.apache.org