Renamed AMD SI workaround defines

2 files changed, 7 insertions, 8 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 2a26193d..fb62ae27 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -36,9 +36,9 @@ option(NETLIB "Enable compilation of the CBLAS Netlib API" OFF)
 option(CUBLAS "Enables performance comparison against cuBLAS on NVIDIA GPUs" OFF)
 
 # Workarounds for bugs
-option(AMD_HAINAN "Enables workaround for bug in AMD SI Hainan GPUs" OFF)
-if(AMD_HAINAN)
-  add_definitions(-DAMD_HAINAN_WORKAROUND)
+option(AMD_SI_EMPTY_KERNEL_WORKAROUND "Enables workaround for bug in AMD Southern Island GPUs" OFF)
+if(AMD_SI_EMPTY_KERNEL_WORKAROUND)
+  add_definitions(-DAMD_SI_EMPTY_KERNEL_WORKAROUND)
 endif()
 
 # Select between an OpenCL API (default) or a CUDA API (beta)
diff --git a/src/clpp11.hpp b/src/clpp11.hpp
index 45ef52d9..093214b3 100644
--- a/src/clpp11.hpp
+++ b/src/clpp11.hpp
@@ -447,11 +447,10 @@ class Program {
 
   // Source-based constructor with memory management
   explicit Program(const Context &context, const std::string &source) {
-    #ifdef AMD_HAINAN_WORKAROUND
+    #ifdef AMD_SI_EMPTY_KERNEL_WORKAROUND
       const std::string source_hainan = source + "\n__kernel void null_kernel() {}\n";
       const char *source_ptr = &source_hainan[0];
       const auto length = source_hainan.length();
-      printf("%s\n", source_hainan.c_str());
     #else
       const char *source_ptr = &source[0];
       const auto length = source.length();
@@ -775,7 +774,7 @@ class Kernel {
     auto status = CL_SUCCESS;
     *kernel_ = clCreateKernel(program->operator()(), name.c_str(), &status);
     CLCudaAPIError::Check(status, "clCreateKernel");
-    #ifdef AMD_HAINAN_WORKAROUND
+    #ifdef AMD_SI_EMPTY_KERNEL_WORKAROUND
       *null_kernel_ = clCreateKernel(program->operator()(), "null_kernel", &status);
       CLCudaAPIError::Check(status, "clCreateKernel");
     #endif
@@ -842,7 +841,7 @@ class Kernel {
                                       static_cast<cl_uint>(waitForEventsPlain.size()),
                                       !waitForEventsPlain.empty() ? waitForEventsPlain.data() : nullptr,
                                       event));
-    #ifdef AMD_HAINAN_WORKAROUND
+    #ifdef AMD_SI_EMPTY_KERNEL_WORKAROUND
       const std::vector<size_t> nullRange = {1};
       CheckError(clEnqueueNDRangeKernel(queue(), *null_kernel_, static_cast<cl_uint>(nullRange.size()),
                                         nullptr, nullRange.data(), nullptr,
@@ -855,7 +854,7 @@ class Kernel {
   const cl_kernel& operator()() const { return *kernel_; }
  private:
   std::shared_ptr<cl_kernel> kernel_;
-  #ifdef AMD_HAINAN_WORKAROUND
+  #ifdef AMD_SI_EMPTY_KERNEL_WORKAROUND
     std::shared_ptr<cl_kernel> null_kernel_;
   #endif