From e2e2910dc1d737796503cb7009435a164b571015 Mon Sep 17 00:00:00 2001
From: JackAKirk <chezjakirk@gmail.com>
Date: Wed, 30 Jun 2021 11:39:44 +0100
Subject: [PATCH]  [SYCL][CUDA] cl_khr_fp16 extension connected to cuda PI.

This change ensures that the fp16 Reduction test case runs for the cuda backend.
The test fp16-with-unnamed-lambda.cpp has been deleted because it has a duplicate in the test suite (in the dir SYCL/Regression).
In both cases the triple is missing on the first line which needs to be added to the llvm-test-suite copy to avoid a test failure now that the test is not skipped for the cuda backend.

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
---
 sycl/plugins/cuda/pi_cuda.cpp                 | 16 +++++++
 .../regression/fp16-with-unnamed-lambda.cpp   | 43 -------------------
 2 files changed, 16 insertions(+), 43 deletions(-)
 delete mode 100644 sycl/test/regression/fp16-with-unnamed-lambda.cpp

diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp
index ec49dc3f04ce3..a21268a46319e 100644
--- a/sycl/plugins/cuda/pi_cuda.cpp
+++ b/sycl/plugins/cuda/pi_cuda.cpp
@@ -1378,6 +1378,22 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
 
     std::string SupportedExtensions = "cl_khr_fp64 ";
 
+    int major = 0;
+    int minor = 0;
+
+    cl::sycl::detail::pi::assertion(
+        cuDeviceGetAttribute(&major,
+                             CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
+                             device->get()) == CUDA_SUCCESS);
+    cl::sycl::detail::pi::assertion(
+        cuDeviceGetAttribute(&minor,
+                             CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
+                             device->get()) == CUDA_SUCCESS);
+
+    if ((major >= 6) || ((major == 5) && (minor >= 3))) {
+      SupportedExtensions += "cl_khr_fp16 ";
+    }
+
     return getInfo(param_value_size, param_value, param_value_size_ret,
                    SupportedExtensions.c_str());
   }
diff --git a/sycl/test/regression/fp16-with-unnamed-lambda.cpp b/sycl/test/regression/fp16-with-unnamed-lambda.cpp
deleted file mode 100644
index 91f3f69bdc3bd..0000000000000
--- a/sycl/test/regression/fp16-with-unnamed-lambda.cpp
+++ /dev/null
@@ -1,43 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-unnamed-lambda %s -o %t.out
-// RUN: %RUN_ON_HOST %t.out
-#include <CL/sycl.hpp>
-
-#include <iostream>
-
-int main() {
-  auto AsyncHandler = [](cl::sycl::exception_list EL) {
-    for (std::exception_ptr const &P : EL) {
-      try {
-        std::rethrow_exception(P);
-      } catch (std::exception const &E) {
-        std::cerr << "Caught async SYCL exception: " << E.what() << std::endl;
-      }
-    }
-  };
-
-  cl::sycl::queue Q(AsyncHandler);
-
-  cl::sycl::device D = Q.get_device();
-  if (!D.has_extension("cl_khr_fp16"))
-    return 0; // Skip the test if halfs are not supported
-
-  cl::sycl::buffer<cl::sycl::cl_half> Buf(1);
-
-  Q.submit([&](cl::sycl::handler &CGH) {
-    auto Acc = Buf.get_access<cl::sycl::access::mode::write>(CGH);
-    CGH.single_task([=]() {
-      Acc[0] = 1;
-    });
-  });
-
-  Q.wait_and_throw();
-
-  auto Acc = Buf.get_access<cl::sycl::access::mode::read>();
-  if (1 != Acc[0]) {
-    std::cerr << "Incorrect result, got: " << Acc[0]
-              << ", expected: 1" << std::endl;
-    return 1;
-  }
-
-  return 0;
-}