Hacking ck_tile fmha Dropout facility

[qianfengz]

This PR makes two changes in ck_tile BlockDropout facility and forward kernel:

1. Add NullBlockDropout, which is used to reduce the vgprs/sgprs when kHasDropout is false (pure inference path)
2. Change in BlockDropout::Run() to reduce the conditional checking for is_store_randval

• From xformers testing, the change is able to improve the fmha forward performance when dropout_prob == 0 visibly, like

  #before the change (last second column shows exec time)
  "f16 16-1024-16-64, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,236,32.0
  "f16 16-1024-16-128, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,508,64.0
  "f16 16-1024-16-160, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,2021,80.0
  "f16 16-1024-16-256, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,1003,128.0
  "b16 1024-197-1-88, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,94,33.859375
  "b16 1024-197-1-80, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,92,30.78125
  "b16 1-4096-16-80, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,318,10.0
  "b16 1-8192-4-128, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,378,8.0
  "b16 16-128-16-16, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,33,1.0
  "b16 16-512-16-256, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,274,64.0
  #after the change (last second column shows exec time)
  "f16 16-1024-16-64, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,235,32.0
  "f16 16-1024-16-128, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,435,64.0
  "f16 16-1024-16-160, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,2008,80.0
  "f16 16-1024-16-256, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,1007,128.0
  "b16 1024-197-1-88, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,94,33.859375
  "b16 1024-197-1-80, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,92,30.78125
  "b16 1-4096-16-80, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,317,10.0
  "b16 1-8192-4-128, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,378,8.0
  "b16 16-128-16-16, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,32,1.0
  "b16 16-512-16-256, p=0.0, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,273,64.0

• From xformers testing, the change is able to improve the fmha forward performance when dropout_prob > 0 obviously, like

#before the change (last second column shows exec time)
"f16 384-197-1-88, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,65,12.697265625
"f16 384-197-1-64, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,36,9.234375
"f16 1024-197-1-64, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,89,24.625
"f16 32-197-16-80, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,94,15.390625
"f16 16-197-16-88, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,45,8.46484375
"f16 150-256-16-64, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,261,75.0
"f16 1-16384-16-40, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,3858,20.0
"f16 2-2048-8-128, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,167,8.0
"f16 16-128-16-256, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,46,16.0
"f16 16-1024-16-32, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,244,16.0
#after  the change (last second column shows exec time_
"f16 384-197-1-88, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,53,12.697265625
"f16 384-197-1-64, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,35,9.234375
"f16 1024-197-1-64, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,80,24.625
"f16 32-197-16-80, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,88,15.390625
"f16 16-197-16-88, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,38,8.46484375
"f16 150-256-16-64, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,252,75.0
"f16 1-16384-16-40, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,3746,20.0
"f16 2-2048-8-128, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,143,8.0
"f16 16-128-16-256, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,36,16.0
"f16 16-1024-16-32, p=0.3, BiasT=NoneType",attention (attn_bias=<class 'NoneType'>),1,ckF,,232,16.0

• The Change is not applied to backward kernel since fmha backward is going to be refactored