Skip to content
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.

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

[CTS] basic/if failed #16

Closed
guaguabao123 opened this issue Sep 11, 2023 · 5 comments
Closed

[CTS] basic/if failed #16

guaguabao123 opened this issue Sep 11, 2023 · 5 comments
Labels
bug Something isn't working

Comments

@guaguabao123
Copy link

  1. run test_basic if failed
  2. simple change case to test
    A. set 32 threads , run test_basic if 32
    B.
    WXWorkLocal_16944375885891
  3. run failed
    input data has 0x1 ,but output data is 0x7fffffff

WXWorkLocal_16944229622649

@guaguabao123
Copy link
Author

@zhoujingya
Copy link
Contributor

image 确定是join指令这边出问题了

@THU-DSP-LAB THU-DSP-LAB deleted a comment from zhoujingfighting Sep 13, 2023
@zhoujingya
Copy link
Contributor

zhoujingya commented Sep 15, 2023


__kernel void test_if(__global int *src, __global int *dst)
{
    int  tid = get_global_id(0);

    if (src[tid] == 0)
        dst[tid] = 0x12345678;
    else if (src[tid] == 1)
        dst[tid] = 0x23456781;
    else
        dst[tid] = 0x7FFFFFFF;
}

上述代码在LLVM IR生成的时候,CPU执行逻辑会在if和else分支,dst[tid]这个内存地址统一存select之后的结果,如0x23456781或者0x7FFFFFFF,但是如https://github.com/THU-DSP-LAB/llvm-project/issues/16#issuecomment-1717686570,这上面贴的图,GPU else与if分支的结果不一样,需要分开存,对应到汇编码就是join语句之后的vmv这个操作,这个操作应该提前到每个分支执行

@zhoujingya
Copy link
Contributor

zhoujingya commented Sep 25, 2023

之前的提交有点问题:4bfde6a
,这个提交会导致barrier与wg_barrier过不了,问题都是一样的,if与else分支的执行结果需要分开存储,之前的没有处理完善

zhoujingya pushed a commit that referenced this issue Sep 25, 2023
When vmv instruction is in join block, need to check the register definition
are in both else&if predecessors or not
Github issue: #16
@zhoujingya
Copy link
Contributor

zhoujingya commented Sep 25, 2023

Solved by this commit 58e57e8

zhoujingya pushed a commit that referenced this issue Sep 26, 2023
When vmv instruction is in join block, need to check the register definition
are in both else&if predecessors or not
Github issue: #16
@zhoujingya zhoujingya reopened this Oct 25, 2023
@zhoujingya zhoujingya added the bug Something isn't working label Oct 25, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working
Projects
None yet
2 participants