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 83 def __str__(self): 84 output = config.package.Package.__str__(self) 85 if hasattr(self,'hipArch'): 86 if self.unifiedMemory: uminfo = ' with unified memory' 87 else: uminfo = '' 88 output += ' HIP arch: '+ self.hipArch + uminfo + '\n' 89 return output 90 91 def checkSizeofVoidP(self): 92 '''Checks if the HIPC compiler agrees with the C compiler on what size of void * should be''' 93 self.log.write('Checking if sizeof(void*) in HIP is the same as with regular compiler\n') 94 size = self.types.checkSizeof('void *', (8, 4), lang='HIP', save=False) 95 if size != self.types.sizes['void-p']: 96 raise RuntimeError('HIP Error: sizeof(void*) with HIP compiler is ' + str(size) + ' which differs from sizeof(void*) with C compiler') 97 return 98 99 def configureTypes(self): 100 import config.setCompilers 101 if not self.getDefaultPrecision() in ['double', 'single']: 102 raise RuntimeError('Must use either single or double precision with HIP') 103 self.checkSizeofVoidP() 104 return 105 106 def checkHIPCDoubleAlign(self): 107 if 'known-hip-align-double' in self.argDB: 108 if not self.argDB['known-hip-align-double']: 109 raise RuntimeError('HIP error: PETSc currently requires that HIP double alignment match the C compiler') 110 else: 111 typedef = 'typedef struct {double a; int b;} teststruct;\n' 112 hip_size = self.types.checkSizeof('teststruct', (16, 12), lang='HIP', codeBegin=typedef, save=False) 113 c_size = self.types.checkSizeof('teststruct', (16, 12), lang='C', codeBegin=typedef, save=False) 114 if c_size != hip_size: 115 raise RuntimeError('HIP compiler error: memory alignment doesn\'t match C compiler (try adding -malign-double to compiler options)') 116 return 117 118 def setFullPathHIPC(self): 119 self.pushLanguage('HIP') 120 HIPC = self.getCompiler() 121 self.popLanguage() 122 self.getExecutable(HIPC,getFullPath=1,resultName='fullPathHIPC') 123 if not hasattr(self,'fullPathHIPC'): 124 raise RuntimeError('Unable to locate the HIPC compiler') 125 126 def getSearchDirectories(self): 127 # Package.getSearchDirectories() return '' by default, so that HIPC's default include path could 128 # be checked. But here we lower priority of '', so that once we validated a header path, it will 129 # be added to HIP_INCLUDE. Other compilers, ex. CC or CXX, might need this path for compilation. 130 yield os.path.dirname(os.path.dirname(self.fullPathHIPC)) # yield /opt/rocm from /opt/rocm/bin/hipcc 131 yield '' 132 133 def configureLibrary(self): 134 self.setFullPathHIPC() 135 self.setHasROCTX() 136 config.package.Package.configureLibrary(self) 137 self.getExecutable('hipconfig',getFullPath=1,resultName='hip_config') 138 if hasattr(self,'hip_config'): 139 try: 140 self.platform = config.package.Package.executeShellCommand([self.hip_config,'--platform'],log=self.log)[0] 141 except RuntimeError: 142 pass 143 144 # Handle the platform issues 145 if not hasattr(self,'platform'): 146 if 'HIP_PLATFORM' in os.environ: 147 self.platform = os.environ['HIP_PLATFORM'] 148 elif hasattr(self,'systemNvcc'): 149 self.platform = 'nvidia' 150 else: 151 self.platform = 'amd' 152 153 self.libraries.pushLanguage('HIP') 154 self.addDefine('HAVE_HIP','1') 155 self.addDefine('HAVE_CUPM','1') # Have either CUDA or HIP 156 if self.platform in ['nvcc','nvidia']: 157 self.pushLanguage('CUDA') 158 petscNvcc = self.getCompiler() 159 cudaFlags = self.getCompilerFlags() 160 self.popLanguage() 161 self.getExecutable(petscNvcc,getFullPath=1,resultName='systemNvcc') 162 if hasattr(self,'systemNvcc'): 163 nvccDir = os.path.dirname(self.systemNvcc) 164 cudaDir = os.path.split(nvccDir)[0] 165 else: 166 raise RuntimeError('Unable to locate CUDA NVCC compiler') 167 self.includedir = ['include',os.path.join(cudaDir,'include')] 168 self.delDefine('HAVE_CUDA') 169 self.addDefine('HAVE_HIPCUDA',1) 170 self.framework.addDefine('__HIP_PLATFORM_NVCC__',1) # deprecated from 4.3.0 171 self.framework.addDefine('__HIP_PLATFORM_NVIDIA__',1) 172 else: 173 self.addDefine('HAVE_HIPROCM',1) 174 if self.hasROCTX: 175 self.framework.addDefine('HAVE_ROCTX',1) 176 self.framework.addDefine('__HIP_PLATFORM_HCC__',1) # deprecated from 4.3.0 177 self.framework.addDefine('__HIP_PLATFORM_AMD__',1) 178 if 'with-hip-arch' in self.framework.clArgDB: 179 if self.argDB['with-hip-arch'].split('_')[-1].lower() == 'apu': # take care of user input gfx942_apu and the like 180 self.unifiedMemory = True 181 self.hipArch = self.argDB['with-hip-arch'].split('_')[0] 182 else: 183 self.getExecutable('rocminfo',getFullPath=1) 184 if hasattr(self,'rocminfo'): 185 try: 186 (out, err, ret) = Configure.executeShellCommand(self.rocminfo + ' | grep -e "AMD Instinct MI300A " -e " gfx" ',timeout = 60, log = self.log, threads = 1) 187 except Exception as e: 188 self.log.write('ROCM utility ' + self.rocminfo + ' failed: '+str(e)+'\n') 189 else: 190 try: 191 s = set([i for i in out.split() if 'gfx' in i]) 192 self.hipArch = list(s)[0] 193 uminfo = '' 194 if 'AMD Instinct MI300A' in out: 195 self.unifiedMemory = True 196 uminfo = ' with unified memory' 197 self.log.write('ROCM utility ' + self.rocminfo + ' said the HIP arch is ' + self.hipArch + uminfo + '\n') 198 except: 199 self.log.write('Unable to parse the ROCM utility ' + self.rocminfo + '\n') 200 if hasattr(self,'hipArch'): 201 self.hipArch.lower() # to have a uniform format even if user set hip arch in weird cases 202 if not self.hipArch.startswith('gfx'): 203 raise RuntimeError('HIP arch name ' + self.hipArch + ' is not in the supported gfxnnn or gfxnnn_apu format') 204 # See https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html, even for MI300A, the LLVM target name is still gfx942 205 self.setCompilers.HIPFLAGS += ' --offload-arch=' + self.hipArch +' ' 206 else: 207 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') 208 209 # Record rocBlas and rocSparse directories as they are needed by Kokkos-Kernels HIP TPL, so that we can handle 210 # a weird (but valid) case --with-hipcc=/opt/rocm-4.5.2/hip/bin/hipcc --with-hip-dir=/opt/rocm-4.5.2 (which 211 # should be better written as --with-hipcc=/opt/rocm-4.5.2/bin/hipcc --with-hip-dir=/opt/rocm-4.5.2 or simply 212 # --with-hip-dir=/opt/rocm-4.5.2) 213 if self.directory: 214 self.hipDir = self.directory 215 else: # directory is '', indicating we are using the compiler's default, so the last resort is to guess the dir from hipcc 216 self.hipDir = os.path.dirname(os.path.dirname(self.fullPathHIPC)) # Ex. peel /opt/rocm-4.5.2/bin/hipcc twice 217 218 self.rocBlasDir = self.hipDir 219 self.rocSparseDir = self.hipDir 220 #self.checkHIPDoubleAlign() 221 self.configureTypes() 222 self.libraries.popLanguage() 223 return 224