Merge branch 'backport-clang-format-cuda' into release

Modules/CMakeCUDACompilerABI.cu

@@ -1,5 +1,5 @@
 #ifndef __CUDACC__
-# error "A C or C++ compiler has been selected for CUDA"
+#error "A C or C++ compiler has been selected for CUDA"
 #endif
 
 #include "CMakeCompilerABI.h"

Tests/Cuda/.clang-format

@@ -0,0 +1,9 @@
+---
+# This configuration requires clang-format 3.8 or higher.
+BasedOnStyle: Mozilla
+AlignOperands: false
+AlwaysBreakAfterReturnType: None
+AlwaysBreakAfterDefinitionReturnType: None
+ColumnLimit: 79
+Standard: Cpp11
+...

Tests/Cuda/Complex/dynamic.cu

@@ -1,7 +1,7 @@
 
-#include <string>
 #include <cuda.h>
 #include <iostream>
+#include <string>
 
 #ifdef _WIN32
 #define EXPORT __declspec(dllexport)
@@ -16,18 +16,15 @@ EXPORT int __host__ cuda_dynamic_host_func(int x)
   return dynamic_base_func(x);
 }
 
-static
+static __global__ void DetermineIfValidCudaDevice()
-__global__
-void DetermineIfValidCudaDevice()
 {
 }
 
 EXPORT void cuda_dynamic_lib_func()
 {
-  DetermineIfValidCudaDevice <<<1,1>>> ();
+  DetermineIfValidCudaDevice<<<1, 1>>>();
   cudaError_t err = cudaGetLastError();
-  if(err == cudaSuccess)
+  if (err == cudaSuccess) {
-    {
     std::cerr << cudaGetErrorString(err) << std::endl;
-    }
+  }
 }

Tests/Cuda/Complex/file1.cu

@@ -5,6 +5,6 @@ result_type __device__ file1_func(int x)
 {
   result_type r;
   r.input = x;
-  r.sum = x*x;
+  r.sum = x * x;
   return r;
 }

Tests/Cuda/Complex/file2.cu

@@ -5,16 +5,12 @@ result_type __device__ file1_func(int x);
 
 result_type_dynamic __device__ file2_func(int x)
 {
-  if(x!=42)
+  if (x != 42) {
-    {
     const result_type r = file1_func(x);
-    const result_type_dynamic rd { r.input, r.sum, true };
+    const result_type_dynamic rd{ r.input, r.sum, true };
     return rd;
-    }
+  } else {
-  else
+    const result_type_dynamic rd{ x, x * x * x, false };
-    {
-    const result_type_dynamic rd { x, x*x*x, false };
     return rd;
-    }
+  }
-
 }

Tests/Cuda/Complex/file3.cu

@@ -7,9 +7,7 @@
 result_type __device__ file1_func(int x);
 result_type_dynamic __device__ file2_func(int x);
 
-static
+static __global__ void file3_kernel(result_type& r, int x)
-__global__
-void file3_kernel(result_type& r, int x)
 {
   r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
@@ -18,12 +16,11 @@ void file3_kernel(result_type& r, int x)
 int file3_launch_kernel(int x)
 {
   result_type r;
-  file3_kernel <<<1,1>>> (r,x);
+  file3_kernel<<<1, 1>>>(r, x);
   cudaError_t err = cudaGetLastError();
-  if(err == cudaSuccess)
+  if (err == cudaSuccess) {
-    {
     std::cerr << cudaGetErrorString(err) << std::endl;
     return x;
-    }
+  }
   return r.sum;
 }

Tests/Cuda/Complex/mixed.cu

@@ -17,9 +17,7 @@ result_type_dynamic __device__ file2_func(int x);
 
 IMPORT void __host__ cuda_dynamic_lib_func();
 
-static
+static __global__ void mixed_kernel(result_type& r, int x)
-__global__
-void mixed_kernel(result_type& r, int x)
 {
   r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
@@ -30,6 +28,6 @@ EXPORT int mixed_launch_kernel(int x)
   cuda_dynamic_lib_func();
 
   result_type r;
-  mixed_kernel <<<1,1>>> (r,x);
+  mixed_kernel<<<1, 1>>>(r, x);
   return r.sum;
 }

Tests/Cuda/ConsumeCompileFeatures/main.cu

@@ -5,14 +5,16 @@ int static_cxx11_func(int);
 
 void test_functions()
 {
-  auto x = static_cxx11_func( int(42) );
+  auto x = static_cxx11_func(int(42));
   std::cout << x << std::endl;
 }
 
-int main(int argc, char **argv)
+int main(int argc, char** argv)
 {
   test_functions();
-  std::cout << "this executable doesn't use cuda code, just call methods defined" << std::endl;
+  std::cout
+    << "this executable doesn't use cuda code, just call methods defined"
+    << std::endl;
   std::cout << "in libraries that have cuda code" << std::endl;
   return 0;
 }

Tests/Cuda/ObjectLibrary/static.cu

@@ -8,14 +8,14 @@ int __host__ file1_sq_func(int x)
   cudaError_t err;
   int nDevices = 0;
   err = cudaGetDeviceCount(&nDevices);
-  if(err != cudaSuccess)
+  if (err != cudaSuccess) {
-  {
     std::cout << "nDevices: " << nDevices << std::endl;
     std::cout << "err: " << err << std::endl;
     return 1;
   }
   std::cout << "this library uses cuda code" << std::endl;
-  std::cout << "you have " << nDevices << " devices that support cuda" << std::endl;
+  std::cout << "you have " << nDevices << " devices that support cuda"
+            << std::endl;
 
   return x * x;
 }

Tests/Cuda/ProperLinkFlags/file1.cu

@@ -6,6 +6,6 @@ result_type __device__ file1_func(int x)
   __ldg(&x);
   result_type r;
   r.input = x;
-  r.sum = x*x;
+  r.sum = x * x;
   return r;
 }

Tests/CudaOnly/.clang-format

@@ -0,0 +1,9 @@
+---
+# This configuration requires clang-format 3.8 or higher.
+BasedOnStyle: Mozilla
+AlignOperands: false
+AlwaysBreakAfterReturnType: None
+AlwaysBreakAfterDefinitionReturnType: None
+ColumnLimit: 79
+Standard: Cpp11
+...

Tests/CudaOnly/EnableStandard/main.cu

@@ -12,11 +12,11 @@ IMPORT int shared_cuda11_func(int);
 
 void test_functions()
 {
-  static_cuda11_func( int(42) );
+  static_cuda11_func(int(42));
-  shared_cuda11_func( int(42) );
+  shared_cuda11_func(int(42));
 }
 
-int main(int argc, char **argv)
+int main(int argc, char** argv)
 {
   test_functions();
   return 0;

Tests/CudaOnly/SeparateCompilation/file1.cu

@@ -5,6 +5,6 @@ result_type __device__ file1_func(int x)
 {
   result_type r;
   r.input = x;
-  r.sum = x*x;
+  r.sum = x * x;
   return r;
 }

Tests/CudaOnly/SeparateCompilation/file2.cu

@@ -5,16 +5,12 @@ result_type __device__ file1_func(int x);
 
 result_type_dynamic __device__ file2_func(int x)
 {
-  if(x!=42)
+  if (x != 42) {
-    {
     const result_type r = file1_func(x);
-    const result_type_dynamic rd { r.input, r.sum, true };
+    const result_type_dynamic rd{ r.input, r.sum, true };
     return rd;
-    }
+  } else {
-  else
+    const result_type_dynamic rd{ x, x * x * x, false };
-    {
-    const result_type_dynamic rd { x, x*x*x, false };
     return rd;
-    }
+  }
-
 }

Tests/CudaOnly/SeparateCompilation/file3.cu

@@ -6,13 +6,10 @@
 result_type __device__ file1_func(int x);
 result_type_dynamic __device__ file2_func(int x);
 
-
+static __global__ void file3_kernel(result_type& r, int x)
-static
-__global__
-void file3_kernel(result_type& r, int x)
 {
-  //call static_func which is a method that is defined in the
+  // call static_func which is a method that is defined in the
-  //static library that is always out of date
+  // static library that is always out of date
   r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
 }
@@ -20,6 +17,6 @@ void file3_kernel(result_type& r, int x)
 result_type file3_launch_kernel(int x)
 {
   result_type r;
-  file3_kernel <<<1,1>>> (r,x);
+  file3_kernel<<<1, 1>>>(r, x);
   return r;
 }

Tests/CudaOnly/SeparateCompilation/file4.cu

@@ -7,12 +7,10 @@
 result_type __device__ file1_func(int x);
 result_type_dynamic __device__ file2_func(int x);
 
-static
+static __global__ void file4_kernel(result_type& r, int x)
-__global__
-void file4_kernel(result_type& r, int x)
 {
-  //call static_func which is a method that is defined in the
+  // call static_func which is a method that is defined in the
-  //static library that is always out of date
+  // static library that is always out of date
   r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
 }
@@ -20,6 +18,6 @@ void file4_kernel(result_type& r, int x)
 int file4_launch_kernel(int x)
 {
   result_type r;
-  file4_kernel <<<1,1>>> (r,x);
+  file4_kernel<<<1, 1>>>(r, x);
   return r.sum;
 }

Tests/CudaOnly/SeparateCompilation/file5.cu

@@ -7,12 +7,10 @@
 result_type __device__ file1_func(int x);
 result_type_dynamic __device__ file2_func(int x);
 
-static
+static __global__ void file5_kernel(result_type& r, int x)
-__global__
-void file5_kernel(result_type& r, int x)
 {
-  //call static_func which is a method that is defined in the
+  // call static_func which is a method that is defined in the
-  //static library that is always out of date
+  // static library that is always out of date
   r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
 }
@@ -20,6 +18,6 @@ void file5_kernel(result_type& r, int x)
 int file5_launch_kernel(int x)
 {
   result_type r;
-  file5_kernel <<<1,1>>> (r,x);
+  file5_kernel<<<1, 1>>>(r, x);
   return r.sum;
 }

Tests/CudaOnly/SeparateCompilation/main.cu

@@ -7,7 +7,7 @@
 int file4_launch_kernel(int x);
 int file5_launch_kernel(int x);
 
-int main(int argc, char **argv)
+int main(int argc, char** argv)
 {
   file4_launch_kernel(42);
   file5_launch_kernel(42);

Tests/CudaOnly/WithDefs/main.notcu

@@ -6,14 +6,12 @@
 #error "PACKED_DEFINE not defined!"
 #endif
 
-static
+static __global__ void DetermineIfValidCudaDevice()
-__global__
-void DetermineIfValidCudaDevice()
 {
 }
 
 #ifdef _MSC_VER
-#pragma pack(push,1)
+#pragma pack(push, 1)
 #undef PACKED_DEFINE
 #define PACKED_DEFINE
 #endif
@@ -32,28 +30,24 @@ struct PACKED_DEFINE result_type
 result_type can_launch_kernel()
 {
   result_type r;
-  DetermineIfValidCudaDevice <<<1,1>>> ();
+  DetermineIfValidCudaDevice<<<1, 1>>>();
   r.valid = (cudaSuccess == cudaGetLastError());
-  if(r.valid)
+  if (r.valid) {
-    {
     r.value = 1;
-    }
+  } else {
-  else
-    {
     r.value = -1;
-    }
+  }
   return r;
 }
 
-int main(int argc, char **argv)
+int main(int argc, char** argv)
 {
   cudaError_t err;
   int nDevices = 0;
   err = cudaGetDeviceCount(&nDevices);
-  if(err != cudaSuccess)
+  if (err != cudaSuccess) {
-    {
+    std::cerr << cudaGetErrorString(err) << std::endl;
-      std::cerr << cudaGetErrorString(err) << std::endl;
+    return 1;
-      return 1;
+  }
-    }
   return 0;
 }

Utilities/Scripts/clang-format.bash

@@ -107,7 +107,7 @@ case "$mode" in
 esac
 
 # Filter sources to which our style should apply.
-$git_ls -z -- '*.c' '*.cc' '*.cpp' '*.cxx' '*.h' '*.hh' '*.hpp' '*.hxx' |
+$git_ls -z -- '*.c' '*.cc' '*.cpp' '*.cxx' '*.h' '*.hh' '*.hpp' '*.hxx' '*.cu' '*.notcu' |
 
   # Exclude lexer/parser generator input and output.
   egrep -z -v '^Source/cmCommandArgumentLexer\.' |