<html><head><meta http-equiv="Content-Type" content="text/html; charset=us-ascii"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""> <div> <blockquote type="cite" class=""><div class="">On Jun 1, 2019, at 10:57 PM, Jan Vesely <<a href="mailto:jan.vesely@rutgers.edu" class="">jan.vesely@rutgers.edu</a>> wrote:</div> <div class=""><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; float: none; display: inline !important;" class="">On Thu, 2019-05-30 at 08:40 -0400, Matt Arsenault wrote:<br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><blockquote type="cite" style="font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; orphans: auto; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; widows: auto; word-spacing: 0px; -webkit-text-size-adjust: auto; -webkit-text-stroke-width: 0px; text-decoration: none;" class="">Ping <blockquote type="cite" class="">On May 23, 2019, at 7:59 PM, <a href="mailto:arsenm2@gmail.com" class="">arsenm2@gmail.com</a> wrote: From: Matt Arsenault <<a href="mailto:arsenm2@gmail.com" class="">arsenm2@gmail.com</a>> --- .../program/execute/call-clobbers-amdgcn.cl | 102 ++++++++++++++++++ 1 file changed, 102 insertions(+) diff --git a/tests/cl/program/execute/call-clobbers-amdgcn.cl b/tests/cl/program/execute/call-clobbers-amdgcn.cl index 18e657ce3..b0a1f8c70 100644 --- a/tests/cl/program/execute/call-clobbers-amdgcn.cl +++ b/tests/cl/program/execute/call-clobbers-amdgcn.cl @@ -19,6 +19,49 @@ dimensions: 1 global_size: 1 0 0 arg_out: 0 buffer int[1] 0xabcd1234 +[test] +name: Conditional call +kernel_name: conditional_call +dimensions: 1 +local_size: 64 0 0 +global_size: 64 0 0 +arg_out: 0 buffer int[64] \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 + +[test] +name: Conditional call partial dispatch +kernel_name: conditional_call +dimensions: 1 +local_size: 16 0 0 +global_size: 16 0 0 +arg_out: 0 buffer int[16] \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \ + 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 + + +[test] +name: Skip call no lanes +kernel_name: skip_call_no_lanes +dimensions: 1 +local_size: 64 0 0 +global_size: 64 0 0 +arg_out: 0 buffer int[64] \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 \ + 123 123 123 123 123 123 123 123 + !*/ #ifndef __AMDGCN__ @@ -65,3 +108,62 @@ kernel void call_clobber_v40(__global int* ret) : "v40"); *ret = tmp; } + +__attribute__((noinline)) +void spill_sgpr_to_csr_vgpr() +{ + __asm volatile( + "s_nop 1" ::: + "v0","v1","v2","v3","v4","v5","v6","v7", + "v8","v9","v10","v11","v12","v13","v14","v15", + "v16","v17","v18","v19","v20","v21","v22","v23", + "v24","v25","v26","v27","v28","v29","v30","v31", + + "s0","s1","s2","s3","s4","s5","s6","s7", + "s8","s9","s10","s11","s12","s13","s14","s15", + "s16","s17","s18","s19","s20","s21","s22","s23", + "s24","s25","s26","s27","s28","s29","s30","s31", + "s32", "s33", "s34", "s35", "s36", "s37", "s38"); +} + +// A CSR VGPR needs to be spilled/restored in the prolog/epilog, but +// all lanes need to be made active to avoid clobbering lanes that did +// not enter the call. +kernel void conditional_call(global int* ret) +{ + __asm volatile("v_mov_b32 v32, 0xabcd1234" : : : "v32"); + + int id = get_local_id(0); + if (id == 0) + { + spill_sgpr_to_csr_vgpr(); + } + + int tmp; + __asm volatile("v_mov_b32 %0, v32" + : "=v"(tmp) + : + : "v32"); + ret[id] = tmp; +} + +__attribute__((noinline)) +void hang_if_all_inactive() +{ + __builtin_amdgcn_s_sendmsghalt(0, 0); +} + +// If all lanes could be dynamically false, the call must not be taken +// in case a side effecting scalar op is called inside. +kernel void skip_call_no_lanes(global int* ret) +{ + int divergent_false; + __asm volatile("v_mov_b32 %0, 0" : "=v"(divergent_false)); + + if (divergent_false) + { + hang_if_all_inactive(); </blockquote></blockquote><br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; float: none; display: inline !important;" class="">this looks like it will hang the GPU on test failure, which is a no-<br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; float: none; display: inline !important;" class="">go.<br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; float: none; display: inline !important;" class="">Jan<br style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""></div></blockquote></div> <div class=""> </div><div class="">Is there a way to specify a timeout? The alternatives require more ABI support</div></body></html>