[Core] Add a device::arch method returning a string describing the ar…

…ch of the device (#679)
libocca · May 10, 2023 · e3932ab · e3932ab
1 parent a25bd3d
commit e3932ab
Show file tree

Hide file tree

Showing 22 changed files with 133 additions and 136 deletions.
diff --git a/include/occa/c/base.h b/include/occa/c/base.h
@@ -25,10 +25,6 @@ occaJson occaDeviceProperties();
 
 void occaFinish();
 
-void occaGetDeviceArchVersion(occaDevice device,
-                              int *archMajorVersion,
-                              int *archMinorVersion);
-
 occaStream occaCreateStream(occaJson props);
 
 occaStream occaGetStream();

diff --git a/include/occa/c/device.h b/include/occa/c/device.h
@@ -15,6 +15,8 @@ const char* occaDeviceMode(occaDevice device);
 
 occaJson occaDeviceGetProperties(occaDevice device);
 
+const char* occaDeviceArch(occaDevice device);
+
 occaJson occaDeviceGetKernelProperties(occaDevice device);
 
 occaJson occaDeviceGetMemoryProperties(occaDevice device);

diff --git a/include/occa/core/device.hpp b/include/occa/core/device.hpp
@@ -283,17 +283,19 @@ namespace occa {
     const occa::json& properties() const;
 
     /**
-     * @startDoc{getDeviceArchVersion}
+     * @startDoc{arch}
      *
      * Description:
-     *   Returns the architecture version of a device.
+     *   Returns a string describing some architecture information of a device.
+     *   The string's contents may vary between backends, even on the same
+     *   physical device
      *
      * Returns:
-     *   The architecture version of a device.
+     *   The architecture version string of a device, such as 'gfx900' or 'sm_70'
      *
      * @endDoc
      */
-    void getDeviceArchVersion(int *archMajorVersion, int *archMinorVersion) const;
+    const std::string& arch() const;
 
     const occa::json& kernelProperties() const;
     occa::json kernelProperties(const occa::json &additionalProps) const;
@@ -746,17 +748,17 @@ namespace occa {
 
     /**
      * @startDoc{unwrap}
-     * 
+     *
      * Description:
      *   Retreives the mode-specific object associated with this [[device]].
      *   The lifetime of the returned object is the same as this device.
-     *   Destruction of the returned object during this device's lifetime results in undefined behavior.   
-     *  
+     *   Destruction of the returned object during this device's lifetime results in undefined behavior.
+     *
      *   > An OCCA application is responsible for correctly converting the returned `void*` pointer to the corresponding mode-specific device type.
-     *   
+     *
      * Returns:
      *   A pointer to the mode-specific object associated with this device.
-     * 
+     *
      * @endDoc
     */
     void* unwrap();

diff --git a/src/c/base.cpp b/src/c/base.cpp
@@ -44,12 +44,6 @@ void occaFinish() {
   occa::finish();
 }
 
-void occaGetDeviceArchVersion(occaDevice device,
-                              int *archMajorVersion,
-                              int *archMinorVersion) {
-  (occa::c::device(device)).getDeviceArchVersion(archMajorVersion, archMinorVersion);
-}
-
 occaStream occaCreateStream(occaJson props) {
   occa::stream stream;
   if (occa::c::isDefault(props)) {

diff --git a/src/c/device.cpp b/src/c/device.cpp
@@ -43,6 +43,10 @@ occaJson occaDeviceGetProperties(occaDevice device) {
   return occa::c::newOccaType(props, false);
 }
 
+const char* occaDeviceArch(occaDevice device) {
+  return occa::c::device(device).arch().c_str();
+}
+
 occaJson occaDeviceGetKernelProperties(occaDevice device) {
   const occa::json &props = occa::c::device(device).kernelProperties();
   return occa::c::newOccaType(props, false);

diff --git a/src/core/device.cpp b/src/core/device.cpp
@@ -176,9 +176,11 @@ namespace occa {
     return modeDevice->properties;
   }
 
-  void device::getDeviceArchVersion(int *archMajorVersion, int *archMinorVersion) const {
-    assertInitialized();
-    modeDevice->getDeviceArchVersion(archMajorVersion, archMinorVersion);
+  const std::string& device::arch() const {
+    static const std::string noArch = "No Arch";
+    return (modeDevice
+            ? modeDevice->arch
+            : noArch);
   }
 
   const occa::json& device::kernelProperties() const {

diff --git a/src/occa/internal/core/device.cpp b/src/occa/internal/core/device.cpp
@@ -142,12 +142,4 @@ namespace occa {
       cachedKernels.erase(it);
     }
   }
-
-  void modeDevice_t::getDeviceArchVersion(int *archMajorVersion,
-                                          int *archMinorVersion) const {
-    if (archMajorVersion != nullptr) *archMajorVersion = 0;
-    if (archMinorVersion != nullptr) *archMinorVersion = 0;
-    std::cout << "getDeviceArchVersion called with device mode = " <<
-          mode << ", but the functionality is not supported.\n";
-  }
 }
diff --git a/src/occa/internal/core/device.hpp b/src/occa/internal/core/device.hpp
@@ -10,6 +10,7 @@ namespace occa {
   class modeDevice_t {
    public:
     std::string mode;
+    std::string arch;
     occa::json properties;
     bool needsLauncherKernel;
 
@@ -70,9 +71,6 @@ namespace occa {
     virtual hash_t hash() const = 0;
     virtual hash_t kernelHash(const occa::json &props) const = 0;
 
-    virtual void getDeviceArchVersion(int *archMajorVersion,
-                                      int *archMinorVersion) const;
-
     //  |---[ Stream ]------------------
     virtual modeStream_t* createStream(const occa::json &props) = 0;
     virtual modeStream_t* wrapStream(void *ptr, const occa::json &props) = 0;

diff --git a/src/occa/internal/modes/cuda/device.cpp b/src/occa/internal/modes/cuda/device.cpp
@@ -59,24 +59,8 @@ namespace occa {
       kernelProps["compiler"] = compiler;
       kernelProps["compiler_flags"] = compilerFlags;
 
-#if CUDA_VERSION < 5000
-      OCCA_CUDA_ERROR("Device: Getting CUDA device arch",
-                      cuDeviceComputeCapability(&archMajorVersion,
-                                                &archMinorVersion,
-                                                cuDevice));
-#else
-      OCCA_CUDA_ERROR("Device: Getting CUDA device major version",
-                      cuDeviceGetAttribute(&archMajorVersion,
-                                           CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
-                                           cuDevice));
-      OCCA_CUDA_ERROR("Device: Getting CUDA device minor version",
-                      cuDeviceGetAttribute(&archMinorVersion,
-                                           CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
-                                           cuDevice));
-#endif
-
-      archMajorVersion = kernelProps.get("arch/major", archMajorVersion);
-      archMinorVersion = kernelProps.get("arch/minor", archMinorVersion);
+      getDeviceArchVersion(cuDevice, archMajorVersion, archMinorVersion);
+      arch = getDeviceArch(cuDevice);
     }
 
     device::~device() {
@@ -128,12 +112,6 @@ namespace occa {
                       cuCtxSetCurrent(cuContext));
     }
 
-    void device::getDeviceArchVersion(int *archMajorVersion_,
-                                      int *archMinorVersion_) const {
-      if (archMajorVersion_ != nullptr) *archMajorVersion_ = archMajorVersion;
-      if (archMinorVersion_ != nullptr) *archMinorVersion_ = archMinorVersion;
-    }
-
     //---[ Stream ]---------------------
     modeStream_t* device::createStream(const occa::json &props) {
       CUstream cuStream = NULL;
@@ -265,9 +243,11 @@ namespace occa {
     void device::setArchCompilerFlags(const occa::json &kernelProps,
                                       std::string &compilerFlags) {
       if (compilerFlags.find("-arch=sm_") == std::string::npos) {
+        int majorVersion = kernelProps.get("arch/major", archMajorVersion);
+        int minorVersion = kernelProps.get("arch/minor", archMinorVersion);
         compilerFlags += " -arch=sm_";
-        compilerFlags += std::to_string(archMajorVersion);
-        compilerFlags += std::to_string(archMinorVersion);
+        compilerFlags += std::to_string(majorVersion);
+        compilerFlags += std::to_string(minorVersion);
       }
     }
 
@@ -329,7 +309,7 @@ namespace occa {
       } else if (verbose) {
           io::stdout << "Output:\n\n" << commandOutput << "\n";
       }
-      
+
       io::sync(binaryFilename);
     }
 

diff --git a/src/occa/internal/modes/cuda/device.hpp b/src/occa/internal/modes/cuda/device.hpp
@@ -42,9 +42,6 @@ namespace occa {
 
       void setCudaContext();
 
-      void getDeviceArchVersion(int *archMajorVersion_,
-                                int *archMinorVersion_) const override;
-
       //---[ Stream ]-------------------
       modeStream_t* createStream(const occa::json &props) override;
       modeStream_t* wrapStream(void* ptr, const occa::json &props) override;

diff --git a/src/occa/internal/modes/cuda/registration.cpp b/src/occa/internal/modes/cuda/registration.cpp
@@ -25,13 +25,15 @@ namespace occa {
         for (int deviceId = 0; deviceId < deviceCount; ++deviceId) {
           const udim_t bytes         = getDeviceMemorySize(getDevice(deviceId));
           const std::string bytesStr = stringifyBytes(bytes);
+          const std::string arch     = getDeviceArch(getDevice(deviceId));
 
           OCCA_CUDA_ERROR("Getting Device Name",
                           cuDeviceGetName(deviceName, 1024, deviceId));
 
           section
             .add("Device Name", deviceName)
             .add("Device ID"  , toString(deviceId))
+            .add("Arch"       , arch)
             .add("Memory"     , bytesStr)
             .addDivider();
         }

diff --git a/src/occa/internal/modes/cuda/utils.cpp b/src/occa/internal/modes/cuda/utils.cpp
@@ -52,6 +52,35 @@ namespace occa {
       return ss.str();
     }
 
+    void getDeviceArchVersion(CUdevice device,
+                              int& archMajorVersion,
+                              int& archMinorVersion) {
+#if CUDA_VERSION < 5000
+      OCCA_CUDA_ERROR("Device: Getting CUDA device arch",
+                      cuDeviceComputeCapability(&archMajorVersion,
+                                                &archMinorVersion,
+                                                device));
+#else
+      OCCA_CUDA_ERROR("Device: Getting CUDA device major version",
+                      cuDeviceGetAttribute(&archMajorVersion,
+                                           CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
+                                           device));
+      OCCA_CUDA_ERROR("Device: Getting CUDA device minor version",
+                      cuDeviceGetAttribute(&archMinorVersion,
+                                           CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
+                                           device));
+#endif
+    }
+
+    std::string getDeviceArch(CUdevice device) {
+      int archMajorVersion=0, archMinorVersion=0;
+      getDeviceArchVersion(device, archMajorVersion, archMinorVersion);
+      std::string arch = std::string("sm_");
+      arch += std::to_string(archMajorVersion);
+      arch += std::to_string(archMinorVersion);
+      return arch;
+    }
+
     void enablePeerToPeer(CUcontext context) {
 #if CUDA_VERSION >= 4000
       OCCA_CUDA_ERROR("Enabling Peer-to-Peer",

diff --git a/src/occa/internal/modes/cuda/utils.hpp b/src/occa/internal/modes/cuda/utils.hpp
@@ -20,6 +20,12 @@ namespace occa {
 
     std::string getVersion();
 
+    void getDeviceArchVersion(CUdevice device,
+                              int& archMajorVersion,
+                              int& archMinorVersion);
+
+    std::string getDeviceArch(CUdevice device);
+
     void enablePeerToPeer(CUcontext context);
     void checkPeerToPeer(CUdevice destDevice,
                          CUdevice srcDevice);

diff --git a/src/occa/internal/modes/dpcpp/device.cpp b/src/occa/internal/modes/dpcpp/device.cpp
@@ -49,6 +49,8 @@ namespace occa
 
       occa::json &kernelProps = properties["kernel"];
       setCompilerLinkerOptions(kernelProps);
+
+      arch = dpcppDevice.get_info<::sycl::info::device::name>();
     }
 
     hash_t device::hash() const
@@ -57,7 +59,7 @@ namespace occa
       {
         std::stringstream ss;
         auto p = dpcppDevice.get_platform();
-        ss << "platform name: " << p.get_info<::sycl::info::platform::name>() 
+        ss << "platform name: " << p.get_info<::sycl::info::platform::name>()
           << " platform vendor: " << p.get_info<::sycl::info::platform::vendor>()
           << " platform version: " << p.get_info<::sycl::info::platform::version>()
           << " device name: " << dpcppDevice.get_info<::sycl::info::device::name>()
@@ -82,8 +84,8 @@ namespace occa
     //---[ Stream ]---------------------
     modeStream_t *device::createStream(const occa::json &props)
     {
-      ::sycl::queue q(dpcppContext, 
-                      dpcppDevice, 
+      ::sycl::queue q(dpcppContext,
+                      dpcppDevice,
                       {::sycl::property::queue::enable_profiling{},
                       ::sycl::property::queue::in_order{}
                       });
@@ -96,7 +98,7 @@ namespace occa
       return new stream(this, props, q);
     }
 
-    // Uses a oneAPI extension to enqueue a barrier. 
+    // Uses a oneAPI extension to enqueue a barrier.
     // When ombined with in-order queues, this provides
     // the execution required for `streamTag`s.
     occa::streamTag device::tagStream()
@@ -123,7 +125,7 @@ namespace occa
       return (dpcppEndTag.endTime() - dpcppStartTag.endTime());
     }
 
-    
+
     //==================================
 
     //---[ Kernel ]---------------------
@@ -237,7 +239,7 @@ namespace occa
       } else if (verbose) {
         io::stdout << "Output:\n\n" << commandOutput << "\n";
       }
-      
+
       io::sync(binaryFilename);
     }
 
@@ -280,7 +282,7 @@ namespace occa
         arguments.erase(arguments.begin());
 
         occa::functionPtr_t kernel_function = sys::dlsym(dl_handle, metadata.name);
-       
+
         kernel *dpcppKernel = new dpcpp::kernel(this,
                                metadata.name,
                                sourceFilename,