source: CLRX/CLRadeonExtender/trunk/samples/CLUtils.cpp @ 3437

Last change on this file since 3437 was 3437, checked in by matszpk, 2 years ago

CLRadeonExtender: Fixed device detection for new Mesa3D 17.2.xx.

File size: 20.6 KB
Line 
1/*
2 *  CLRadeonExtender - Unofficial OpenCL Radeon Extensions Library
3 *  Copyright (C) 2014-2017 Mateusz Szpakowski
4 *
5 *  This library is free software; you can redistribute it and/or
6 *  modify it under the terms of the GNU Lesser General Public
7 *  License as published by the Free Software Foundation; either
8 *  version 2.1 of the License, or (at your option) any later version.
9 *
10 *  This library is distributed in the hope that it will be useful,
11 *  but WITHOUT ANY WARRANTY; without even the implied warranty of
12 *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
13 *  Lesser General Public License for more details.
14 *
15 *  You should have received a copy of the GNU Lesser General Public
16 *  License along with this library; if not, write to the Free Software
17 *  Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA
18 */
19
20#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
21#define CL_USE_DEPRECATED_OPENCL_2_0_APIS
22
23#include <algorithm>
24#include <iostream>
25#include <fstream>
26#include <vector>
27#include <cstring>
28#include <memory>
29#include <CL/cl.h>
30#include <CLRX/amdasm/Assembler.h>
31#include "CLUtils.h"
32
33using namespace CLRX;
34
35static char* stripCString(char* str)
36{
37    while (*str==' ') str++;
38    char* last = str+::strlen(str);
39    while (last!=str && (*last==0||*last==' '))
40        last--;
41    if (*last!=0) last[1] = 0;
42    return str;
43}
44
45/* parse args from command line and handle options:
46 * print help, list OpenCL devices, get choosen device and choosen OpenCL standard
47 * return true if sample should exit */
48bool CLFacade::parseArgs(const char* progName, const char* usagePart, int argc,
49                  const char** argv, cl_uint& deviceIndex, cxuint& useCL)
50{
51    if (argc >= 2 && ::strcmp(argv[1], "-?")==0)
52    {
53        std::cout << "Usage: " << progName << " [DEVICE_INDEX[cl1|old|cl2]] " <<
54                usagePart << "\n" "Print device list: " << progName << " -L" << "\n"
55                "Print help: " << progName << " -?\n"
56                "'cl2' after DEVICE_INDEX enables AMD OpenCL 2.0 mode\n"
57                "'cl1' or 'old' after DEVICE_INDEX force old AMD OpenCL 1.2 mode"
58                << std::endl;
59        return true;
60    }
61   
62    cl_uint platformsNum;
63    std::unique_ptr<cl_platform_id[]> platforms;
64    cl_int error = 0;
65    error = clGetPlatformIDs(0, nullptr, &platformsNum);
66    if (error != CL_SUCCESS)
67        throw CLError(error, "clGetPlatformIDs");
68    platforms.reset(new cl_platform_id[platformsNum]);
69    error = clGetPlatformIDs(platformsNum, platforms.get(), nullptr);
70    if (error != CL_SUCCESS)
71        throw CLError(error, "clGetPlatformIDs");
72   
73    cl_platform_id choosenPlatform = nullptr;
74    /// find platform with AMD or GalliumCompute devices
75    for (cl_uint i = 0; i < platformsNum; i++)
76    {
77        size_t platformNameSize;
78        std::unique_ptr<char[]> platformName;
79        error = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, 0, nullptr,
80                           &platformNameSize);
81        if (error != CL_SUCCESS)
82            throw CLError(error, "clGetPlatformInfo");
83        platformName.reset(new char[platformNameSize]);
84        error = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, platformNameSize,
85                   platformName.get(), nullptr);
86        if (error != CL_SUCCESS)
87            throw CLError(error, "clGetPlatformInfo");
88       
89        // find correct platform (supported GalliumCompute and AMD APP)
90        const char* splatformName = stripCString(platformName.get());
91        if (::strcmp(splatformName, "AMD Accelerated Parallel Processing")==0 ||
92            ::strcmp(splatformName, "Clover")==0)
93        {
94            choosenPlatform = platforms[i];
95            break;
96        }
97    }
98   
99    if (choosenPlatform==nullptr)
100        throw Exception("PlatformNotFound");
101   
102    if (argc >= 2 && ::strcmp(argv[1], "-L")==0)
103    {
104        // list devices, before it get GPU devices
105        cl_uint devicesNum;
106        std::unique_ptr<cl_device_id[]> devices;
107        error = clGetDeviceIDs(choosenPlatform, CL_DEVICE_TYPE_GPU, 0,
108                               nullptr, &devicesNum);
109        if (error != CL_SUCCESS)
110            throw CLError(error, "clGetDeviceIDs");
111       
112        devices.reset(new cl_device_id[devicesNum]);
113        error = clGetDeviceIDs(choosenPlatform, CL_DEVICE_TYPE_GPU,
114                        devicesNum, devices.get(), nullptr);
115        if (error != CL_SUCCESS)
116            throw CLError(error, "clGetDeviceIDs");
117       
118        for (cl_uint i = 0; i < devicesNum; i++)
119        {
120            cl_device_id device = devices[i];
121            // get device and print that
122            size_t deviceNameSize;
123            std::unique_ptr<char[]> deviceName;
124            error = clGetDeviceInfo(device, CL_DEVICE_NAME, 0, nullptr, &deviceNameSize);
125            if (error != CL_SUCCESS)
126                throw CLError(error, "clGetDeviceInfoName");
127            deviceName.reset(new char[deviceNameSize]);
128            error = clGetDeviceInfo(device, CL_DEVICE_NAME, deviceNameSize,
129                                     deviceName.get(), nullptr);
130            if (error != CL_SUCCESS)
131                throw CLError(error, "clGetDeviceInfoName");
132            std::cout << "Device: " << i << " - " << deviceName.get() << "\n";
133        }
134        std::cout.flush();
135        return true;
136    }
137    else if (argc >= 2)
138    {
139        const char* end;
140        useCL = 0;
141        deviceIndex = cstrtovCStyle<cl_uint>(argv[1], nullptr, end);
142        if (strcasecmp(end, "cl2")==0)
143            useCL = 2;
144        else if (strcasecmp(end, "cl1")==0 || strcasecmp(end, "old")==0)
145            useCL = 1;
146    }
147    return false;
148}
149
150static const char* binaryFormatNamesTbl[] =
151{
152    "AMD OpenCL 1.2", "GalliumCompute", "Raw code", "AMD OpenCL 2.0"
153};
154
155CLFacade::CLFacade(cl_uint deviceIndex, const char* sourceCode, const char* kernelNames,
156            cxuint useCL)
157{
158try
159{
160    context = nullptr;
161    queue = nullptr;
162    program = nullptr;
163   
164    cl_uint platformsNum;
165    std::unique_ptr<cl_platform_id[]> platforms;
166    cl_int error = 0;
167    error = clGetPlatformIDs(0, nullptr, &platformsNum);
168    if (error != CL_SUCCESS)
169        throw CLError(error, "clGetPlatformIDs");
170    platforms.reset(new cl_platform_id[platformsNum]);
171    error = clGetPlatformIDs(platformsNum, platforms.get(), nullptr);
172    if (error != CL_SUCCESS)
173        throw CLError(error, "clGetPlatformIDs");
174   
175    cxuint amdappVersion = 0;
176   
177    BinaryFormat binaryFormat = BinaryFormat::GALLIUM;
178    cl_platform_id choosenPlatform = nullptr;
179    bool defaultCL2ForDriver = false;
180    /// find platform with AMD devices
181    for (cl_uint i = 0; i < platformsNum; i++)
182    {
183        size_t platformNameSize;
184        std::unique_ptr<char[]> platformName;
185        error = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, 0, nullptr,
186                           &platformNameSize);
187        if (error != CL_SUCCESS)
188            throw CLError(error, "clGetPlatformInfo");
189        platformName.reset(new char[platformNameSize]);
190        error = clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, platformNameSize,
191                   platformName.get(), nullptr);
192        if (error != CL_SUCCESS)
193            throw CLError(error, "clGetPlatformInfo");
194       
195        const char* splatformName = stripCString(platformName.get());
196        if (::strcmp(splatformName, "AMD Accelerated Parallel Processing")==0 ||
197            ::strcmp(splatformName, "Clover")==0)
198        {
199            choosenPlatform = platforms[i];
200            binaryFormat = ::strcmp(platformName.get(), "Clover")==0 ?
201                    BinaryFormat::GALLIUM : BinaryFormat::AMD;
202           
203            if (binaryFormat == BinaryFormat::AMD)
204            {
205                // get amdappVersion
206                size_t platformVersionSize;
207                std::unique_ptr<char[]> platformVersion;
208                error = clGetPlatformInfo(choosenPlatform, CL_PLATFORM_VERSION, 0, nullptr,
209                                        &platformVersionSize);
210                if (error != CL_SUCCESS)
211                    throw CLError(error, "clGetPlatformInfoVersion");
212                platformVersion.reset(new char[platformVersionSize]);
213                error = clGetPlatformInfo(choosenPlatform, CL_PLATFORM_VERSION,
214                                platformVersionSize, platformVersion.get(), nullptr);
215                if (error != CL_SUCCESS)
216                    throw CLError(error, "clGetPlatformInfoVersion");
217               
218                const char* amdappPart = strstr(platformVersion.get(), "AMD-APP (");
219                if (amdappPart!=nullptr)
220                {
221                    // parse AMDAPP version
222                    try
223                    {
224                        const char* majorVerPart = amdappPart+9;
225                        const char* minorVerPart;
226                        const char* end;
227                        cxuint majorVersion = cstrtoui(majorVerPart, nullptr,
228                                        minorVerPart);
229                       
230                        if (*minorVerPart!=0)
231                        {
232                            minorVerPart++; // skip '.'
233                            cxuint minorVersion = cstrtoui(minorVerPart, nullptr, end);
234                            amdappVersion = majorVersion*100U + minorVersion;
235                        }
236                    }
237                    catch(const ParseException& ex)
238                    { } // ignore error
239                }
240            }
241           
242            if (binaryFormat == BinaryFormat::AMD && useCL==2)
243                binaryFormat = BinaryFormat::AMDCL2;
244            // for driver 2004.6 OpenCL 2.0 binary format is default
245            if (binaryFormat == BinaryFormat::AMD && amdappVersion >= 200406)
246                defaultCL2ForDriver = true;
247            break;
248        }
249    }
250   
251    if (choosenPlatform==nullptr)
252        throw Exception("PlatformNotFound");
253   
254    // find device
255    cl_uint devicesNum;
256    std::unique_ptr<cl_device_id[]> devices;
257    error = clGetDeviceIDs(choosenPlatform, CL_DEVICE_TYPE_GPU, 0, nullptr, &devicesNum);
258    if (error != CL_SUCCESS)
259        throw CLError(error, "clGetDeviceIDs");
260   
261    if (deviceIndex >= devicesNum)
262        throw CLError(0, "DeviceIndexOutOfRange");
263       
264    devices.reset(new cl_device_id[devicesNum]);
265    error = clGetDeviceIDs(choosenPlatform, CL_DEVICE_TYPE_GPU,
266                    devicesNum, devices.get(), nullptr);
267    if (error != CL_SUCCESS)
268        throw CLError(error, "clGetDeviceIDs");
269   
270    device = devices[deviceIndex];
271   
272    cl_uint bits = 32;
273    if (binaryFormat != BinaryFormat::GALLIUM)
274    {
275        // get address Bits from device info (for AMDAPP)
276        error = clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint),
277                                 &bits, nullptr);
278        if (error != CL_SUCCESS)
279            throw CLError(error, "clGetDeviceAddressBits");
280    }
281   
282    // get workGroupSize and Compute Units of device
283    error = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t),
284                             &maxWorkGroupSize, nullptr);
285    if (error != CL_SUCCESS)
286        throw CLError(error, "clGetDeviceMaxWorkGroupSize");
287    error = clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint),
288                             &computeUnits, nullptr);
289    if (error != CL_SUCCESS)
290        throw CLError(error, "clGetDeviceMaxComputeUnits");
291   
292    /// create context
293    context = clCreateContext(nullptr, 1, &device, nullptr, nullptr, &error);
294    if (context==nullptr)
295        throw CLError(error, "clCreateContext");
296   
297    // get device and print that
298    size_t deviceNameSize;
299    std::unique_ptr<char[]> deviceName;
300    error = clGetDeviceInfo(device, CL_DEVICE_NAME, 0, nullptr, &deviceNameSize);
301    if (error != CL_SUCCESS)
302        throw CLError(error, "clGetDeviceInfoName");
303    deviceName.reset(new char[deviceNameSize]);
304    error = clGetDeviceInfo(device, CL_DEVICE_NAME, deviceNameSize,
305                             deviceName.get(), nullptr);
306    if (error != CL_SUCCESS)
307        throw CLError(error, "clGetDeviceInfoName");
308    std::cout << "Device: " << deviceIndex << " - " << deviceName.get() << std::endl;
309   
310    // get device version - used for getting Mesa3D version and LLVM version
311    size_t deviceVersionSize;
312    std::unique_ptr<char[]> deviceVersion;
313    error = clGetDeviceInfo(device, CL_DEVICE_VERSION, 0, nullptr, &deviceVersionSize);
314    if (error != CL_SUCCESS)
315        throw CLError(error, "clGetDeviceInfoVersion");
316    deviceVersion.reset(new char[deviceVersionSize]);
317    error = clGetDeviceInfo(device, CL_DEVICE_VERSION, deviceVersionSize,
318                            deviceVersion.get(), nullptr);
319    if (error != CL_SUCCESS)
320        throw CLError(error, "clGetDeviceInfoVersion");
321   
322    // get bits from device name (LLVM version)
323    cxuint llvmVersion = 0;
324    cxuint mesaVersion = 0;
325   
326    if (binaryFormat == BinaryFormat::GALLIUM)
327    {
328        const char* llvmPart = strstr(deviceName.get(), "LLVM ");
329        if (llvmPart!=nullptr)
330        {
331            try
332            {
333                // parse LLVM version
334                const char* majorVerPart = llvmPart+5;
335                const char* minorVerPart;
336                const char* end;
337                cxuint majorVersion = cstrtoui(majorVerPart, nullptr, minorVerPart);
338                if (*minorVerPart!=0)
339                {
340                    minorVerPart++; // skip '.'
341                    cxuint minorVersion = cstrtoui(minorVerPart, nullptr, end);
342                    llvmVersion = majorVersion*10000U + minorVersion*100U;
343#if HAVE_64BIT
344                    if (majorVersion*10000U + minorVersion*100U >= 30900U)
345                        bits = 64; // use 64-bit
346#endif
347                }
348            }
349            catch(const ParseException& ex)
350            { } // ignore error
351        }
352       
353        const char* mesaPart = strstr(deviceVersion.get(), "Mesa ");
354        if (mesaPart==nullptr)
355            mesaPart = strstr(deviceVersion.get(), "MESA ");
356        if (mesaPart!=nullptr)
357        {
358            try
359            {
360                // parse Mesa3D version
361                const char* majorVerPart = mesaPart+5;
362                const char* minorVerPart;
363                const char* end;
364                cxuint majorVersion = cstrtoui(majorVerPart, nullptr, minorVerPart);
365                if (*minorVerPart!=0)
366                {
367                    minorVerPart++; // skip '.'
368                    cxuint minorVersion = cstrtoui(minorVerPart, nullptr, end);
369                    mesaVersion = majorVersion*10000U + minorVersion*100U;
370                }
371            }
372            catch(const ParseException& ex)
373            { } // ignore error
374        }
375    }
376    /* assemble source code */
377    /// determine device type
378    char* sdeviceName = stripCString(deviceName.get());
379    char* devNamePtr = sdeviceName;
380    if (binaryFormat==BinaryFormat::GALLIUM)
381    {
382        char* sptr = ::strstr(sdeviceName, "(AMD ");
383        // if form 'AMD Radeon xxx (AMD CODENAME /...)
384        if (sptr != nullptr) // if found 'AMD ';
385            devNamePtr = sptr+5;
386        else
387        {
388            // if form 'AMD CODENAME (....
389            sptr = ::strstr(sdeviceName, "AMD ");
390            if (sptr != nullptr) // if found 'AMD ';
391                devNamePtr = sptr+4;
392        }
393    }
394    char* devNameEnd = devNamePtr;
395    while (isAlnum(*devNameEnd)) devNameEnd++;
396    *devNameEnd = 0; // finish at first word
397    const GPUDeviceType devType = getGPUDeviceTypeFromName(devNamePtr);
398    /* change binary format to AMDCL2 if default for this driver version and
399     * architecture >= GCN 1.1 */
400    bool useLegacy = false;
401    if (defaultCL2ForDriver &&
402        getGPUArchitectureFromDeviceType(devType) >= GPUArchitecture::GCN1_1)
403    {
404        if (useCL!=1) // if not cl1/old
405            binaryFormat = BinaryFormat::AMDCL2;
406        else // use legacy
407            useLegacy = true;
408    }
409
410    std::cout << "BinaryFormat: " << binaryFormatNamesTbl[cxuint(binaryFormat)] << "\n"
411        "Bitness: " << bits << std::endl;
412   
413    /// create command queue
414    queue = clCreateCommandQueue(context, device, 0, &error);
415    if (queue==nullptr)
416        throw CLError(error, "clCreateCommandQueue");
417   
418    Array<cxbyte> binary;
419    {
420        /* assemble source code */
421        /// determine device type
422        ArrayIStream astream(::strlen(sourceCode), sourceCode);
423        // by default assembler put logs to stderr
424        Assembler assembler("", astream, 0, binaryFormat, devType);
425        assembler.set64Bit(bits==64);
426        // setting version (LLVM and driverVersion)
427        if (binaryFormat == BinaryFormat::GALLIUM && llvmVersion != 0)
428            assembler.setLLVMVersion(llvmVersion);
429        if (binaryFormat == BinaryFormat::GALLIUM && mesaVersion != 0)
430            assembler.setDriverVersion(mesaVersion);
431        else if ((binaryFormat == BinaryFormat::AMD ||
432                binaryFormat == BinaryFormat::AMDCL2) && amdappVersion != 0)
433            assembler.setDriverVersion(amdappVersion);
434        assembler.assemble();
435        assembler.writeBinary(binary);
436    }
437   
438    size_t binarySize = binary.size();
439    const cxbyte* binaryContent = binary.data();
440    program = clCreateProgramWithBinary(context, 1, &device, &binarySize,
441                        &binaryContent, nullptr, &error);
442    if (program==nullptr)
443        throw CLError(error, "clCreateProgramWithBinary");
444    // build program
445    error = clBuildProgram(program, 1, &device,
446               (binaryFormat==BinaryFormat::AMDCL2) ? "-cl-std=CL2.0" :
447               (useLegacy ? "-legacy" : ""),
448               nullptr, nullptr);
449    if (error != CL_SUCCESS)
450    {
451        /* get build logs */
452        size_t buildLogSize;
453        std::unique_ptr<char[]> buildLog;
454        cl_int lerror = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
455                           0, nullptr, &buildLogSize);
456        if (lerror == CL_SUCCESS)
457        {
458            buildLog.reset(new char[buildLogSize]);
459            lerror = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
460                           buildLogSize, buildLog.get(), nullptr);
461            if (lerror == CL_SUCCESS) // print build log
462                std::cerr << "BuildLog:\n" << buildLog.get() << std::endl;
463        }
464        throw CLError(error, "clBuildProgram");
465    }
466   
467    if (kernelNames!=nullptr)
468    try
469    {
470        for (const char* kn = kernelNames; *kn!=0;)
471        {
472            const char* knameStart = kn;
473            while (*kn!=0 && *kn!=' ') kn++;
474            std::string kernelName(knameStart, kn);
475            cl_kernel kernel = clCreateKernel(program, kernelName.c_str(), &error);
476            if (error != CL_SUCCESS)
477                throw CLError(error, "clCreateKernel");
478            kernels.push_back(kernel);
479            while (*kn==' ') kn++; // skip spaces
480        }
481    }
482    catch(...)
483    {
484        for (cl_kernel kernel: kernels)
485            clReleaseKernel(kernel);
486        throw;
487    }
488}
489catch(...)
490{
491    if (program!=nullptr)
492        clReleaseProgram(program);
493    if (queue!=nullptr)
494        clReleaseCommandQueue(queue);
495    if (context!=nullptr)
496        clReleaseContext(context);
497    throw;
498}
499}
500
501CLFacade::~CLFacade()
502{
503    for (cl_mem memObj: memObjects)
504        clReleaseMemObject(memObj);
505    for (cl_kernel kernel: kernels)
506        clReleaseKernel(kernel);
507    if (program!=nullptr)
508        clReleaseProgram(program);
509    if (queue!=nullptr)
510        clReleaseCommandQueue(queue);
511    if (context!=nullptr)
512        clReleaseContext(context);
513}
514
515// get work group size and work group size multiple from kernel
516void CLFacade::getKernelInfo(cl_kernel kernel, size_t& workGroupSize,
517               size_t& workGroupSizeMultiple)
518{
519    cl_int error = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE,
520                 sizeof(size_t), &workGroupSize, nullptr);
521    if (error != CL_SUCCESS)
522        throw CLError(error, "clGetKernelWorkGroupSize");
523    error = clGetKernelWorkGroupInfo(kernel, device,
524                 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
525                 sizeof(size_t), &workGroupSizeMultiple, nullptr);
526    if (error != CL_SUCCESS || workGroupSizeMultiple == 1)
527        workGroupSizeMultiple = 64; // fix for GalliumCompute
528}
529
530void CLFacade::callNDRangeKernel(cl_kernel kernel, cl_uint workDim, const size_t* offset,
531               const size_t* workSize, const size_t* localSize)
532{
533    cl_event event;
534    cl_int error = clEnqueueNDRangeKernel(queue, kernel, workDim, offset, workSize,
535                                   localSize, 0, nullptr, &event);
536    if (error != CL_SUCCESS)
537        throw CLError(error, "clEnqueueNDRangeKernel");
538    error = clWaitForEvents(1, &event); // waiting for finish kernel
539    if (error != CL_SUCCESS)
540    {
541        clReleaseEvent(event);
542        throw CLError(error, "clWaitForEvents");
543    }
544    clReleaseEvent(event);
545}
Note: See TracBrowser for help on using the repository browser.