Revert "Update Cutlass to v2.11 (pytorch#94188)"

This reverts commit a0f9abd. Reverted pytorch#94188 on behalf of https://github.com/ezyang due to bouncing this to derisk branch cut
juhofuriosa · Feb 13, 2023 · 36dfbb0 · 36dfbb0
1 parent f70ba23
commit 36dfbb0
Show file tree

Hide file tree

Showing 6 changed files with 7 additions and 4 deletions.
diff --git a/BUILD.bazel b/BUILD.bazel
@@ -414,6 +414,7 @@ cc_library(
 torch_cuda_half_options = [
     "-DCUDA_HAS_FP16=1",
     "-D__CUDA_NO_HALF_OPERATORS__",
+    "-D__CUDA_NO_HALF_CONVERSIONS__",
     "-D__CUDA_NO_BFLOAT16_CONVERSIONS__",
     "-D__CUDA_NO_HALF2_OPERATORS__",
 ]

diff --git a/aten/src/ATen/native/cuda/KernelUtils.cuh b/aten/src/ATen/native/cuda/KernelUtils.cuh
@@ -49,14 +49,14 @@ __device__ __forceinline__ void fastSpecializedAtomicAdd(
 
   if (low_byte && index < (numel - 1)) {
     __half2 value2;
-    value2.x = static_cast<__half>(value);
+    value2.x = value;
     value2.y = __int2half_rz(0);
     atomicAdd(reinterpret_cast<__half2*>(target_addr), value2);
 
   } else if (!low_byte && index > 0) {
     __half2 value2;
     value2.x = __int2half_rz(0);
-    value2.y = static_cast<__half>(value);
+    value2.y = value;
     atomicAdd(reinterpret_cast<__half2*>(target_addr - 1), value2);
 
   } else {

diff --git a/aten/src/ATen/test/cuda_half_test.cu b/aten/src/ATen/test/cuda_half_test.cu
@@ -21,7 +21,7 @@ __device__ void test(){
 
   __half a = __float2half(3.0f);
   __half b = __float2half(2.0f);
-  __half c = Half(a) - Half(b);
+  __half c = a - Half(b);
   assert(static_cast<Half>(c) == Half(1.0));
 
   // asserting if the  functions used on

diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
@@ -1653,6 +1653,7 @@ if(NOT INTERN_BUILD_MOBILE)
   message(STATUS "Found CUDA with FP16 support, compiling with torch.cuda.HalfTensor")
   string(APPEND CMAKE_CUDA_FLAGS " -DCUDA_HAS_FP16=1"
                                  " -D__CUDA_NO_HALF_OPERATORS__"
+                                 " -D__CUDA_NO_HALF_CONVERSIONS__"
                                  " -D__CUDA_NO_HALF2_OPERATORS__"
                                  " -D__CUDA_NO_BFLOAT16_CONVERSIONS__")
 

diff --git a/third_party/cutlass b/third_party/cutlass
diff --git a/torch/utils/cpp_extension.py b/torch/utils/cpp_extension.py
@@ -225,6 +225,7 @@ def _join_rocm_home(*paths) -> str:
 
 COMMON_NVCC_FLAGS = [
     '-D__CUDA_NO_HALF_OPERATORS__',
+    '-D__CUDA_NO_HALF_CONVERSIONS__',
     '-D__CUDA_NO_BFLOAT16_CONVERSIONS__',
     '-D__CUDA_NO_HALF2_OPERATORS__',
     '--expt-relaxed-constexpr'