Extend CK gemm/sdpa support to gfx950 (#45)

Update CK for gfx950 (#49)

Author: alugorey

aten/src/ATen/Context.cpp

@@ -416,7 +416,7 @@ void Context::setROCmFAPreferredBackend(at::ROCmFABackend b) {
   if(b == at::ROCmFABackend::Ck) {
     static const bool ck_unsupported = []() {
       static const std::vector<std::string> archs = {
-          "gfx90a",  "gfx942"
+          "gfx90a",  "gfx942", "gfx950"
       };
       for (auto index: c10::irange(detail::getCUDAHooks().deviceCount())) {
         if (!detail::getCUDAHooks().isGPUArch(archs, index)) {

aten/src/ATen/native/transformers/hip/flash_attn/ck/fmha_bwd.hpp

@@ -453,4 +453,5 @@ struct fmha_bwd_traits
     bool is_deterministic;
     // TODO: padding check is inside this api
 };
+template <int Version = 2>
 float fmha_bwd(fmha_bwd_traits, fmha_bwd_args, const ck_tile::stream_config&);

aten/src/ATen/native/transformers/hip/flash_attn/ck/launch_kernel_pt.hpp

@@ -14,7 +14,7 @@ __launch_bounds__(MaxThreadPerBlock, MinBlockPerCu)
 #endif
     __global__ void kentry_pt(Args... args)
 {
-#if (defined(__gfx90a__) || defined(__gfx942__))
+#if (defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__))
     Kernel{}(args...);
 #else
     CUDA_KERNEL_ASSERT(false && "Fatal! Attempting to call a CK SDPA kernel on unsupported hardware");