1# Portions of this code are under: 2# Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. 3 4import config.package 5import os 6 7class Configure(config.package.Package): 8 def __init__(self, framework): 9 config.package.Package.__init__(self, framework) 10 11 self.minversion = '5.0.0' 12 # Check version from rocm-core here, as HIP_VERSION_PATCH (e.g., 31061 from hip_version.h) is not necessarily the AMD advertised patch version, e.g., in 5.6.0 13 self.versionname = 'ROCM_VERSION_MAJOR.ROCM_VERSION_MINOR.ROCM_VERSION_PATCH' 14 self.versioninclude = ['rocm-core/rocm_version.h', 'rocm_version.h'] 15 self.requiresversion = 1 16 self.functions = ['hipInit'] 17 self.includes = ['hip/hip_runtime.h'] 18 self.includedir = ['include'] 19 # PETSc does not use hipsparse or hipblas, but dependencies can (e.g., magma) 20 self.liblist = [['libhipsparse.a','libhipblas.a','libhipsolver.a','librocsparse.a','librocsolver.a','librocblas.a','librocrand.a','libamdhip64.a','libhsa-runtime64.a'], 21 ['hipsparse.lib','hipblas.lib','hipsolver.lib','rocsparse.lib','rocsolver.lib','rocblas.lib','rocrand.lib','amdhip64.lib','hsa-runtime64.lib'],] 22 self.precisions = ['single','double'] 23 self.buildLanguages = ['HIP'] 24 self.devicePackage = 1 25 self.fullPathHIPC = '' 26 self.hasROCTX = None 27 self.unifiedMemory = False 28 self.skipMPIDependency= 1 29 return 30 31 def setupHelp(self, help): 32 import nargs 33 config.package.Package.setupHelp(self, help) 34 help.addArgument('HIP', '-with-hip-arch', nargs.ArgString(None, None, 'AMD GPU architecture for code generation, for example gfx908, (this may be used by external packages)')) 35 return 36 37 def setupDependencies(self, framework): 38 config.package.Package.setupDependencies(self, framework) 39 self.setCompilers = framework.require('config.setCompilers',self) 40 self.headers = framework.require('config.headers',self) 41 return 42 43 def setHasROCTX(self): 44 if self.hasROCTX is not None: 45 return 46 47 if not hasattr(self,'platform'): 48 if 'HIP_PLATFORM' in os.environ: 49 self.platform = os.environ['HIP_PLATFORM'] 50 elif hasattr(self,'systemNvcc'): 51 self.platform = 'nvidia' 52 else: 53 self.platform = 'amd' 54 55 if self.platform == 'amd': 56 for prefix in self.getSearchDirectories(): 57 if not self.version_tuple: 58 self.checkVersion() 59 if self.version_tuple[0] >= 6 and self.version_tuple[1] >= 4: 60 if os.path.exists(os.path.join(prefix, 'rocprofiler-sdk-roctx', 'roctx.h')): 61 self.includes = ['hip/hip_runtime.h', 'rocprofiler-sdk-roctx/roctx.h'] 62 self.liblist[0] += ['librocprofiler-sdk-roctx.a'] 63 self.liblist[1] += ['rocprofiler-sdk-roctx.lib'] 64 self.hasROCTX = True 65 elif self.version_tuple[0] >= 6: 66 if os.path.exists(os.path.join(prefix, 'include', 'roctracer')): 67 self.includes = ['hip/hip_runtime.h', 'roctracer/roctx.h'] 68 self.liblist[0] += ['libroctx64.a'] 69 self.liblist[1] += ['roctx64.lib'] 70 self.hasROCTX = True 71 else: 72 if os.path.exists(os.path.join(prefix, 'roctracer')): 73 self.includes = ['hip/hip_runtime.h', 'roctx.h'] 74 self.includedir = ['include', os.path.join('roctracer', 'include')] 75 self.liblist[0] += ['libroctx64.a'] 76 self.liblist[1] += ['roctx64.lib'] 77 self.hasROCTX = True 78 79 if not self.hasROCTX: 80 self.hasROCTX = False 81 82 def __str__(self): 83 output = config.package.Package.__str__(self) 84 if hasattr(self,'hipArch'): 85 if self.unifiedMemory: uminfo = ' with unified memory' 86 else: uminfo = '' 87 output += ' HIP arch: '+ self.hipArch + uminfo + '\n' 88 return output 89 90 def checkSizeofVoidP(self): 91 '''Checks if the HIPC compiler agrees with the C compiler on what size of void * should be''' 92 self.log.write('Checking if sizeof(void*) in HIP is the same as with regular compiler\n') 93 size = self.types.checkSizeof('void *', (8, 4), lang='HIP', save=False) 94 if size != self.types.sizes['void-p']: 95 raise RuntimeError('HIP Error: sizeof(void*) with HIP compiler is ' + str(size) + ' which differs from sizeof(void*) with C compiler') 96 return 97 98 def configureTypes(self): 99 import config.setCompilers 100 if not self.getDefaultPrecision() in ['double', 'single']: 101 raise RuntimeError('Must use either single or double precision with HIP') 102 self.checkSizeofVoidP() 103 return 104 105 def checkHIPCDoubleAlign(self): 106 if 'known-hip-align-double' in self.argDB: 107 if not self.argDB['known-hip-align-double']: 108 raise RuntimeError('HIP error: PETSc currently requires that HIP double alignment match the C compiler') 109 else: 110 typedef = 'typedef struct {double a; int b;} teststruct;\n' 111 hip_size = self.types.checkSizeof('teststruct', (16, 12), lang='HIP', codeBegin=typedef, save=False) 112 c_size = self.types.checkSizeof('teststruct', (16, 12), lang='C', codeBegin=typedef, save=False) 113 if c_size != hip_size: 114 raise RuntimeError('HIP compiler error: memory alignment doesn\'t match C compiler (try adding -malign-double to compiler options)') 115 return 116 117 def setFullPathHIPC(self): 118 self.pushLanguage('HIP') 119 HIPC = self.getCompiler() 120 self.popLanguage() 121 self.getExecutable(HIPC,getFullPath=1,resultName='fullPathHIPC') 122 if not hasattr(self,'fullPathHIPC'): 123 raise RuntimeError('Unable to locate the HIPC compiler') 124 125 def getSearchDirectories(self): 126 # Package.getSearchDirectories() return '' by default, so that HIPC's default include path could 127 # be checked. But here we lower priority of '', so that once we validated a header path, it will 128 # be added to HIP_INCLUDE. Other compilers, ex. CC or CXX, might need this path for compilation. 129 yield os.path.dirname(os.path.dirname(self.fullPathHIPC)) # yield /opt/rocm from /opt/rocm/bin/hipcc 130 yield '' 131 132 def configureLibrary(self): 133 self.setFullPathHIPC() 134 self.setHasROCTX() 135 config.package.Package.configureLibrary(self) 136 self.getExecutable('hipconfig',getFullPath=1,resultName='hip_config') 137 if hasattr(self,'hip_config'): 138 try: 139 self.platform = config.package.Package.executeShellCommand([self.hip_config,'--platform'],log=self.log)[0] 140 except RuntimeError: 141 pass 142 143 # Handle the platform issues 144 if not hasattr(self,'platform'): 145 if 'HIP_PLATFORM' in os.environ: 146 self.platform = os.environ['HIP_PLATFORM'] 147 elif hasattr(self,'systemNvcc'): 148 self.platform = 'nvidia' 149 else: 150 self.platform = 'amd' 151 152 self.libraries.pushLanguage('HIP') 153 self.addDefine('HAVE_CUPM','1') # Have either CUDA or HIP 154 if self.platform in ['nvcc','nvidia']: 155 self.pushLanguage('CUDA') 156 petscNvcc = self.getCompiler() 157 cudaFlags = self.getCompilerFlags() 158 self.popLanguage() 159 self.getExecutable(petscNvcc,getFullPath=1,resultName='systemNvcc') 160 if hasattr(self,'systemNvcc'): 161 nvccDir = os.path.dirname(self.systemNvcc) 162 cudaDir = os.path.split(nvccDir)[0] 163 else: 164 raise RuntimeError('Unable to locate CUDA NVCC compiler') 165 self.includedir = ['include',os.path.join(cudaDir,'include')] 166 self.delDefine('HAVE_CUDA') 167 self.addDefine('HAVE_HIPCUDA',1) 168 self.framework.addDefine('__HIP_PLATFORM_NVCC__',1) # deprecated from 4.3.0 169 self.framework.addDefine('__HIP_PLATFORM_NVIDIA__',1) 170 else: 171 self.addDefine('HAVE_HIPROCM',1) 172 if self.hasROCTX: 173 self.framework.addDefine('HAVE_ROCTX',1) 174 self.framework.addDefine('__HIP_PLATFORM_HCC__',1) # deprecated from 4.3.0 175 self.framework.addDefine('__HIP_PLATFORM_AMD__',1) 176 if 'with-hip-arch' in self.framework.clArgDB: 177 if self.argDB['with-hip-arch'].split('_')[-1].lower() == 'apu': # take care of user input gfx942_apu and the like 178 self.unifiedMemory = True 179 self.hipArch = self.argDB['with-hip-arch'].split('_')[0] 180 else: 181 self.getExecutable('rocminfo',getFullPath=1) 182 if hasattr(self,'rocminfo'): 183 try: 184 (out, err, ret) = Configure.executeShellCommand(self.rocminfo + ' | grep -e "AMD Instinct MI300A " -e " gfx" ',timeout = 60, log = self.log, threads = 1) 185 except Exception as e: 186 self.log.write('ROCM utility ' + self.rocminfo + ' failed: '+str(e)+'\n') 187 else: 188 try: 189 s = set([i for i in out.split() if 'gfx' in i]) 190 self.hipArch = list(s)[0] 191 uminfo = '' 192 if 'AMD Instinct MI300A' in out: 193 self.unifiedMemory = True 194 uminfo = ' with unified memory' 195 self.log.write('ROCM utility ' + self.rocminfo + ' said the HIP arch is ' + self.hipArch + uminfo + '\n') 196 except: 197 self.log.write('Unable to parse the ROCM utility ' + self.rocminfo + '\n') 198 if hasattr(self,'hipArch'): 199 self.hipArch.lower() # to have a uniform format even if user set hip arch in weird cases 200 if not self.hipArch.startswith('gfx'): 201 raise RuntimeError('HIP arch name ' + self.hipArch + ' is not in the supported gfxnnn or gfxnnn_apu format') 202 # See https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html, even for MI300A, the LLVM target name is still gfx942 203 self.setCompilers.HIPFLAGS += ' --offload-arch=' + self.hipArch +' ' 204 else: 205 raise RuntimeError('You must set --with-hip-arch=gfx942_apu, gfx942, gfx900, gfx906, gfx908, gfx90a etc or make ROCM utility "rocminfo" available on your PATH') 206 207 # Record rocBlas and rocSparse directories as they are needed by Kokkos-Kernels HIP TPL, so that we can handle 208 # a weird (but valid) case --with-hipcc=/opt/rocm-4.5.2/hip/bin/hipcc --with-hip-dir=/opt/rocm-4.5.2 (which 209 # should be better written as --with-hipcc=/opt/rocm-4.5.2/bin/hipcc --with-hip-dir=/opt/rocm-4.5.2 or simply 210 # --with-hip-dir=/opt/rocm-4.5.2) 211 if self.directory: 212 self.hipDir = self.directory 213 else: # directory is '', indicating we are using the compiler's default, so the last resort is to guess the dir from hipcc 214 self.hipDir = os.path.dirname(os.path.dirname(self.fullPathHIPC)) # Ex. peel /opt/rocm-4.5.2/bin/hipcc twice 215 216 self.rocBlasDir = self.hipDir 217 self.rocSparseDir = self.hipDir 218 #self.checkHIPDoubleAlign() 219 self.configureTypes() 220 self.libraries.popLanguage() 221 return 222