From d9d9c6939911e8287f3b9c9c976aa9564d45375d Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 12 Feb 2020 13:16:06 -0800 Subject: [PATCH] Replace cl_* integral types with standard types. cl_bool -> bool cl_int -> int32_t cl_uint -> uint32_t cl_long -> int64_t cl_ulong -> uint64_t cl_float -> float cl_double -> double cl_bitfield -> uint64_t Change-Id: I840c8993b55f98f5b745d21e27f5f28233647a58 --- rocclr/device/blit.cpp | 16 +-- rocclr/device/blit.hpp | 2 +- rocclr/device/device.cpp | 12 +- rocclr/device/device.hpp | 172 +++++++++++++------------- rocclr/device/devkernel.hpp | 4 +- rocclr/device/devprogram.cpp | 6 +- rocclr/device/devprogram.hpp | 12 +- rocclr/device/gpu/gpublit.cpp | 72 +++++------ rocclr/device/gpu/gpucounters.hpp | 6 +- rocclr/device/gpu/gpudebugmanager.cpp | 6 +- rocclr/device/gpu/gpudebugmanager.hpp | 6 +- rocclr/device/gpu/gpudevice.cpp | 34 ++--- rocclr/device/gpu/gpudevice.hpp | 2 +- rocclr/device/gpu/gpukernel.cpp | 8 +- rocclr/device/gpu/gpukernel.hpp | 4 +- rocclr/device/gpu/gpuprogram.cpp | 4 +- rocclr/device/gpu/gpuvirtual.cpp | 8 +- rocclr/device/hwdebug.hpp | 4 +- rocclr/device/pal/palblit.cpp | 72 +++++------ rocclr/device/pal/palconstbuf.cpp | 2 +- rocclr/device/pal/palcounters.hpp | 6 +- rocclr/device/pal/paldebugmanager.cpp | 6 +- rocclr/device/pal/paldebugmanager.hpp | 6 +- rocclr/device/pal/paldevice.cpp | 34 ++--- rocclr/device/pal/paldevice.hpp | 2 +- rocclr/device/pal/palvirtual.cpp | 2 +- rocclr/device/rocm/rocblit.cpp | 80 ++++++------ rocclr/device/rocm/roccounters.cpp | 6 +- rocclr/device/rocm/roccounters.hpp | 6 +- rocclr/device/rocm/rocdevice.cpp | 16 +-- rocclr/device/rocm/rocvirtual.cpp | 2 +- rocclr/include/vdi_agent_amd.h | 44 +++---- rocclr/platform/agent.cpp | 32 ++--- rocclr/platform/agent.hpp | 14 +-- rocclr/platform/command.cpp | 20 +-- rocclr/platform/command.hpp | 56 ++++----- rocclr/platform/kernel.cpp | 2 +- rocclr/platform/kernel.hpp | 2 +- rocclr/platform/memory.cpp | 16 +-- rocclr/platform/memory.hpp | 6 +- rocclr/platform/program.cpp | 20 +-- rocclr/platform/program.hpp | 8 +- 42 files changed, 419 insertions(+), 419 deletions(-) diff --git a/rocclr/device/blit.cpp b/rocclr/device/blit.cpp index 9caf081738..7bbf73a192 100644 --- a/rocclr/device/blit.cpp +++ b/rocclr/device/blit.cpp @@ -596,17 +596,17 @@ bool HostBlitManager::fillImage(device::Memory& memory, const void* pattern, size_t devSlicePitch; void* newpattern = const_cast(pattern); - cl_float4 fFillColor; + float fFillColor[4]; // Converting a linear RGB floating-point color value to a normalized 8-bit unsigned integer sRGB // value so that the cpu path can treat sRGB as RGB for host transfer. if (memory.owner()->asImage()->getImageFormat().image_channel_order == CL_sRGBA) { float* fColor = static_cast(newpattern); - fFillColor.s[0] = sRGBmap(fColor[0]) / 255.0f; - fFillColor.s[1] = sRGBmap(fColor[1]) / 255.0f; - fFillColor.s[2] = sRGBmap(fColor[2]) / 255.0f; - fFillColor.s[3] = fColor[3]; - newpattern = static_cast(&fFillColor); + fFillColor[0] = sRGBmap(fColor[0]) / 255.0f; + fFillColor[1] = sRGBmap(fColor[1]) / 255.0f; + fFillColor[2] = sRGBmap(fColor[2]) / 255.0f; + fFillColor[3] = fColor[3]; + newpattern = static_cast(&fFillColor[0]); } // Map memory @@ -657,7 +657,7 @@ bool HostBlitManager::fillImage(device::Memory& memory, const void* pattern, return true; } -cl_uint HostBlitManager::sRGBmap(float fc) const { +uint32_t HostBlitManager::sRGBmap(float fc) const { double c = (double)fc; #ifdef ATI_OS_LINUX @@ -675,6 +675,6 @@ cl_uint HostBlitManager::sRGBmap(float fc) const { else c = (1055.0 / 1000.0) * pow(c, 5.0 / 12.0) - (55.0 / 1000.0); - return (cl_uint)(c * 255.0 + 0.5); + return (uint32_t)(c * 255.0 + 0.5); } } // namespace gpu diff --git a/rocclr/device/blit.hpp b/rocclr/device/blit.hpp index e8d1f2dbaf..7865a036f9 100644 --- a/rocclr/device/blit.hpp +++ b/rocclr/device/blit.hpp @@ -343,7 +343,7 @@ class HostBlitManager : public device::BlitManager { bool entire = false //!< Entire buffer will be updated ) const; - cl_uint sRGBmap(float fc) const; + uint32_t sRGBmap(float fc) const; protected: VirtualDevice& vDev_; //!< Virtual device object diff --git a/rocclr/device/device.cpp b/rocclr/device/device.cpp index 14d5910bd1..6b0930f0f2 100644 --- a/rocclr/device/device.cpp +++ b/rocclr/device/device.cpp @@ -389,10 +389,10 @@ size_t Device::numDevices(cl_device_type type, bool offlineDevices) { return result; } -bool Device::getDeviceIDs(cl_device_type deviceType, cl_uint numEntries, cl_device_id* devices, - cl_uint* numDevices, bool offlineDevices) { +bool Device::getDeviceIDs(cl_device_type deviceType, uint32_t numEntries, cl_device_id* devices, + uint32_t* numDevices, bool offlineDevices) { if (numDevices != nullptr && devices == nullptr) { - *numDevices = (cl_uint)amd::Device::numDevices(deviceType, offlineDevices); + *numDevices = (uint32_t)amd::Device::numDevices(deviceType, offlineDevices); return (*numDevices > 0) ? true : false; } assert(devices != nullptr && "check the code above"); @@ -404,7 +404,7 @@ bool Device::getDeviceIDs(cl_device_type deviceType, cl_uint numEntries, cl_devi } auto it = ret.cbegin(); - cl_uint count = std::min(numEntries, (cl_uint)ret.size()); + uint32_t count = std::min(numEntries, (uint32_t)ret.size()); while (count--) { *devices++ = as_cl(*it++); @@ -414,7 +414,7 @@ bool Device::getDeviceIDs(cl_device_type deviceType, cl_uint numEntries, cl_devi *devices++ = (cl_device_id)0; } - *not_null(numDevices) = (cl_uint)ret.size(); + *not_null(numDevices) = (uint32_t)ret.size(); return true; } @@ -673,7 +673,7 @@ bool ClBinary::createElfBinary(bool doencrypt, Program::type_t type) { } else { // char OpenCLVersion[256]; // size_t sz; - // cl_int ret= clGetPlatformInfo(AMD_PLATFORM, CL_PLATFORM_VERSION, 256, OpenCLVersion, &sz); + // int32_t ret= clGetPlatformInfo(AMD_PLATFORM, CL_PLATFORM_VERSION, 256, OpenCLVersion, &sz); // if (ret == CL_SUCCESS) { // buildVerInfo.append(OpenCLVersion, sz); // } diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index 1debd6b45a..83328007a8 100644 --- a/rocclr/device/device.hpp +++ b/rocclr/device/device.hpp @@ -194,14 +194,14 @@ struct Info : public amd::EmbeddedObject { cl_device_type type_; //! A unique device vendor identifier. - cl_uint vendorId_; + uint32_t vendorId_; //! The number of parallel compute cores on the compute device. - cl_uint maxComputeUnits_; + uint32_t maxComputeUnits_; //! Maximum dimensions that specify the global and local work-item IDs // used by the data-parallel execution model. - cl_uint maxWorkItemDimensions_; + uint32_t maxWorkItemDimensions_; //! Maximum number of work-items that can be specified in each dimension // to clEnqueueNDRangeKernel. @@ -218,58 +218,58 @@ struct Info : public amd::EmbeddedObject { //! Number of shader engines in physical GPU size_t numberOfShaderEngines; - //! cl_uint Preferred native vector width size for built-in scalar types + //! uint32_t Preferred native vector width size for built-in scalar types // that can be put into vectors. - cl_uint preferredVectorWidthChar_; - cl_uint preferredVectorWidthShort_; - cl_uint preferredVectorWidthInt_; - cl_uint preferredVectorWidthLong_; - cl_uint preferredVectorWidthFloat_; - cl_uint preferredVectorWidthDouble_; - cl_uint preferredVectorWidthHalf_; + uint32_t preferredVectorWidthChar_; + uint32_t preferredVectorWidthShort_; + uint32_t preferredVectorWidthInt_; + uint32_t preferredVectorWidthLong_; + uint32_t preferredVectorWidthFloat_; + uint32_t preferredVectorWidthDouble_; + uint32_t preferredVectorWidthHalf_; //! Returns the native ISA vector width. The vector width is defined as the // number of scalar elements that can be stored in the vector. - cl_uint nativeVectorWidthChar_; - cl_uint nativeVectorWidthShort_; - cl_uint nativeVectorWidthInt_; - cl_uint nativeVectorWidthLong_; - cl_uint nativeVectorWidthFloat_; - cl_uint nativeVectorWidthDouble_; - cl_uint nativeVectorWidthHalf_; + uint32_t nativeVectorWidthChar_; + uint32_t nativeVectorWidthShort_; + uint32_t nativeVectorWidthInt_; + uint32_t nativeVectorWidthLong_; + uint32_t nativeVectorWidthFloat_; + uint32_t nativeVectorWidthDouble_; + uint32_t nativeVectorWidthHalf_; //! Maximum configured engine clock frequency of the device in MHz. - cl_uint maxEngineClockFrequency_; + uint32_t maxEngineClockFrequency_; //! Maximum configured memory clock frequency of the device in MHz. - cl_uint maxMemoryClockFrequency_; + uint32_t maxMemoryClockFrequency_; //! Memory bus width in bits. - cl_uint vramBusBitWidth_; + uint32_t vramBusBitWidth_; //! Size of L2 Cache in bytes. - cl_uint l2CacheSize_; + uint32_t l2CacheSize_; //! Timestamp frequency in Hz. - cl_uint timeStampFrequency_; + uint32_t timeStampFrequency_; //! Describes the address spaces supported by the device. - cl_uint addressBits_; + uint32_t addressBits_; //! Max number of simultaneous image objects that can be read by a // kernel. - cl_uint maxReadImageArgs_; + uint32_t maxReadImageArgs_; //! Max number of simultaneous image objects that can be written to // by a kernel. - cl_uint maxWriteImageArgs_; + uint32_t maxWriteImageArgs_; //! Max number of simultaneous image objects that can be read/written to // by a kernel. - cl_uint maxReadWriteImageArgs_; + uint32_t maxReadWriteImageArgs_; //! Max size of memory object allocation in bytes. - cl_ulong maxMemAllocSize_; + uint64_t maxMemAllocSize_; //! Max width of 2D image in pixels. size_t image2DMaxWidth_; @@ -287,20 +287,20 @@ struct Info : public amd::EmbeddedObject { size_t image3DMaxDepth_; //! Describes whether images are supported - cl_bool imageSupport_; + bool imageSupport_; //! Max size in bytes of the arguments that can be passed to a kernel. size_t maxParameterSize_; //! Maximum number of samplers that can be used in a kernel. - cl_uint maxSamplers_; + uint32_t maxSamplers_; //! Describes the alignment in bits of the base address of any // allocated memory object. - cl_uint memBaseAddrAlign_; + uint32_t memBaseAddrAlign_; //! The smallest alignment in bytes which can be used for any data type. - cl_uint minDataTypeAlignSize_; + uint32_t minDataTypeAlignSize_; //! Describes single precision floating point capability of the device. cl_device_fp_config halfFPConfig_; @@ -311,53 +311,53 @@ struct Info : public amd::EmbeddedObject { cl_device_mem_cache_type globalMemCacheType_; //! Size of global memory cache line in bytes. - cl_uint globalMemCacheLineSize_; + uint32_t globalMemCacheLineSize_; //! Size of global memory cache in bytes. - cl_ulong globalMemCacheSize_; + uint64_t globalMemCacheSize_; //! Size of global device memory in bytes. - cl_ulong globalMemSize_; + uint64_t globalMemSize_; //! Max size in bytes of a constant buffer allocation. - cl_ulong maxConstantBufferSize_; + uint64_t maxConstantBufferSize_; //! Preferred size in bytes of a constant buffer allocation. - cl_ulong preferredConstantBufferSize_; + uint64_t preferredConstantBufferSize_; //! Max number of arguments declared - cl_uint maxConstantArgs_; + uint32_t maxConstantArgs_; //! This is used to determine the type of local memory that is available cl_device_local_mem_type localMemType_; //! Size of local memory arena in bytes. - cl_ulong localMemSize_; + uint64_t localMemSize_; //! If enabled, implies that all the memories, caches, registers etc. in // the device implement error correction. - cl_bool errorCorrectionSupport_; + bool errorCorrectionSupport_; //! CL_TRUE if the device and the host have a unified memory subsystem and // is CL_FALSE otherwise. - cl_bool hostUnifiedMemory_; + bool hostUnifiedMemory_; //! Describes the resolution of device timer. size_t profilingTimerResolution_; //! Timer starting point offset to Epoch. - cl_ulong profilingTimerOffset_; + uint64_t profilingTimerOffset_; //! CL_TRUE if device is a little endian device. - cl_bool littleEndian_; + bool littleEndian_; //! If enabled, implies that commands can be submitted to command-queues // created on this device. - cl_bool available_; + bool available_; //! if the implementation does not have a compiler available to compile // the program source. - cl_bool compilerAvailable_; + bool compilerAvailable_; //! Describes the execution capabilities of the device. cl_device_exec_capabilities executionCapabilities_; @@ -366,13 +366,13 @@ struct Info : public amd::EmbeddedObject { cl_device_svm_capabilities svmCapabilities_; //! Preferred alignment for OpenCL fine-grained SVM atomic types. - cl_uint preferredPlatformAtomicAlignment_; + uint32_t preferredPlatformAtomicAlignment_; //! Preferred alignment for OpenCL global atomic types. - cl_uint preferredGlobalAtomicAlignment_; + uint32_t preferredGlobalAtomicAlignment_; //! Preferred alignment for OpenCL local atomic types. - cl_uint preferredLocalAtomicAlignment_; + uint32_t preferredLocalAtomicAlignment_; //! Describes the command-queue properties supported of the host queue. cl_command_queue_properties queueProperties_; @@ -402,7 +402,7 @@ struct Info : public amd::EmbeddedObject { const char* extensions_; //! Returns if device linker is available - cl_bool linkerAvailable_; + bool linkerAvailable_; //! Returns the list of built-in kernels, supported by the device const char* builtInKernels_; @@ -415,21 +415,21 @@ struct Info : public amd::EmbeddedObject { //! Returns CL_TRUE if the devices preference is for the user to be //! responsible for synchronization - cl_bool preferredInteropUserSync_; + bool preferredInteropUserSync_; //! Returns maximum size of the internal buffer that holds the output //! of printf calls from a kernel size_t printfBufferSize_; //! Indicates maximum number of supported global atomic counters - cl_uint maxAtomicCounters_; + uint32_t maxAtomicCounters_; //! Returns the topology for the device cl_device_topology_amd deviceTopology_; //! Semaphore information - cl_uint maxSemaphores_; - cl_uint maxSemaphoreSize_; + uint32_t maxSemaphores_; + uint32_t maxSemaphoreSize_; //! Returns the SKU board name for the device char boardName_[128]; @@ -437,47 +437,47 @@ struct Info : public amd::EmbeddedObject { //! Number of SIMD (Single Instruction Multiple Data) units per compute unit //! that execute in parallel. All work items from the same work group must be //! executed by SIMDs in the same compute unit. - cl_uint simdPerCU_; - cl_uint cuPerShaderArray_; //!< Number of CUs per shader array + uint32_t simdPerCU_; + uint32_t cuPerShaderArray_; //!< Number of CUs per shader array //! The maximum number of work items from the same work group that can be //! executed by a SIMD in parallel - cl_uint simdWidth_; + uint32_t simdWidth_; //! The number of instructions that a SIMD can execute in parallel - cl_uint simdInstructionWidth_; + uint32_t simdInstructionWidth_; //! The number of workitems per wavefront - cl_uint wavefrontWidth_; + uint32_t wavefrontWidth_; //! Available number of SGPRs - cl_uint availableSGPRs_; + uint32_t availableSGPRs_; //! Number of global memory channels - cl_uint globalMemChannels_; + uint32_t globalMemChannels_; //! Number of banks in each global memory channel - cl_uint globalMemChannelBanks_; + uint32_t globalMemChannelBanks_; //! Width in bytes of each of global memory bank - cl_uint globalMemChannelBankWidth_; + uint32_t globalMemChannelBankWidth_; //! Local memory size per CU - cl_uint localMemSizePerCU_; + uint32_t localMemSizePerCU_; //! Number of banks of local memory - cl_uint localMemBanks_; + uint32_t localMemBanks_; //! The core engine GFXIP version - cl_uint gfxipVersion_; + uint32_t gfxipVersion_; //! Number of available async queues - cl_uint numAsyncQueues_; + uint32_t numAsyncQueues_; //! Number of available real time queues - cl_uint numRTQueues_; + uint32_t numRTQueues_; //! Number of available real time compute units - cl_uint numRTCUs_; + uint32_t numRTCUs_; //! Thread trace enable - cl_bool threadTraceEnable_; + bool threadTraceEnable_; //! ECC protected GPRs support (only available Vega20+) - cl_bool sramEccEnabled_; + bool sramEccEnabled_; //! Image pitch alignment for image2d_from_buffer - cl_uint imagePitchAlignment_; + uint32_t imagePitchAlignment_; //! Image base address alignment for image2d_from_buffer - cl_uint imageBaseAddressAlignment_; + uint32_t imageBaseAddressAlignment_; //! Describes whether buffers from images are supported - cl_bool bufferFromImageSupport_; + bool bufferFromImageSupport_; //! Returns the supported SPIR versions for the device const char* spirVersions_; @@ -485,22 +485,22 @@ struct Info : public amd::EmbeddedObject { //! OpenCL20 device info fields: //! The max number of pipe objects that can be passed as arguments to a kernel - cl_uint maxPipeArgs_; + uint32_t maxPipeArgs_; //! The max number of reservations that can be active for a pipe per work-item in a kernel - cl_uint maxPipeActiveReservations_; + uint32_t maxPipeActiveReservations_; //! The max size of pipe packet in bytes - cl_uint maxPipePacketSize_; + uint32_t maxPipePacketSize_; //! The command-queue properties supported of the device queue. cl_command_queue_properties queueOnDeviceProperties_; //! The preferred size of the device queue in bytes - cl_uint queueOnDevicePreferredSize_; + uint32_t queueOnDevicePreferredSize_; //! The max size of the device queue in bytes - cl_uint queueOnDeviceMaxSize_; + uint32_t queueOnDeviceMaxSize_; //! The maximum number of device queues - cl_uint maxOnDeviceQueues_; + uint32_t maxOnDeviceQueues_; //! The maximum number of events in use on a device queue - cl_uint maxOnDeviceEvents_; + uint32_t maxOnDeviceEvents_; //! The maximum size of global scope variables size_t maxGlobalVariableSize_; @@ -513,12 +513,12 @@ struct Info : public amd::EmbeddedObject { uint32_t pcieRevisionId_; //! Max numbers of threads per CU - cl_uint maxThreadsPerCU_; + uint32_t maxThreadsPerCU_; //! GPU device supports a launch of cooperative groups - cl_bool cooperativeGroups_; + bool cooperativeGroups_; //! GPU device supports a launch of cooperative groups on multiple devices - cl_bool cooperativeMultiDeviceGroups_; + bool cooperativeMultiDeviceGroups_; }; //! Device settings @@ -1233,9 +1233,9 @@ class Device : public RuntimeObject { ); static bool getDeviceIDs(cl_device_type deviceType, //!< Device type - cl_uint numEntries, //!< Number of entries in the array + uint32_t numEntries, //!< Number of entries in the array cl_device_id* devices, //!< Array of the device ID(s) - cl_uint* numDevices, //!< Number of available devices + uint32_t* numDevices, //!< Number of available devices bool offlineDevices //!< Report offline devices ); @@ -1360,7 +1360,7 @@ class Device : public RuntimeObject { HwDebugManager* hwDebugMgr() const { return hwDebugMgr_; } //! Initialize the Hardware Debug Manager - virtual cl_int hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage) { + virtual int32_t hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage) { return CL_SUCCESS; } diff --git a/rocclr/device/devkernel.hpp b/rocclr/device/devkernel.hpp index a0a798019a..b9a26290e0 100644 --- a/rocclr/device/devkernel.hpp +++ b/rocclr/device/devkernel.hpp @@ -364,9 +364,9 @@ class Kernel : public amd::HeapObject { struct WorkGroupInfo : public amd::EmbeddedObject { size_t size_; //!< kernel workgroup size size_t compileSize_[3]; //!< kernel compiled workgroup size - cl_ulong localMemSize_; //!< amount of used local memory + uint64_t localMemSize_; //!< amount of used local memory size_t preferredSizeMultiple_; //!< preferred multiple for launch - cl_ulong privateMemSize_; //!< amount of used private memory + uint64_t privateMemSize_; //!< amount of used private memory size_t scratchRegs_; //!< amount of used scratch registers size_t wavefrontPerSIMD_; //!< number of wavefronts per SIMD size_t wavefrontSize_; //!< number of threads per wavefront diff --git a/rocclr/device/devprogram.cpp b/rocclr/device/devprogram.cpp index 859fe0c87f..562dcf31f1 100644 --- a/rocclr/device/devprogram.cpp +++ b/rocclr/device/devprogram.cpp @@ -1452,7 +1452,7 @@ bool Program::finiBuild(bool isBuildGood) { } // ================================================================================================ -cl_int Program::compile(const std::string& sourceCode, +int32_t Program::compile(const std::string& sourceCode, const std::vector& headers, const char** headerIncludeNames, const char* origOptions, amd::option::Options* options) { @@ -1540,7 +1540,7 @@ cl_int Program::compile(const std::string& sourceCode, } // ================================================================================================ -cl_int Program::link(const std::vector& inputPrograms, const char* origLinkOptions, +int32_t Program::link(const std::vector& inputPrograms, const char* origLinkOptions, amd::option::Options* linkOptions) { lastBuildOptionsArg_ = origLinkOptions ? origLinkOptions : ""; if (linkOptions) { @@ -1643,7 +1643,7 @@ cl_int Program::link(const std::vector& inputPrograms, const char* ori } // ================================================================================================ -cl_int Program::build(const std::string& sourceCode, const char* origOptions, +int32_t Program::build(const std::string& sourceCode, const char* origOptions, amd::option::Options* options) { uint64_t start_time = 0; if (options->oVariables->EnableBuildTiming) { diff --git a/rocclr/device/devprogram.hpp b/rocclr/device/devprogram.hpp index 3a04ea2391..2f39234e8a 100644 --- a/rocclr/device/devprogram.hpp +++ b/rocclr/device/devprogram.hpp @@ -115,8 +115,8 @@ class Program : public amd::HeapObject { std::string lastBuildOptionsArg_; mutable std::string buildLog_; //!< build log. - cl_int buildStatus_; //!< build status. - cl_int buildError_; //!< build error + int32_t buildStatus_; //!< build status. + int32_t buildError_; //!< build error const char* machineTarget_; //!< Machine target for this program aclTargetInfo info_; //!< The info target for this binary. @@ -148,16 +148,16 @@ class Program : public amd::HeapObject { amd::option::Options* getCompilerOptions() const { return programOptions_; } //! Compile the device program. - cl_int compile(const std::string& sourceCode, const std::vector& headers, + int32_t compile(const std::string& sourceCode, const std::vector& headers, const char** headerIncludeNames, const char* origOptions, amd::option::Options* options); //! Builds the device program. - cl_int link(const std::vector& inputPrograms, const char* origOptions, + int32_t link(const std::vector& inputPrograms, const char* origOptions, amd::option::Options* options); //! Builds the device program. - cl_int build(const std::string& sourceCode, const char* origOptions, + int32_t build(const std::string& sourceCode, const char* origOptions, amd::option::Options* options); //! Returns the device object, associated with this program. @@ -177,7 +177,7 @@ class Program : public amd::HeapObject { cl_build_status buildStatus() const { return buildStatus_; } //! Return the build error. - cl_int buildError() const { return buildError_; } + int32_t buildError() const { return buildError_; } //! Return the symbols vector. const kernels_t& kernels() const { return kernels_; } diff --git a/rocclr/device/gpu/gpublit.cpp b/rocclr/device/gpu/gpublit.cpp index f2a90887a6..21b9494b02 100644 --- a/rocclr/device/gpu/gpublit.cpp +++ b/rocclr/device/gpu/gpublit.cpp @@ -761,8 +761,8 @@ bool KernelBlitManager::createProgram(Device& device) { // The following data structures will be used for the view creations. // Some formats has to be converted before a kernel blit operation struct FormatConvertion { - cl_uint clOldType_; - cl_uint clNewType_; + uint32_t clOldType_; + uint32_t clNewType_; }; // The list of rejected data formats and corresponding conversion @@ -949,7 +949,7 @@ bool KernelBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Mem return result; } -void CalcRowSlicePitches(cl_ulong* pitch, const cl_int* copySize, size_t rowPitch, +void CalcRowSlicePitches(uint64_t* pitch, const int32_t* copySize, size_t rowPitch, size_t slicePitch, const Memory& mem) { size_t memFmtSize = memoryFormatSize(mem.cal()->format_).size_; bool img1Darray = (mem.cal()->dimension_ == GSL_MOA_TEXTURE_1D_ARRAY) ? true : false; @@ -1115,18 +1115,18 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, const MemFormatStruct& memFmt = memoryFormatSize(gpuMem(dstMemory).cal()->format_); // 1 element granularity for writes by default - cl_int granularity = 1; + int32_t granularity = 1; if (memFmt.size_ == 2) { granularity = 2; } else if (memFmt.size_ >= 4) { granularity = 4; } CondLog(((srcOrigin[0] % granularity) != 0), "Unaligned offset in blit!"); - cl_ulong srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2], 0}; + uint64_t srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2], 0}; setArgument(kernels_[blitType], 2, sizeof(srcOrg), srcOrg); - cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1], (cl_int)dstOrigin[2], 0}; - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t dstOrg[4] = {(int32_t)dstOrigin[0], (int32_t)dstOrigin[1], (int32_t)dstOrigin[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 3, sizeof(dstOrg), dstOrg); setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); @@ -1134,11 +1134,11 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, // Program memory format uint multiplier = memFmt.size_ / sizeof(uint32_t); multiplier = (multiplier == 0) ? 1 : multiplier; - cl_uint format[4] = {memFmt.components_, memFmt.size_ / memFmt.components_, multiplier, 0}; + uint32_t format[4] = {memFmt.components_, memFmt.size_ / memFmt.components_, multiplier, 0}; setArgument(kernels_[blitType], 5, sizeof(format), format); // Program row and slice pitches - cl_ulong pitch[4] = {0}; + uint64_t pitch[4] = {0}; CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, gpuMem(dstMemory)); setArgument(kernels_[blitType], 6, sizeof(pitch), pitch); @@ -1424,31 +1424,31 @@ bool KernelBlitManager::copyImageToBufferKernel(device::Memory& srcMemory, setArgument(kernels_[blitType], 2, sizeof(cl_mem), &mem); setArgument(kernels_[blitType], 3, sizeof(cl_mem), &mem); - cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1], (cl_int)srcOrigin[2], 0}; - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t srcOrg[4] = {(int32_t)srcOrigin[0], (int32_t)srcOrigin[1], (int32_t)srcOrigin[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 4, sizeof(srcOrg), srcOrg); const MemFormatStruct& memFmt = memoryFormatSize(gpuMem(srcMemory).cal()->format_); // 1 element granularity for writes by default - cl_int granularity = 1; + int32_t granularity = 1; if (memFmt.size_ == 2) { granularity = 2; } else if (memFmt.size_ >= 4) { granularity = 4; } CondLog(((dstOrigin[0] % granularity) != 0), "Unaligned offset in blit!"); - cl_ulong dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2], 0}; + uint64_t dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2], 0}; setArgument(kernels_[blitType], 5, sizeof(dstOrg), dstOrg); setArgument(kernels_[blitType], 6, sizeof(copySize), copySize); // Program memory format uint multiplier = memFmt.size_ / sizeof(uint32_t); multiplier = (multiplier == 0) ? 1 : multiplier; - cl_uint format[4] = {memFmt.components_, memFmt.size_ / memFmt.components_, multiplier, 0}; + uint32_t format[4] = {memFmt.components_, memFmt.size_ / memFmt.components_, multiplier, 0}; setArgument(kernels_[blitType], 7, sizeof(format), format); // Program row and slice pitches - cl_ulong pitch[4] = {0}; + uint64_t pitch[4] = {0}; CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, gpuMem(srcMemory)); setArgument(kernels_[blitType], 8, sizeof(pitch), pitch); @@ -1565,14 +1565,14 @@ bool KernelBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dst setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); // Program source origin - cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1], (cl_int)srcOrigin[2], 0}; + int32_t srcOrg[4] = {(int32_t)srcOrigin[0], (int32_t)srcOrigin[1], (int32_t)srcOrigin[2], 0}; setArgument(kernels_[blitType], 2, sizeof(srcOrg), srcOrg); // Program destinaiton origin - cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1], (cl_int)dstOrigin[2], 0}; + int32_t dstOrg[4] = {(int32_t)dstOrigin[0], (int32_t)dstOrigin[1], (int32_t)dstOrigin[2], 0}; setArgument(kernels_[blitType], 3, sizeof(dstOrg), dstOrg); - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); // Create ND range object for the kernel's execution @@ -1806,11 +1806,11 @@ bool KernelBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory setArgument(kernels_[blitType], 0, sizeof(cl_mem), &mem); mem = &gpuMem(dstMemory); setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); - cl_ulong src[4] = {srcRect.rowPitch_, srcRect.slicePitch_, srcRect.start_, 0}; + uint64_t src[4] = {srcRect.rowPitch_, srcRect.slicePitch_, srcRect.start_, 0}; setArgument(kernels_[blitType], 2, sizeof(src), src); - cl_ulong dst[4] = {dstRect.rowPitch_, dstRect.slicePitch_, dstRect.start_, 0}; + uint64_t dst[4] = {dstRect.rowPitch_, dstRect.slicePitch_, dstRect.start_, 0}; setArgument(kernels_[blitType], 3, sizeof(dst), dst); - cl_ulong copySize[4] = {size[0], size[1], size[2], CopyRectAlignment[i]}; + uint64_t copySize[4] = {size[0], size[1], size[2], CopyRectAlignment[i]}; setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); // Create ND range object for the kernel's execution @@ -2036,7 +2036,7 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, } else { uint fillType = FillBuffer; size_t globalWorkOffset[3] = {0, 0, 0}; - cl_ulong fillSize = size[0] / patternSize; + uint64_t fillSize = size[0] / patternSize; size_t globalWorkSize = amd::alignUp(fillSize, 256); size_t localWorkSize = 256; bool dwordAligned = ((patternSize % sizeof(uint32_t)) == 0) ? true : false; @@ -2058,12 +2058,12 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, memcpy(constBuf, pattern, patternSize); gpuCB->unmap(&gpu()); setArgument(kernels_[fillType], 2, sizeof(cl_mem), &gpuCB); - cl_ulong offset = origin[0]; + uint64_t offset = origin[0]; if (dwordAligned) { patternSize /= sizeof(uint32_t); offset /= sizeof(uint32_t); } - setArgument(kernels_[fillType], 3, sizeof(cl_uint), &patternSize); + setArgument(kernels_[fillType], 3, sizeof(uint32_t), &patternSize); setArgument(kernels_[fillType], 4, sizeof(offset), &offset); setArgument(kernels_[fillType], 5, sizeof(fillSize), &fillSize); @@ -2114,7 +2114,7 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds } } - cl_uint remain; + uint32_t remain; if (blitType == BlitCopyBufferAligned) { size.c[0] /= CopyBuffAlignment[i]; } else { @@ -2146,20 +2146,20 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds mem = &gpuMem(dstMemory); setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); // Program source origin - cl_ulong srcOffset = srcOrigin[0] / CopyBuffAlignment[i]; + uint64_t srcOffset = srcOrigin[0] / CopyBuffAlignment[i]; ; setArgument(kernels_[blitType], 2, sizeof(srcOffset), &srcOffset); // Program destinaiton origin - cl_ulong dstOffset = dstOrigin[0] / CopyBuffAlignment[i]; + uint64_t dstOffset = dstOrigin[0] / CopyBuffAlignment[i]; ; setArgument(kernels_[blitType], 3, sizeof(dstOffset), &dstOffset); - cl_ulong copySize = size[0]; + uint64_t copySize = size[0]; setArgument(kernels_[blitType], 4, sizeof(copySize), ©Size); if (blitType == BlitCopyBufferAligned) { - cl_int alignment = CopyBuffAlignment[i]; + int32_t alignment = CopyBuffAlignment[i]; setArgument(kernels_[blitType], 5, sizeof(alignment), &alignment); } else { setArgument(kernels_[blitType], 5, sizeof(remain), &remain); @@ -2206,7 +2206,7 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, dim = 3; void* newpattern = const_cast(pattern); - cl_uint4 iFillColor; + uint32_t4 iFillColor; bool rejected = false; bool releaseView = false; @@ -2235,7 +2235,7 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, iFillColor.s[0] = sRGBmap(fColor[0]); iFillColor.s[1] = sRGBmap(fColor[1]); iFillColor.s[2] = sRGBmap(fColor[2]); - iFillColor.s[3] = (cl_uint)(fColor[3] * 255.0f); + iFillColor.s[3] = (uint32_t)(fColor[3] * 255.0f); newpattern = static_cast(&iFillColor); for (uint i = 0; i < RejectedFormatChannelTotal; ++i) { if (RejectedOrder[i].clOldType_ == newFormat.image_channel_order) { @@ -2281,12 +2281,12 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, // Program kernels arguments for the blit operation Memory* mem = memView; setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); - setArgument(kernels_[fillType], 1, sizeof(cl_float4), newpattern); - setArgument(kernels_[fillType], 2, sizeof(cl_int4), newpattern); - setArgument(kernels_[fillType], 3, sizeof(cl_uint4), newpattern); + setArgument(kernels_[fillType], 1, sizeof(float4), newpattern); + setArgument(kernels_[fillType], 2, sizeof(int32_t4), newpattern); + setArgument(kernels_[fillType], 3, sizeof(uint32_t4), newpattern); - cl_int fillOrigin[4] = {(cl_int)origin[0], (cl_int)origin[1], (cl_int)origin[2], 0}; - cl_int fillSize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t fillOrigin[4] = {(int32_t)origin[0], (int32_t)origin[1], (int32_t)origin[2], 0}; + int32_t fillSize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[fillType], 4, sizeof(fillOrigin), fillOrigin); setArgument(kernels_[fillType], 5, sizeof(fillSize), fillSize); diff --git a/rocclr/device/gpu/gpucounters.hpp b/rocclr/device/gpu/gpucounters.hpp index 35bff250e3..8cb50ceac0 100644 --- a/rocclr/device/gpu/gpucounters.hpp +++ b/rocclr/device/gpu/gpucounters.hpp @@ -81,9 +81,9 @@ class PerfCounter : public device::PerfCounter { //! Constructor for the GPU PerfCounter object PerfCounter(const Device& device, //!< A GPU device object const VirtualGPU& gpu, //!< Virtual GPU device object - cl_uint blockIndex, //!< HW block index - cl_uint counterIndex, //!< Counter index within the block - cl_uint eventIndex) //!< Event index for profiling + uint32_t blockIndex, //!< HW block index + uint32_t counterIndex, //!< Counter index within the block + uint32_t eventIndex) //!< Event index for profiling : gpuDevice_(device), gpu_(gpu), calRef_(NULL), diff --git a/rocclr/device/gpu/gpudebugmanager.cpp b/rocclr/device/gpu/gpudebugmanager.cpp index 72e031c3b2..1a84353bce 100644 --- a/rocclr/device/gpu/gpudebugmanager.cpp +++ b/rocclr/device/gpu/gpudebugmanager.cpp @@ -120,7 +120,7 @@ void GpuDebugManager::mapKernelCode(void* aqlCodeInfo) const { codeInfo->aqlCodeSize_ = aqlCodeSize_; } -cl_int GpuDebugManager::registerDebugger(amd::Context* context, uintptr_t messageStorage) { +int32_t GpuDebugManager::registerDebugger(amd::Context* context, uintptr_t messageStorage) { if (!device()->settings().enableHwDebug_) { LogError("debugmanager: Register debugger error - HW DEBUG is not enable"); return CL_DEBUGGER_REGISTER_FAILURE_AMD; @@ -228,7 +228,7 @@ DebugEvent GpuDebugManager::createDebugEvent(const bool autoReset) { return 0; } -cl_int GpuDebugManager::waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const { +int32_t GpuDebugManager::waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const { if (osEventTimedWait(pEvent, timeOut)) { return CL_SUCCESS; } else { @@ -291,7 +291,7 @@ void GpuDebugManager::setGlobalMemory(amd::Memory* memObj, uint32_t offset, void globalMem->unmap(NULL); } -cl_int GpuDebugManager::createRuntimeTrapHandler() { +int32_t GpuDebugManager::createRuntimeTrapHandler() { size_t codeSize = 0; const uint32_t* rtTrapCode = NULL; diff --git a/rocclr/device/gpu/gpudebugmanager.hpp b/rocclr/device/gpu/gpudebugmanager.hpp index 56ec313a35..61bdda86ab 100644 --- a/rocclr/device/gpu/gpudebugmanager.hpp +++ b/rocclr/device/gpu/gpudebugmanager.hpp @@ -66,13 +66,13 @@ class GpuDebugManager : public amd::HwDebugManager { DebugEvent createDebugEvent(const bool autoReset); //! Wait for the debug event - cl_int waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const; + int32_t waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const; //! Destroy the debug event void destroyDebugEvent(DebugEvent* pEvent); //! Register the debugger - cl_int registerDebugger(amd::Context* context, uintptr_t messageStorage); + int32_t registerDebugger(amd::Context* context, uintptr_t messageStorage); //! Unregister the debugger void unregisterDebugger(); @@ -105,7 +105,7 @@ class GpuDebugManager : public amd::HwDebugManager { void setupTrapInformation(DebugToolInfo* toolInfo); //! Create runtime trap handler - cl_int createRuntimeTrapHandler(); + int32_t createRuntimeTrapHandler(); protected: const VirtualGPU* vGpu() const { return vGpu_; } diff --git a/rocclr/device/gpu/gpudevice.cpp b/rocclr/device/gpu/gpudevice.cpp index a1509536a1..526ff1e041 100644 --- a/rocclr/device/gpu/gpudevice.cpp +++ b/rocclr/device/gpu/gpudevice.cpp @@ -351,7 +351,7 @@ void NullDevice::fillDeviceInfo(const CALdeviceattribs& calAttr, const gslMemInf info_.vramBusBitWidth_ = calAttr.memBusWidth; info_.l2CacheSize_ = 0; info_.maxParameterSize_ = 1024; - info_.minDataTypeAlignSize_ = sizeof(cl_long16); + info_.minDataTypeAlignSize_ = sizeof(int64_t16); info_.singleFPConfig_ = CL_FP_ROUND_TO_NEAREST | CL_FP_ROUND_TO_ZERO | CL_FP_ROUND_TO_INF | CL_FP_INF_NAN | CL_FP_FMA; @@ -377,20 +377,20 @@ void NullDevice::fillDeviceInfo(const CALdeviceattribs& calAttr, const gslMemInf #if defined(ATI_OS_LINUX) info_.globalMemSize_ = - (static_cast(std::min(GPU_MAX_HEAP_SIZE, 100u)) * + (static_cast(std::min(GPU_MAX_HEAP_SIZE, 100u)) * // globalMemSize is the actual available size for app on Linux // Because Linux base driver doesn't support paging - static_cast(memInfo.cardMemAvailableBytes + memInfo.cardExtMemAvailableBytes) / + static_cast(memInfo.cardMemAvailableBytes + memInfo.cardExtMemAvailableBytes) / 100u); #else - info_.globalMemSize_ = (static_cast(std::min(GPU_MAX_HEAP_SIZE, 100u)) * - static_cast(calAttr.localRAM) / 100u) * + info_.globalMemSize_ = (static_cast(std::min(GPU_MAX_HEAP_SIZE, 100u)) * + static_cast(calAttr.localRAM) / 100u) * Mi; #endif int uswcPercentAvailable = (calAttr.uncachedRemoteRAM > 1536 && IS_WINDOWS) ? 75 : 50; if (settings().apuSystem_) { info_.globalMemSize_ += - (static_cast(calAttr.uncachedRemoteRAM) * Mi * uswcPercentAvailable) / 100; + (static_cast(calAttr.uncachedRemoteRAM) * Mi * uswcPercentAvailable) / 100; } // We try to calculate the largest available memory size from @@ -400,32 +400,32 @@ void NullDevice::fillDeviceInfo(const CALdeviceattribs& calAttr, const gslMemInf // application progresses. #if defined(BRAHMA) && defined(ATI_BITS_64) info_.maxMemAllocSize_ = - std::max(cl_ulong(memInfo.cardMemAvailableBytes), cl_ulong(memInfo.cardExtMemAvailableBytes)); + std::max(uint64_t(memInfo.cardMemAvailableBytes), uint64_t(memInfo.cardExtMemAvailableBytes)); #else - info_.maxMemAllocSize_ = std::max(cl_ulong(memInfo.cardLargestFreeBlockBytes), - cl_ulong(memInfo.cardExtLargestFreeBlockBytes)); + info_.maxMemAllocSize_ = std::max(uint64_t(memInfo.cardLargestFreeBlockBytes), + uint64_t(memInfo.cardExtLargestFreeBlockBytes)); #endif if (settings().apuSystem_) { info_.maxMemAllocSize_ = std::max( - (static_cast(calAttr.uncachedRemoteRAM) * Mi * uswcPercentAvailable) / 100, + (static_cast(calAttr.uncachedRemoteRAM) * Mi * uswcPercentAvailable) / 100, info_.maxMemAllocSize_); } info_.maxMemAllocSize_ = - cl_ulong(info_.maxMemAllocSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); + uint64_t(info_.maxMemAllocSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); //! \note Force max single allocation size. //! 4GB limit for the blit kernels and 64 bit optimizations. info_.maxMemAllocSize_ = - std::min(info_.maxMemAllocSize_, static_cast(settings().maxAllocSize_)); + std::min(info_.maxMemAllocSize_, static_cast(settings().maxAllocSize_)); - if (info_.maxMemAllocSize_ < cl_ulong(128 * Mi)) { + if (info_.maxMemAllocSize_ < uint64_t(128 * Mi)) { LogError( "We are unable to get a heap large enough to support the OpenCL minimum " "requirement for FULL_PROFILE"); } - info_.maxMemAllocSize_ = std::max(cl_ulong(128 * Mi), info_.maxMemAllocSize_); + info_.maxMemAllocSize_ = std::max(uint64_t(128 * Mi), info_.maxMemAllocSize_); // Clamp max single alloc size to the globalMemSize since it's // reduced by default @@ -441,7 +441,7 @@ void NullDevice::fillDeviceInfo(const CALdeviceattribs& calAttr, const gslMemInf } else { info_.addressBits_ = 32; // Limit total size with 3GB for 32 bit - info_.globalMemSize_ = std::min(info_.globalMemSize_, cl_ulong(3 * Gi)); + info_.globalMemSize_ = std::min(info_.globalMemSize_, uint64_t(3 * Gi)); } // Alignment in BITS of the base address of any allocated memory object @@ -2237,8 +2237,8 @@ void Device::SrdManager::fillResourceList(std::vector& memList) { } } -cl_int Device::hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage) { - cl_int status = hwDebugMgr_->registerDebugger(context, messageStorage); +int32_t Device::hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage) { + int32_t status = hwDebugMgr_->registerDebugger(context, messageStorage); if (CL_SUCCESS != status) { delete hwDebugMgr_; diff --git a/rocclr/device/gpu/gpudevice.hpp b/rocclr/device/gpu/gpudevice.hpp index fb1a943e14..6b36b2ab23 100644 --- a/rocclr/device/gpu/gpudevice.hpp +++ b/rocclr/device/gpu/gpudevice.hpp @@ -538,7 +538,7 @@ class Device : public NullDevice, public CALGSLDevice { SrdManager& srds() const { return *srdManager_; } //! Initial the Hardware Debug Manager - cl_int hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage); + int32_t hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage); private: //! Disable copy constructor diff --git a/rocclr/device/gpu/gpukernel.cpp b/rocclr/device/gpu/gpukernel.cpp index 5f40217478..35a7c6d66a 100644 --- a/rocclr/device/gpu/gpukernel.cpp +++ b/rocclr/device/gpu/gpukernel.cpp @@ -399,9 +399,9 @@ size_t KernelArg::size(bool gpuLayer) const { case PointerHwPrivate: return (gpuLayer) ? sizeof(uint32_t) * size_ : 0; case Float: - return sizeof(cl_float) * amd::nextPowerOfTwo(size_); + return sizeof(float) * amd::nextPowerOfTwo(size_); case Double: - return sizeof(cl_double) * amd::nextPowerOfTwo(size_); + return sizeof(double) * amd::nextPowerOfTwo(size_); case Char: case UChar: return sizeof(cl_char) * amd::nextPowerOfTwo(size_); @@ -410,10 +410,10 @@ size_t KernelArg::size(bool gpuLayer) const { return sizeof(cl_short) * amd::nextPowerOfTwo(size_); case Int: case UInt: - return sizeof(cl_uint) * amd::nextPowerOfTwo(size_); + return sizeof(uint32_t) * amd::nextPowerOfTwo(size_); case Long: case ULong: - return sizeof(cl_ulong) * amd::nextPowerOfTwo(size_); + return sizeof(uint64_t) * amd::nextPowerOfTwo(size_); case Struct: case Union: return (gpuLayer) ? amd::alignUp(size_, 16) : size_; diff --git a/rocclr/device/gpu/gpukernel.hpp b/rocclr/device/gpu/gpukernel.hpp index 4a2d01de62..22777f93ad 100644 --- a/rocclr/device/gpu/gpukernel.hpp +++ b/rocclr/device/gpu/gpukernel.hpp @@ -425,7 +425,7 @@ class NullKernel : public device::Kernel { const NullProgram& nullProg() const { return reinterpret_cast(prog_); } //! Returns the kernel's build error - const cl_int buildError() const { return buildError_; } + const int32_t buildError() const { return buildError_; } //! Returns the kernel's flags uint flags() const { return flags_; } @@ -467,7 +467,7 @@ class NullKernel : public device::Kernel { //! Returns UAV raw index for this kernel uint uavRaw() const { return uavRaw_; } - cl_int buildError_; //!< Kernel's build error + int32_t buildError_; //!< Kernel's build error std::string ilSource_; //!< IL source code of this kernel const NullDevice& gpuDev_; //!< GPU device object diff --git a/rocclr/device/gpu/gpuprogram.cpp b/rocclr/device/gpu/gpuprogram.cpp index bcc32fd89f..f4fcfc8fc3 100644 --- a/rocclr/device/gpu/gpuprogram.cpp +++ b/rocclr/device/gpu/gpuprogram.cpp @@ -223,7 +223,7 @@ bool NullProgram::linkImpl(amd::option::Options* options) { if (!llvmBinary_.empty()) { // Compile llvm binary to the IL source code // This is link/OPT/Codegen part of compiler. - cl_int iErr = compileBinaryToIL(options); + int32_t iErr = compileBinaryToIL(options); if (iErr != CL_SUCCESS) { buildLog_ += "Error: Compilation from LLVMIR binary to IL text failed!"; LogError(buildLog_.c_str()); @@ -588,7 +588,7 @@ bool NullProgram::linkImpl(const std::vector& inputPrograms, // Compile llvm binary to the IL source code // This is link/OPT/Codegen part of compiler. - cl_int iErr = compileBinaryToIL(options); + int32_t iErr = compileBinaryToIL(options); if (iErr != CL_SUCCESS) { buildLog_ += "Error: Compilation from LLVMIR binary to IL text failed!"; LogError(buildLog_.c_str()); diff --git a/rocclr/device/gpu/gpuvirtual.cpp b/rocclr/device/gpu/gpuvirtual.cpp index aeced97474..a7d4fa5bfc 100644 --- a/rocclr/device/gpu/gpuvirtual.cpp +++ b/rocclr/device/gpu/gpuvirtual.cpp @@ -1424,7 +1424,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& vcmd) { std::vector& svmPointers = vcmd.svmPointers(); if (vcmd.pfnFreeFunc() == NULL) { // pointers allocated using clSVMAlloc - for (cl_uint i = 0; i < svmPointers.size(); ++i) { + for (uint32_t i = 0; i < svmPointers.size(); ++i) { dev().svmFree(svmPointers[i]); } } else { @@ -2418,10 +2418,10 @@ void VirtualGPU::submitMakeBuffersResident(amd::MakeBuffersResidentCommand& vcmd amd::ScopedLock lock(execution()); profilingBegin(vcmd); std::vector memObjects = vcmd.memObjects(); - cl_uint numObjects = memObjects.size(); + uint32_t numObjects = memObjects.size(); gslMemObject* pGSLMemObjects = new gslMemObject[numObjects]; - for (cl_uint i = 0; i < numObjects; ++i) { + for (uint32_t i = 0; i < numObjects; ++i) { gpu::Memory* gpuMemory = dev().getGpuMemory(memObjects[i]); pGSLMemObjects[i] = gpuMemory->gslResource(); gpuMemory->syncCacheFromHost(*this); @@ -2436,7 +2436,7 @@ void VirtualGPU::submitMakeBuffersResident(amd::MakeBuffersResidentCommand& vcmd vcmd.setStatus(CL_INVALID_OPERATION); } else { cl_bus_address_amd* busAddr = vcmd.busAddress(); - for (cl_uint i = 0; i < numObjects; ++i) { + for (uint32_t i = 0; i < numObjects; ++i) { busAddr[i].surface_bus_address = surfBusAddr[i]; busAddr[i].marker_bus_address = markerBusAddr[i]; } diff --git a/rocclr/device/hwdebug.hpp b/rocclr/device/hwdebug.hpp index 7f0ab9d069..9f90227d24 100644 --- a/rocclr/device/hwdebug.hpp +++ b/rocclr/device/hwdebug.hpp @@ -180,13 +180,13 @@ class HwDebugManager { virtual DebugEvent createDebugEvent(const bool autoReset) = 0; //! Wait for the debug event - virtual cl_int waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const = 0; + virtual int32_t waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const = 0; //! Destroy the debug event virtual void destroyDebugEvent(DebugEvent* pEvent) = 0; //! Register the debugger - virtual cl_int registerDebugger(amd::Context* context, uintptr_t pMessageStorage) = 0; + virtual int32_t registerDebugger(amd::Context* context, uintptr_t pMessageStorage) = 0; //! Unregister the debugger virtual void unregisterDebugger() = 0; diff --git a/rocclr/device/pal/palblit.cpp b/rocclr/device/pal/palblit.cpp index 63b57f2278..fca550a227 100644 --- a/rocclr/device/pal/palblit.cpp +++ b/rocclr/device/pal/palblit.cpp @@ -748,8 +748,8 @@ bool KernelBlitManager::createProgram(Device& device) { // The following data structures will be used for the view creations. // Some formats has to be converted before a kernel blit operation struct FormatConvertion { - cl_uint clOldType_; - cl_uint clNewType_; + uint32_t clOldType_; + uint32_t clNewType_; }; // The list of rejected data formats and corresponding conversion @@ -936,7 +936,7 @@ bool KernelBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Mem return result; } -void CalcRowSlicePitches(cl_ulong* pitch, const cl_int* copySize, size_t rowPitch, +void CalcRowSlicePitches(uint64_t* pitch, const int32_t* copySize, size_t rowPitch, size_t slicePitch, const Memory& mem) { uint32_t memFmtSize = mem.elementSize(); bool img1Darray = (mem.desc().topology_ == CL_MEM_OBJECT_IMAGE1D_ARRAY) ? true : false; @@ -1118,18 +1118,18 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, uint32_t components = gpuMem(dstMemory).numComponents(); // 1 element granularity for writes by default - cl_int granularity = 1; + int32_t granularity = 1; if (memFmtSize == 2) { granularity = 2; } else if (memFmtSize >= 4) { granularity = 4; } CondLog(((srcOrigin[0] % granularity) != 0), "Unaligned offset in blit!"); - cl_ulong srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2], 0}; + uint64_t srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2], 0}; setArgument(kernels_[blitType], 2, sizeof(srcOrg), srcOrg); - cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1], (cl_int)dstOrigin[2], 0}; - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t dstOrg[4] = {(int32_t)dstOrigin[0], (int32_t)dstOrigin[1], (int32_t)dstOrigin[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; if (swapLayer) { dstOrg[2] = dstOrg[1]; @@ -1144,11 +1144,11 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, // Program memory format uint multiplier = memFmtSize / sizeof(uint32_t); multiplier = (multiplier == 0) ? 1 : multiplier; - cl_uint format[4] = {components, memFmtSize / components, multiplier, 0}; + uint32_t format[4] = {components, memFmtSize / components, multiplier, 0}; setArgument(kernels_[blitType], 5, sizeof(format), format); // Program row and slice pitches - cl_ulong pitch[4] = {0}; + uint64_t pitch[4] = {0}; CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, gpuMem(dstMemory)); setArgument(kernels_[blitType], 6, sizeof(pitch), pitch); @@ -1441,8 +1441,8 @@ bool KernelBlitManager::copyImageToBufferKernel(device::Memory& srcMemory, setArgument(kernels_[blitType], 2, sizeof(cl_mem), &mem); setArgument(kernels_[blitType], 3, sizeof(cl_mem), &mem); - cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1], (cl_int)srcOrigin[2], 0}; - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t srcOrg[4] = {(int32_t)srcOrigin[0], (int32_t)srcOrigin[1], (int32_t)srcOrigin[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; if (swapLayer) { srcOrg[2] = srcOrg[1]; srcOrg[1] = 0; @@ -1454,25 +1454,25 @@ bool KernelBlitManager::copyImageToBufferKernel(device::Memory& srcMemory, uint32_t components = gpuMem(srcMemory).numComponents(); // 1 element granularity for writes by default - cl_int granularity = 1; + int32_t granularity = 1; if (memFmtSize == 2) { granularity = 2; } else if (memFmtSize >= 4) { granularity = 4; } CondLog(((dstOrigin[0] % granularity) != 0), "Unaligned offset in blit!"); - cl_ulong dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2], 0}; + uint64_t dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2], 0}; setArgument(kernels_[blitType], 5, sizeof(dstOrg), dstOrg); setArgument(kernels_[blitType], 6, sizeof(copySize), copySize); // Program memory format uint multiplier = memFmtSize / sizeof(uint32_t); multiplier = (multiplier == 0) ? 1 : multiplier; - cl_uint format[4] = {components, memFmtSize / components, multiplier, 0}; + uint32_t format[4] = {components, memFmtSize / components, multiplier, 0}; setArgument(kernels_[blitType], 7, sizeof(format), format); // Program row and slice pitches - cl_ulong pitch[4] = {0}; + uint64_t pitch[4] = {0}; CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, gpuMem(srcMemory)); setArgument(kernels_[blitType], 8, sizeof(pitch), pitch); @@ -1586,7 +1586,7 @@ bool KernelBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dst setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); // Program source origin - cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1], (cl_int)srcOrigin[2], 0}; + int32_t srcOrg[4] = {(int32_t)srcOrigin[0], (int32_t)srcOrigin[1], (int32_t)srcOrigin[2], 0}; if ((gpuMem(srcMemory).desc().topology_ == CL_MEM_OBJECT_IMAGE1D_ARRAY) && dev().settings().gfx10Plus_) { srcOrg[3] = 1; @@ -1594,14 +1594,14 @@ bool KernelBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dst setArgument(kernels_[blitType], 2, sizeof(srcOrg), srcOrg); // Program destinaiton origin - cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1], (cl_int)dstOrigin[2], 0}; + int32_t dstOrg[4] = {(int32_t)dstOrigin[0], (int32_t)dstOrigin[1], (int32_t)dstOrigin[2], 0}; if ((gpuMem(dstMemory).desc().topology_ == CL_MEM_OBJECT_IMAGE1D_ARRAY) && dev().settings().gfx10Plus_) { dstOrg[3] = 1; } setArgument(kernels_[blitType], 3, sizeof(dstOrg), dstOrg); - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); // Create ND range object for the kernel's execution @@ -1848,11 +1848,11 @@ bool KernelBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory setArgument(kernels_[blitType], 0, sizeof(cl_mem), &mem); mem = &gpuMem(dstMemory); setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); - cl_ulong src[4] = {srcRect.rowPitch_, srcRect.slicePitch_, srcRect.start_, 0}; + uint64_t src[4] = {srcRect.rowPitch_, srcRect.slicePitch_, srcRect.start_, 0}; setArgument(kernels_[blitType], 2, sizeof(src), src); - cl_ulong dst[4] = {dstRect.rowPitch_, dstRect.slicePitch_, dstRect.start_, 0}; + uint64_t dst[4] = {dstRect.rowPitch_, dstRect.slicePitch_, dstRect.start_, 0}; setArgument(kernels_[blitType], 3, sizeof(dst), dst); - cl_ulong copySize[4] = {size[0], size[1], size[2], CopyRectAlignment[i]}; + uint64_t copySize[4] = {size[0], size[1], size[2], CopyRectAlignment[i]}; setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); // Create ND range object for the kernel's execution @@ -2077,7 +2077,7 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, } else { uint fillType = FillBuffer; size_t globalWorkOffset[3] = {0, 0, 0}; - cl_ulong fillSize = size[0] / patternSize; + uint64_t fillSize = size[0] / patternSize; size_t globalWorkSize = amd::alignUp(fillSize, 256); size_t localWorkSize = 256; bool dwordAligned = ((patternSize % sizeof(uint32_t)) == 0) ? true : false; @@ -2097,12 +2097,12 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, gpuCB.unmap(&gpu()); Memory* pGpuCB = &gpuCB; setArgument(kernels_[fillType], 2, sizeof(cl_mem), &pGpuCB); - cl_ulong offset = origin[0]; + uint64_t offset = origin[0]; if (dwordAligned) { patternSize /= sizeof(uint32_t); offset /= sizeof(uint32_t); } - setArgument(kernels_[fillType], 3, sizeof(cl_uint), &patternSize); + setArgument(kernels_[fillType], 3, sizeof(uint32_t), &patternSize); setArgument(kernels_[fillType], 4, sizeof(offset), &offset); setArgument(kernels_[fillType], 5, sizeof(fillSize), &fillSize); @@ -2153,7 +2153,7 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds } } - cl_uint remain; + uint32_t remain; if (blitType == BlitCopyBufferAligned) { size.c[0] /= CopyBuffAlignment[i]; } else { @@ -2172,18 +2172,18 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds mem = &gpuMem(dstMemory); setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); // Program source origin - cl_ulong srcOffset = srcOrigin[0] / CopyBuffAlignment[i]; + uint64_t srcOffset = srcOrigin[0] / CopyBuffAlignment[i]; setArgument(kernels_[blitType], 2, sizeof(srcOffset), &srcOffset); // Program destinaiton origin - cl_ulong dstOffset = dstOrigin[0] / CopyBuffAlignment[i]; + uint64_t dstOffset = dstOrigin[0] / CopyBuffAlignment[i]; setArgument(kernels_[blitType], 3, sizeof(dstOffset), &dstOffset); - cl_ulong copySize = size[0]; + uint64_t copySize = size[0]; setArgument(kernels_[blitType], 4, sizeof(copySize), ©Size); if (blitType == BlitCopyBufferAligned) { - cl_int alignment = CopyBuffAlignment[i]; + int32_t alignment = CopyBuffAlignment[i]; setArgument(kernels_[blitType], 5, sizeof(alignment), &alignment); } else { setArgument(kernels_[blitType], 5, sizeof(remain), &remain); @@ -2232,7 +2232,7 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, dim = 3; void* newpattern = const_cast(pattern); - cl_uint4 iFillColor; + uint32_t4 iFillColor; bool rejected = false; bool releaseView = false; @@ -2254,7 +2254,7 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, iFillColor.s[0] = sRGBmap(fColor[0]); iFillColor.s[1] = sRGBmap(fColor[1]); iFillColor.s[2] = sRGBmap(fColor[2]); - iFillColor.s[3] = (cl_uint)(fColor[3] * 255.0f); + iFillColor.s[3] = (uint32_t)(fColor[3] * 255.0f); newpattern = static_cast(&iFillColor); for (uint i = 0; i < RejectedFormatChannelTotal; ++i) { if (RejectedOrder[i].clOldType_ == newFormat.image_channel_order) { @@ -2308,12 +2308,12 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, // Program kernels arguments for the blit operation Memory* mem = memView; setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); - setArgument(kernels_[fillType], 1, sizeof(cl_float4), newpattern); - setArgument(kernels_[fillType], 2, sizeof(cl_int4), newpattern); - setArgument(kernels_[fillType], 3, sizeof(cl_uint4), newpattern); + setArgument(kernels_[fillType], 1, sizeof(float4), newpattern); + setArgument(kernels_[fillType], 2, sizeof(int32_t4), newpattern); + setArgument(kernels_[fillType], 3, sizeof(uint32_t4), newpattern); - cl_int fillOrigin[4] = {(cl_int)origin[0], (cl_int)origin[1], (cl_int)origin[2], 0}; - cl_int fillSize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t fillOrigin[4] = {(int32_t)origin[0], (int32_t)origin[1], (int32_t)origin[2], 0}; + int32_t fillSize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; if (swapLayer) { fillOrigin[2] = fillOrigin[1]; fillOrigin[1] = 0; diff --git a/rocclr/device/pal/palconstbuf.cpp b/rocclr/device/pal/palconstbuf.cpp index fe43142160..5243e3ff68 100644 --- a/rocclr/device/pal/palconstbuf.cpp +++ b/rocclr/device/pal/palconstbuf.cpp @@ -72,7 +72,7 @@ bool ManagedBuffer::create(Resource::MemoryType type) { // ================================================================================================ address ManagedBuffer::reserve(uint32_t size, uint64_t* gpu_address) { // Align to the maximum data size available in OpenCL - static constexpr uint32_t MemAlignment = sizeof(cl_double16); + static constexpr uint32_t MemAlignment = sizeof(double16); // Align reserve size on the vector's boundary uint32_t count = amd::alignUp(size, MemAlignment); diff --git a/rocclr/device/pal/palcounters.hpp b/rocclr/device/pal/palcounters.hpp index 332e4d4526..d3adc3b6c9 100644 --- a/rocclr/device/pal/palcounters.hpp +++ b/rocclr/device/pal/palcounters.hpp @@ -99,9 +99,9 @@ class PerfCounter : public device::PerfCounter { //! Constructor for the GPU PerfCounter object PerfCounter(const Device& device, //!< A GPU device object PalCounterReference* palRef, //!< Counter Reference - cl_uint blockIndex, //!< HW block index - cl_uint counterIndex, //!< Counter index within the block - cl_uint eventIndex) //!< Event index for profiling + uint32_t blockIndex, //!< HW block index + uint32_t counterIndex, //!< Counter index within the block + uint32_t eventIndex) //!< Event index for profiling : gpuDevice_(device), palRef_(palRef) { info_.blockIndex_ = blockIndex; info_.counterIndex_ = counterIndex; diff --git a/rocclr/device/pal/paldebugmanager.cpp b/rocclr/device/pal/paldebugmanager.cpp index 38986ff2a5..11419842ae 100644 --- a/rocclr/device/pal/paldebugmanager.cpp +++ b/rocclr/device/pal/paldebugmanager.cpp @@ -121,7 +121,7 @@ void GpuDebugManager::mapKernelCode(void* aqlCodeInfo) const { codeInfo->aqlCodeSize_ = aqlCodeSize_; } -cl_int GpuDebugManager::registerDebugger(amd::Context* context, uintptr_t messageStorage) { +int32_t GpuDebugManager::registerDebugger(amd::Context* context, uintptr_t messageStorage) { if (!device()->settings().enableHwDebug_) { LogError("debugmanager: Register debugger error - HW DEBUG is not enable"); return CL_DEBUGGER_REGISTER_FAILURE_AMD; @@ -234,7 +234,7 @@ DebugEvent GpuDebugManager::createDebugEvent(const bool autoReset) { return 0; } -cl_int GpuDebugManager::waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const { +int32_t GpuDebugManager::waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const { Unimplemented(); /* if (osEventTimedWait(pEvent, timeOut)) { @@ -307,7 +307,7 @@ void GpuDebugManager::setGlobalMemory(amd::Memory* memObj, uint32_t offset, void globalMem->unmap(nullptr); } -cl_int GpuDebugManager::createRuntimeTrapHandler() { +int32_t GpuDebugManager::createRuntimeTrapHandler() { size_t codeSize = 0; const uint32_t* rtTrapCode = nullptr; diff --git a/rocclr/device/pal/paldebugmanager.hpp b/rocclr/device/pal/paldebugmanager.hpp index a766652a82..6c800997a8 100644 --- a/rocclr/device/pal/paldebugmanager.hpp +++ b/rocclr/device/pal/paldebugmanager.hpp @@ -65,13 +65,13 @@ class GpuDebugManager : public amd::HwDebugManager { DebugEvent createDebugEvent(const bool autoReset); //! Wait for the debug event - cl_int waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const; + int32_t waitDebugEvent(DebugEvent pEvent, uint32_t timeOut) const; //! Destroy the debug event void destroyDebugEvent(DebugEvent* pEvent); //! Register the debugger - cl_int registerDebugger(amd::Context* context, uintptr_t messageStorage); + int32_t registerDebugger(amd::Context* context, uintptr_t messageStorage); //! Unregister the debugger void unregisterDebugger(); @@ -107,7 +107,7 @@ class GpuDebugManager : public amd::HwDebugManager { void setupTrapInformation(DebugToolInfo* toolInfo); //! Create runtime trap handler - cl_int createRuntimeTrapHandler(); + int32_t createRuntimeTrapHandler(); const pal::Device* device() const { return reinterpret_cast(device_); } diff --git a/rocclr/device/pal/paldevice.cpp b/rocclr/device/pal/paldevice.cpp index d6e71d2d04..570f127906 100644 --- a/rocclr/device/pal/paldevice.cpp +++ b/rocclr/device/pal/paldevice.cpp @@ -400,7 +400,7 @@ void NullDevice::fillDeviceInfo(const Pal::DeviceProperties& palProp, info_.vramBusBitWidth_ = palProp.gpuMemoryProperties.performance.vramBusBitWidth; info_.l2CacheSize_ = palProp.gfxipProperties.shaderCore.tccSizeInBytes; info_.maxParameterSize_ = 1024; - info_.minDataTypeAlignSize_ = sizeof(cl_long16); + info_.minDataTypeAlignSize_ = sizeof(int64_t16); info_.singleFPConfig_ = CL_FP_ROUND_TO_NEAREST | CL_FP_ROUND_TO_ZERO | CL_FP_ROUND_TO_INF | CL_FP_INF_NAN | CL_FP_FMA; @@ -432,49 +432,49 @@ void NullDevice::fillDeviceInfo(const Pal::DeviceProperties& palProp, heaps[Pal::GpuHeapLocal].physicalHeapSize + heaps[Pal::GpuHeapInvisible].physicalHeapSize; } - info_.globalMemSize_ = (static_cast(std::min(GPU_MAX_HEAP_SIZE, 100u)) * - static_cast(localRAM) / 100u); + info_.globalMemSize_ = (static_cast(std::min(GPU_MAX_HEAP_SIZE, 100u)) * + static_cast(localRAM) / 100u); uint uswcPercentAvailable = - ((static_cast(heaps[Pal::GpuHeapGartUswc].heapSize) / Mi) > 1536 && IS_WINDOWS) + ((static_cast(heaps[Pal::GpuHeapGartUswc].heapSize) / Mi) > 1536 && IS_WINDOWS) ? 75 : 50; if (settings().apuSystem_) { info_.globalMemSize_ += - (static_cast(heaps[Pal::GpuHeapGartUswc].heapSize) * uswcPercentAvailable) / 100; + (static_cast(heaps[Pal::GpuHeapGartUswc].heapSize) * uswcPercentAvailable) / 100; } // Find the largest heap form FB memory if (GPU_ADD_HBCC_SIZE) { - info_.maxMemAllocSize_ = std::max(cl_ulong(heaps[Pal::GpuHeapLocal].heapSize), - cl_ulong(heaps[Pal::GpuHeapInvisible].heapSize)); + info_.maxMemAllocSize_ = std::max(uint64_t(heaps[Pal::GpuHeapLocal].heapSize), + uint64_t(heaps[Pal::GpuHeapInvisible].heapSize)); } else { - info_.maxMemAllocSize_ = std::max(cl_ulong(heaps[Pal::GpuHeapLocal].physicalHeapSize), - cl_ulong(heaps[Pal::GpuHeapInvisible].physicalHeapSize)); + info_.maxMemAllocSize_ = std::max(uint64_t(heaps[Pal::GpuHeapLocal].physicalHeapSize), + uint64_t(heaps[Pal::GpuHeapInvisible].physicalHeapSize)); } #if defined(ATI_OS_WIN) if (settings().apuSystem_) { info_.maxMemAllocSize_ = std::max( - (static_cast(heaps[Pal::GpuHeapGartUswc].heapSize) * uswcPercentAvailable) / 100, + (static_cast(heaps[Pal::GpuHeapGartUswc].heapSize) * uswcPercentAvailable) / 100, info_.maxMemAllocSize_); } #endif info_.maxMemAllocSize_ = - cl_ulong(info_.maxMemAllocSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); + uint64_t(info_.maxMemAllocSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); //! \note Force max single allocation size. //! 4GB limit for the blit kernels and 64 bit optimizations. info_.maxMemAllocSize_ = - std::min(info_.maxMemAllocSize_, static_cast(settings().maxAllocSize_)); + std::min(info_.maxMemAllocSize_, static_cast(settings().maxAllocSize_)); - if (info_.maxMemAllocSize_ < cl_ulong(128 * Mi)) { + if (info_.maxMemAllocSize_ < uint64_t(128 * Mi)) { LogError( "We are unable to get a heap large enough to support the OpenCL minimum " "requirement for FULL_PROFILE"); } - info_.maxMemAllocSize_ = std::max(cl_ulong(128 * Mi), info_.maxMemAllocSize_); + info_.maxMemAllocSize_ = std::max(uint64_t(128 * Mi), info_.maxMemAllocSize_); // Clamp max single alloc size to the globalMemSize since it's // reduced by default @@ -490,7 +490,7 @@ void NullDevice::fillDeviceInfo(const Pal::DeviceProperties& palProp, } else { info_.addressBits_ = (settings().useLightning_) ? 64 : 32; // Limit total size with 3GB for 32 bit - info_.globalMemSize_ = std::min(info_.globalMemSize_, cl_ulong(3 * Gi)); + info_.globalMemSize_ = std::min(info_.globalMemSize_, uint64_t(3 * Gi)); } // Alignment in BITS of the base address of any allocated memory object @@ -2409,8 +2409,8 @@ void Device::SrdManager::fillResourceList(VirtualGPU& gpu) { } } -cl_int Device::hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage) { - cl_int status = hwDebugMgr_->registerDebugger(context, messageStorage); +int32_t Device::hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage) { + int32_t status = hwDebugMgr_->registerDebugger(context, messageStorage); if (CL_SUCCESS != status) { delete hwDebugMgr_; diff --git a/rocclr/device/pal/paldevice.hpp b/rocclr/device/pal/paldevice.hpp index 20e24ce946..d1c21353b9 100644 --- a/rocclr/device/pal/paldevice.hpp +++ b/rocclr/device/pal/paldevice.hpp @@ -481,7 +481,7 @@ class Device : public NullDevice { SrdManager& srds() const { return *srdManager_; } //! Initial the Hardware Debug Manager - cl_int hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage); + int32_t hwDebugManagerInit(amd::Context* context, uintptr_t messageStorage); //! Returns PAL device properties const Pal::DeviceProperties& properties() const { return properties_; } diff --git a/rocclr/device/pal/palvirtual.cpp b/rocclr/device/pal/palvirtual.cpp index a09556cf4f..53dc7e52d7 100644 --- a/rocclr/device/pal/palvirtual.cpp +++ b/rocclr/device/pal/palvirtual.cpp @@ -2075,7 +2075,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& vcmd) { std::vector& svmPointers = vcmd.svmPointers(); if (vcmd.pfnFreeFunc() == nullptr) { // pointers allocated using clSVMAlloc - for (cl_uint i = 0; i < svmPointers.size(); ++i) { + for (uint32_t i = 0; i < svmPointers.size(); ++i) { dev().svmFree(svmPointers[i]); } } else { diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index 548c2a0897..3a7b36e0ac 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -808,8 +808,8 @@ bool KernelBlitManager::createProgram(Device& device) { // The following data structures will be used for the view creations. // Some formats has to be converted before a kernel blit operation struct FormatConvertion { - cl_uint clOldType_; - cl_uint clNewType_; + uint32_t clOldType_; + uint32_t clNewType_; }; // The list of rejected data formats and corresponding conversion @@ -874,7 +874,7 @@ bool KernelBlitManager::copyBufferToImage(device::Memory& srcMemory, device::Mem return result; } -void CalcRowSlicePitches(cl_ulong* pitch, const cl_int* copySize, size_t rowPitch, +void CalcRowSlicePitches(uint64_t* pitch, const int32_t* copySize, size_t rowPitch, size_t slicePitch, const Memory& mem) { amd::Image* image = static_cast(mem.owner()); uint32_t memFmtSize = image->getImageFormat().getElementSize(); @@ -985,18 +985,18 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, uint32_t components = dstImage->getImageFormat().getNumChannels(); // 1 element granularity for writes by default - cl_int granularity = 1; + int32_t granularity = 1; if (memFmtSize == 2) { granularity = 2; } else if (memFmtSize >= 4) { granularity = 4; } CondLog(((srcOrigin[0] % granularity) != 0), "Unaligned offset in blit!"); - cl_ulong srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2], 0}; + uint64_t srcOrg[4] = {srcOrigin[0] / granularity, srcOrigin[1], srcOrigin[2], 0}; setArgument(kernels_[blitType], 2, sizeof(srcOrg), srcOrg); - cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1], (cl_int)dstOrigin[2], 0}; - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t dstOrg[4] = {(int32_t)dstOrigin[0], (int32_t)dstOrigin[1], (int32_t)dstOrigin[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 3, sizeof(dstOrg), dstOrg); setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); @@ -1004,11 +1004,11 @@ bool KernelBlitManager::copyBufferToImageKernel(device::Memory& srcMemory, // Program memory format uint multiplier = memFmtSize / sizeof(uint32_t); multiplier = (multiplier == 0) ? 1 : multiplier; - cl_uint format[4] = {components, memFmtSize / components, multiplier, 0}; + uint32_t format[4] = {components, memFmtSize / components, multiplier, 0}; setArgument(kernels_[blitType], 5, sizeof(format), format); // Program row and slice pitches - cl_ulong pitch[4] = {0}; + uint64_t pitch[4] = {0}; CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, gpuMem(dstMemory)); setArgument(kernels_[blitType], 6, sizeof(pitch), pitch); @@ -1164,32 +1164,32 @@ bool KernelBlitManager::copyImageToBufferKernel(device::Memory& srcMemory, setArgument(kernels_[blitType], 2, sizeof(cl_mem), &mem); setArgument(kernels_[blitType], 3, sizeof(cl_mem), &mem); - cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1], (cl_int)srcOrigin[2], 0}; - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t srcOrg[4] = {(int32_t)srcOrigin[0], (int32_t)srcOrigin[1], (int32_t)srcOrigin[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 4, sizeof(srcOrg), srcOrg); uint32_t memFmtSize = srcImage->getImageFormat().getElementSize(); uint32_t components = srcImage->getImageFormat().getNumChannels(); // 1 element granularity for writes by default - cl_int granularity = 1; + int32_t granularity = 1; if (memFmtSize == 2) { granularity = 2; } else if (memFmtSize >= 4) { granularity = 4; } CondLog(((dstOrigin[0] % granularity) != 0), "Unaligned offset in blit!"); - cl_ulong dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2], 0}; + uint64_t dstOrg[4] = {dstOrigin[0] / granularity, dstOrigin[1], dstOrigin[2], 0}; setArgument(kernels_[blitType], 5, sizeof(dstOrg), dstOrg); setArgument(kernels_[blitType], 6, sizeof(copySize), copySize); // Program memory format uint multiplier = memFmtSize / sizeof(uint32_t); multiplier = (multiplier == 0) ? 1 : multiplier; - cl_uint format[4] = {components, memFmtSize / components, multiplier, 0}; + uint32_t format[4] = {components, memFmtSize / components, multiplier, 0}; setArgument(kernels_[blitType], 7, sizeof(format), format); // Program row and slice pitches - cl_ulong pitch[4] = {0}; + uint64_t pitch[4] = {0}; CalcRowSlicePitches(pitch, copySize, rowPitch, slicePitch, gpuMem(srcMemory)); setArgument(kernels_[blitType], 8, sizeof(pitch), pitch); @@ -1308,14 +1308,14 @@ bool KernelBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dst setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); // Program source origin - cl_int srcOrg[4] = {(cl_int)srcOrigin[0], (cl_int)srcOrigin[1], (cl_int)srcOrigin[2], 0}; + int32_t srcOrg[4] = {(int32_t)srcOrigin[0], (int32_t)srcOrigin[1], (int32_t)srcOrigin[2], 0}; setArgument(kernels_[blitType], 2, sizeof(srcOrg), srcOrg); // Program destinaiton origin - cl_int dstOrg[4] = {(cl_int)dstOrigin[0], (cl_int)dstOrigin[1], (cl_int)dstOrigin[2], 0}; + int32_t dstOrg[4] = {(int32_t)dstOrigin[0], (int32_t)dstOrigin[1], (int32_t)dstOrigin[2], 0}; setArgument(kernels_[blitType], 3, sizeof(dstOrg), dstOrg); - cl_int copySize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t copySize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); // Create ND range object for the kernel's execution @@ -1549,11 +1549,11 @@ bool KernelBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory setArgument(kernels_[blitType], 0, sizeof(cl_mem), &mem); mem = as_cl(dstMemory.owner()); setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); - cl_ulong src[4] = {srcRect.rowPitch_, srcRect.slicePitch_, srcRect.start_, 0}; + uint64_t src[4] = {srcRect.rowPitch_, srcRect.slicePitch_, srcRect.start_, 0}; setArgument(kernels_[blitType], 2, sizeof(src), src); - cl_ulong dst[4] = {dstRect.rowPitch_, dstRect.slicePitch_, dstRect.start_, 0}; + uint64_t dst[4] = {dstRect.rowPitch_, dstRect.slicePitch_, dstRect.start_, 0}; setArgument(kernels_[blitType], 3, sizeof(dst), dst); - cl_ulong copySize[4] = {size[0], size[1], size[2], CopyRectAlignment[i]}; + uint64_t copySize[4] = {size[0], size[1], size[2], CopyRectAlignment[i]}; setArgument(kernels_[blitType], 4, sizeof(copySize), copySize); // Create ND range object for the kernel's execution @@ -1772,7 +1772,7 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, } else { uint fillType = FillBuffer; size_t globalWorkOffset[3] = {0, 0, 0}; - cl_ulong fillSize = size[0] / patternSize; + uint64_t fillSize = size[0] / patternSize; size_t globalWorkSize = amd::alignUp(fillSize, 256); size_t localWorkSize = 256; bool dwordAligned = ((patternSize % sizeof(uint32_t)) == 0) ? true : false; @@ -1795,12 +1795,12 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, mem = as_cl(gpuCB->owner()); setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem); - cl_ulong offset = origin[0]; + uint64_t offset = origin[0]; if (dwordAligned) { patternSize /= sizeof(uint32_t); offset /= sizeof(uint32_t); } - setArgument(kernels_[fillType], 3, sizeof(cl_uint), &patternSize); + setArgument(kernels_[fillType], 3, sizeof(uint32_t), &patternSize); setArgument(kernels_[fillType], 4, sizeof(offset), &offset); setArgument(kernels_[fillType], 5, sizeof(fillSize), &fillSize); @@ -1854,7 +1854,7 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds } } - cl_uint remain; + uint32_t remain; if (blitType == BlitCopyBufferAligned) { size.c[0] /= CopyBuffAlignment[i]; } else { @@ -1873,20 +1873,20 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds mem = as_cl(dstMemory.owner()); setArgument(kernels_[blitType], 1, sizeof(cl_mem), &mem); // Program source origin - cl_ulong srcOffset = srcOrigin[0] / CopyBuffAlignment[i]; + uint64_t srcOffset = srcOrigin[0] / CopyBuffAlignment[i]; ; setArgument(kernels_[blitType], 2, sizeof(srcOffset), &srcOffset); // Program destinaiton origin - cl_ulong dstOffset = dstOrigin[0] / CopyBuffAlignment[i]; + uint64_t dstOffset = dstOrigin[0] / CopyBuffAlignment[i]; ; setArgument(kernels_[blitType], 3, sizeof(dstOffset), &dstOffset); - cl_ulong copySize = size[0]; + uint64_t copySize = size[0]; setArgument(kernels_[blitType], 4, sizeof(copySize), ©Size); if (blitType == BlitCopyBufferAligned) { - cl_int alignment = CopyBuffAlignment[i]; + int32_t alignment = CopyBuffAlignment[i]; setArgument(kernels_[blitType], 5, sizeof(alignment), &alignment); } else { setArgument(kernels_[blitType], 5, sizeof(remain), &remain); @@ -1935,7 +1935,7 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, dim = 3; void* newpattern = const_cast(pattern); - cl_uint4 iFillColor; + uint32_t iFillColor[4]; bool rejected = false; bool releaseView = false; @@ -1955,11 +1955,11 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, // Converting a linear RGB floating-point color value to a 8-bit unsigned integer sRGB value // because hw is not support write_imagef for sRGB. float* fColor = static_cast(newpattern); - iFillColor.s[0] = sRGBmap(fColor[0]); - iFillColor.s[1] = sRGBmap(fColor[1]); - iFillColor.s[2] = sRGBmap(fColor[2]); - iFillColor.s[3] = (cl_uint)(fColor[3] * 255.0f); - newpattern = static_cast(&iFillColor); + iFillColor[0] = sRGBmap(fColor[0]); + iFillColor[1] = sRGBmap(fColor[1]); + iFillColor[2] = sRGBmap(fColor[2]); + iFillColor[3] = (uint32_t)(fColor[3] * 255.0f); + newpattern = static_cast(&iFillColor[0]); for (uint i = 0; i < RejectedFormatChannelTotal; ++i) { if (RejectedOrder[i].clOldType_ == newFormat.image_channel_order) { newFormat.image_channel_order = RejectedOrder[i].clNewType_; @@ -2008,12 +2008,12 @@ bool KernelBlitManager::fillImage(device::Memory& memory, const void* pattern, // Program kernels arguments for the blit operation cl_mem mem = as_cl(memView->owner()); setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); - setArgument(kernels_[fillType], 1, sizeof(cl_float4), newpattern); - setArgument(kernels_[fillType], 2, sizeof(cl_int4), newpattern); - setArgument(kernels_[fillType], 3, sizeof(cl_uint4), newpattern); + setArgument(kernels_[fillType], 1, sizeof(float[4]), newpattern); + setArgument(kernels_[fillType], 2, sizeof(int32_t[4]), newpattern); + setArgument(kernels_[fillType], 3, sizeof(uint32_t[4]), newpattern); - cl_int fillOrigin[4] = {(cl_int)origin[0], (cl_int)origin[1], (cl_int)origin[2], 0}; - cl_int fillSize[4] = {(cl_int)size[0], (cl_int)size[1], (cl_int)size[2], 0}; + int32_t fillOrigin[4] = {(int32_t)origin[0], (int32_t)origin[1], (int32_t)origin[2], 0}; + int32_t fillSize[4] = {(int32_t)size[0], (int32_t)size[1], (int32_t)size[2], 0}; setArgument(kernels_[fillType], 4, sizeof(fillOrigin), fillOrigin); setArgument(kernels_[fillType], 5, sizeof(fillSize), fillSize); diff --git a/rocclr/device/rocm/roccounters.cpp b/rocclr/device/rocm/roccounters.cpp index a6e2eea11d..93a8251ee3 100644 --- a/rocclr/device/rocm/roccounters.cpp +++ b/rocclr/device/rocm/roccounters.cpp @@ -419,9 +419,9 @@ static const std::array, 139 //! Constructor for the ROC PerfCounter object PerfCounter::PerfCounter(const Device& device, //!< A ROC device object - cl_uint blockIndex, //!< HW block index - cl_uint counterIndex, //!< Counter index (Counter register) within the block - cl_uint eventIndex) //!< Event index (Counter selection) for profiling + uint32_t blockIndex, //!< HW block index + uint32_t counterIndex, //!< Counter index (Counter register) within the block + uint32_t eventIndex) //!< Event index (Counter selection) for profiling : roc_device_(device), profileRef_(nullptr) { diff --git a/rocclr/device/rocm/roccounters.hpp b/rocclr/device/rocm/roccounters.hpp index 38963f8bcf..5719cde663 100644 --- a/rocclr/device/rocm/roccounters.hpp +++ b/rocclr/device/rocm/roccounters.hpp @@ -50,9 +50,9 @@ class PerfCounter : public device::PerfCounter { //! Constructor for the ROC PerfCounter object PerfCounter(const Device& device, //!< A ROC device object - cl_uint blockIndex, //!< HW block index - cl_uint counterIndex, //!< Counter index (Counter register) within the block - cl_uint eventIndex); //!< Event index (Counter selection) for profiling + uint32_t blockIndex, //!< HW block index + uint32_t counterIndex, //!< Counter index (Counter register) within the block + uint32_t eventIndex); //!< Event index (Counter selection) for profiling //! Destructor for the ROCM PerfCounter object virtual ~PerfCounter(); diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 95b102c55a..3d0a50e0a0 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -1129,13 +1129,13 @@ bool Device::populateOCLDeviceConstants() { } assert(global_segment_size > 0); - info_.globalMemSize_ = static_cast(global_segment_size); + info_.globalMemSize_ = static_cast(global_segment_size); gpuvm_segment_max_alloc_ = - cl_ulong(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); + uint64_t(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); assert(gpuvm_segment_max_alloc_ > 0); - info_.maxMemAllocSize_ = static_cast(gpuvm_segment_max_alloc_); + info_.maxMemAllocSize_ = static_cast(gpuvm_segment_max_alloc_); if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info(gpuvm_segment_, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, @@ -1147,10 +1147,10 @@ bool Device::populateOCLDeviceConstants() { } else { // We suppose half of physical memory can be used by GPU in APU system info_.globalMemSize_ = - cl_ulong(sysconf(_SC_PAGESIZE)) * cl_ulong(sysconf(_SC_PHYS_PAGES)) / 2; - info_.globalMemSize_ = std::max(info_.globalMemSize_, cl_ulong(1 * Gi)); + uint64_t(sysconf(_SC_PAGESIZE)) * uint64_t(sysconf(_SC_PHYS_PAGES)) / 2; + info_.globalMemSize_ = std::max(info_.globalMemSize_, uint64_t(1 * Gi)); info_.maxMemAllocSize_ = - cl_ulong(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); + uint64_t(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u); if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info( @@ -1203,8 +1203,8 @@ bool Device::populateOCLDeviceConstants() { info_.hostUnifiedMemory_ = CL_TRUE; } info_.memBaseAddrAlign_ = - 8 * (flagIsDefault(MEMOBJ_BASE_ADDR_ALIGN) ? sizeof(cl_long16) : MEMOBJ_BASE_ADDR_ALIGN); - info_.minDataTypeAlignSize_ = sizeof(cl_long16); + 8 * (flagIsDefault(MEMOBJ_BASE_ADDR_ALIGN) ? sizeof(int64_t[16]) : MEMOBJ_BASE_ADDR_ALIGN); + info_.minDataTypeAlignSize_ = sizeof(int64_t[16]); info_.maxConstantArgs_ = 8; info_.preferredConstantBufferSize_ = 16 * Ki; diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 6c3e9dbf8d..530451cdeb 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -1137,7 +1137,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) { const std::vector& svmPointers = cmd.svmPointers(); if (cmd.pfnFreeFunc() == nullptr) { // pointers allocated using clSVMAlloc - for (cl_uint i = 0; i < svmPointers.size(); i++) { + for (uint32_t i = 0; i < svmPointers.size(); i++) { amd::SvmBuffer::free(cmd.context(), svmPointers[i]); } } else { diff --git a/rocclr/include/vdi_agent_amd.h b/rocclr/include/vdi_agent_amd.h index c1ef2c1706..6bae5262ed 100644 --- a/rocclr/include/vdi_agent_amd.h +++ b/rocclr/include/vdi_agent_amd.h @@ -56,8 +56,8 @@ typedef void(CL_CALLBACK* acEventCreate_fn)(vdi_agent* /* agent */, cl_event /* typedef void(CL_CALLBACK* acEventFree_fn)(vdi_agent* /* agent */, cl_event /* event */); typedef void(CL_CALLBACK* acEventStatusChanged_fn)(vdi_agent* /* agent */, cl_event /* event */, - cl_int /* execution_status */, - cl_long /* epoch_time_stamp */); + int32_t /* execution_status */, + int64_t /* epoch_time_stamp */); /* Memory Object Callbacks */ @@ -67,7 +67,7 @@ typedef void(CL_CALLBACK* acMemObjectFree_fn)(vdi_agent* /* agent */, cl_mem /* typedef void(CL_CALLBACK* acMemObjectAcquired_fn)(vdi_agent* /* agent */, cl_mem /* memobj */, cl_device_id /* device */, - cl_long /* elapsed_time */); + int64_t /* elapsed_time */); /* Sampler Callbacks */ @@ -90,7 +90,7 @@ typedef void(CL_CALLBACK* acKernelCreate_fn)(vdi_agent* /* agent */, cl_kernel / typedef void(CL_CALLBACK* acKernelFree_fn)(vdi_agent* /* agent */, cl_kernel /* kernel */); typedef void(CL_CALLBACK* acKernelSetArg_fn)(vdi_agent* /* agent */, cl_kernel /* kernel */, - cl_int /* arg_index */, size_t /* size */, + int32_t /* arg_index */, size_t /* size */, const void* /* value_ptr */); typedef struct _vdi_agent_callbacks { @@ -128,55 +128,55 @@ typedef struct _vdi_agent_callbacks { } vdi_agent_callbacks; -typedef cl_uint vdi_agent_capability_action; +typedef uint32_t vdi_agent_capability_action; #define VDI_AGENT_ADD_CAPABILITIES 0x0 #define VDI_AGENT_RELINQUISH_CAPABILITIES 0x1 typedef struct _vdi_agent_capabilities { - cl_bitfield canGenerateContextEvents : 1; - cl_bitfield canGenerateCommandQueueEvents : 1; - cl_bitfield canGenerateEventEvents : 1; - cl_bitfield canGenerateMemObjectEvents : 1; - cl_bitfield canGenerateSamplerEvents : 1; - cl_bitfield canGenerateProgramEvents : 1; - cl_bitfield canGenerateKernelEvents : 1; + uint64_t canGenerateContextEvents : 1; + uint64_t canGenerateCommandQueueEvents : 1; + uint64_t canGenerateEventEvents : 1; + uint64_t canGenerateMemObjectEvents : 1; + uint64_t canGenerateSamplerEvents : 1; + uint64_t canGenerateProgramEvents : 1; + uint64_t canGenerateKernelEvents : 1; } vdi_agent_capabilities; struct _vdi_agent { - cl_int(CL_API_CALL* GetVersionNumber)(vdi_agent* /* agent */, cl_int* /* version_ret */); + int32_t(CL_API_CALL* GetVersionNumber)(vdi_agent* /* agent */, int32_t* /* version_ret */); - cl_int(CL_API_CALL* GetPlatform)(vdi_agent* /* agent */, cl_platform_id* /* platform_id_ret */); + int32_t(CL_API_CALL* GetPlatform)(vdi_agent* /* agent */, cl_platform_id* /* platform_id_ret */); - cl_int(CL_API_CALL* GetTime)(vdi_agent* /* agent */, cl_long* /* time_nanos */); + int32_t(CL_API_CALL* GetTime)(vdi_agent* /* agent */, int64_t* /* time_nanos */); - cl_int(CL_API_CALL* SetCallbacks)(vdi_agent* /* agent */, + int32_t(CL_API_CALL* SetCallbacks)(vdi_agent* /* agent */, const vdi_agent_callbacks* /* callbacks */, size_t /* size */); - cl_int(CL_API_CALL* GetPotentialCapabilities)(vdi_agent* /* agent */, + int32_t(CL_API_CALL* GetPotentialCapabilities)(vdi_agent* /* agent */, vdi_agent_capabilities* /* capabilities */); - cl_int(CL_API_CALL* GetCapabilities)(vdi_agent* /* agent */, + int32_t(CL_API_CALL* GetCapabilities)(vdi_agent* /* agent */, vdi_agent_capabilities* /* capabilities */); - cl_int(CL_API_CALL* SetCapabilities)(vdi_agent* /* agent */, + int32_t(CL_API_CALL* SetCapabilities)(vdi_agent* /* agent */, const vdi_agent_capabilities* /* capabilities */, vdi_agent_capability_action /* action */); - cl_int(CL_API_CALL* GetICDDispatchTable)(vdi_agent* /* agent */, + int32_t(CL_API_CALL* GetICDDispatchTable)(vdi_agent* /* agent */, cl_icd_dispatch_table* /* table */, size_t /* size */); - cl_int(CL_API_CALL* SetICDDispatchTable)(vdi_agent* /* agent */, + int32_t(CL_API_CALL* SetICDDispatchTable)(vdi_agent* /* agent */, const cl_icd_dispatch_table* /* table */, size_t /* size */); /* add Kernel/Program helper functions, etc... */ }; -extern cl_int CL_CALLBACK vdiAgent_OnLoad(vdi_agent* /* agent */); +extern int32_t CL_CALLBACK vdiAgent_OnLoad(vdi_agent* /* agent */); extern void CL_CALLBACK vdiAgent_OnUnload(vdi_agent* /* agent */); diff --git a/rocclr/platform/agent.cpp b/rocclr/platform/agent.cpp index 4958b2fc98..72c9e48a48 100644 --- a/rocclr/platform/agent.cpp +++ b/rocclr/platform/agent.cpp @@ -32,7 +32,7 @@ namespace amd { -typedef cl_int(CL_CALLBACK* vdiAgent_OnLoad_fn)(vdi_agent* agent); +typedef int32_t(CL_CALLBACK* vdiAgent_OnLoad_fn)(vdi_agent* agent); typedef void(CL_CALLBACK* vdiAgent_OnUnload_fn)(vdi_agent* agent); Agent::Agent(const char* moduleName) : ready_(false) { @@ -78,13 +78,13 @@ Agent::~Agent() { } } -cl_int Agent::setCallbacks(const vdi_agent_callbacks* callbacks, size_t size) { +int32_t Agent::setCallbacks(const vdi_agent_callbacks* callbacks, size_t size) { // FIXME_lmoriche: check size memcpy(&callbacks_, callbacks, size); return CL_SUCCESS; } -cl_int Agent::getCapabilities(vdi_agent_capabilities* caps) { +int32_t Agent::getCapabilities(vdi_agent_capabilities* caps) { if (caps == NULL) { return CL_INVALID_VALUE; } @@ -148,7 +148,7 @@ static inline bool operator!=(const vdi_agent_capabilities& lhs, const vdi_agent return !(lhs == rhs); } -cl_int Agent::setCapabilities(const vdi_agent_capabilities* caps, bool install) { +int32_t Agent::setCapabilities(const vdi_agent_capabilities* caps, bool install) { ScopedLock sl(capabilitiesLock_); if (caps == NULL || *caps != (*caps & potentialCapabilities_)) { @@ -211,7 +211,7 @@ void Agent::tearDown() { namespace agent { -static cl_int CL_API_CALL GetVersionNumber(vdi_agent* agent, cl_int* version_ret) { +static int32_t CL_API_CALL GetVersionNumber(vdi_agent* agent, int32_t* version_ret) { if (version_ret == NULL) { return CL_INVALID_VALUE; } @@ -219,7 +219,7 @@ static cl_int CL_API_CALL GetVersionNumber(vdi_agent* agent, cl_int* version_ret return CL_SUCCESS; } -static cl_int CL_API_CALL GetPlatform(vdi_agent* agent, cl_platform_id* platform_id_ret) { +static int32_t CL_API_CALL GetPlatform(vdi_agent* agent, cl_platform_id* platform_id_ret) { if (platform_id_ret == NULL) { return CL_INVALID_VALUE; } @@ -227,7 +227,7 @@ static cl_int CL_API_CALL GetPlatform(vdi_agent* agent, cl_platform_id* platform return CL_SUCCESS; } -static cl_int CL_API_CALL GetTime(vdi_agent* agent, cl_long* time_nanos) { +static int32_t CL_API_CALL GetTime(vdi_agent* agent, int64_t* time_nanos) { if (time_nanos == NULL) { return CL_INVALID_VALUE; } @@ -235,12 +235,12 @@ static cl_int CL_API_CALL GetTime(vdi_agent* agent, cl_long* time_nanos) { return CL_SUCCESS; } -static cl_int CL_API_CALL SetCallbacks(vdi_agent* agent, const vdi_agent_callbacks* callbacks, +static int32_t CL_API_CALL SetCallbacks(vdi_agent* agent, const vdi_agent_callbacks* callbacks, size_t size) { return Agent::get(agent)->setCallbacks(callbacks, size); } -static cl_int CL_API_CALL GetPotentialCapabilities(vdi_agent* agent, +static int32_t CL_API_CALL GetPotentialCapabilities(vdi_agent* agent, vdi_agent_capabilities* capabilities) { if (capabilities == NULL) { return CL_INVALID_VALUE; @@ -250,24 +250,24 @@ static cl_int CL_API_CALL GetPotentialCapabilities(vdi_agent* agent, return CL_SUCCESS; } -static cl_int CL_API_CALL GetCapabilities(vdi_agent* agent, vdi_agent_capabilities* capabilities) { +static int32_t CL_API_CALL GetCapabilities(vdi_agent* agent, vdi_agent_capabilities* capabilities) { return Agent::get(agent)->getCapabilities(capabilities); } -static cl_int CL_API_CALL SetCapabilities(vdi_agent* agent, +static int32_t CL_API_CALL SetCapabilities(vdi_agent* agent, const vdi_agent_capabilities* capabilities, vdi_agent_capability_action action) { return Agent::get(agent)->setCapabilities(capabilities, action == VDI_AGENT_ADD_CAPABILITIES); } -static cl_int CL_API_CALL GetICDDispatchTable(vdi_agent* agent, cl_icd_dispatch_table* table, +static int32_t CL_API_CALL GetICDDispatchTable(vdi_agent* agent, cl_icd_dispatch_table* table, size_t size) { // FIXME_lmoriche: check size memcpy(table, amd::ICDDispatchedObject::icdVendorDispatch_, size); return CL_SUCCESS; } -static cl_int CL_API_CALL SetICDDispatchTable(vdi_agent* agent, const cl_icd_dispatch_table* table, +static int32_t CL_API_CALL SetICDDispatchTable(vdi_agent* agent, const cl_icd_dispatch_table* table, size_t size) { // FIXME_lmoriche: check size memcpy(amd::ICDDispatchedObject::icdVendorDispatch_, table, size); @@ -340,7 +340,7 @@ void Agent::postEventFree(cl_event event) { } } -void Agent::postEventStatusChanged(cl_event event, cl_int status, cl_long ts) { +void Agent::postEventStatusChanged(cl_event event, int32_t status, int64_t ts) { for (Agent* agent = list_; agent != NULL; agent = agent->next_) { acEventStatusChanged_fn callback = agent->callbacks_.EventStatusChanged; if (callback != NULL && agent->canGenerateEventEvents()) { @@ -367,7 +367,7 @@ void Agent::postMemObjectFree(cl_mem memobj) { } } -void Agent::postMemObjectAcquired(cl_mem memobj, cl_device_id device, cl_long elapsed) { +void Agent::postMemObjectAcquired(cl_mem memobj, cl_device_id device, int64_t elapsed) { for (Agent* agent = list_; agent != NULL; agent = agent->next_) { acMemObjectAcquired_fn callback = agent->callbacks_.MemObjectAcquired; if (callback != NULL && agent->canGenerateMemObjectEvents()) { @@ -439,7 +439,7 @@ void Agent::postKernelFree(cl_kernel kernel) { } } -void Agent::postKernelSetArg(cl_kernel kernel, cl_int index, size_t size, const void* value_ptr) { +void Agent::postKernelSetArg(cl_kernel kernel, int32_t index, size_t size, const void* value_ptr) { for (Agent* agent = list_; agent != NULL; agent = agent->next_) { acKernelSetArg_fn callback = agent->callbacks_.KernelSetArg; if (callback != NULL && agent->canGenerateKernelEvents()) { diff --git a/rocclr/platform/agent.hpp b/rocclr/platform/agent.hpp index 416ed7b10f..1a02f446c1 100644 --- a/rocclr/platform/agent.hpp +++ b/rocclr/platform/agent.hpp @@ -77,15 +77,15 @@ class Agent : public _vdi_agent { //! Post an event destruction event static void postEventFree(cl_event event); //! Post and event status change event. - static void postEventStatusChanged(cl_event event, cl_int execution_status, - cl_long epoch_timestamp); + static void postEventStatusChanged(cl_event event, int32_t execution_status, + int64_t epoch_timestamp); //! Post a memory object creation event static void postMemObjectCreate(cl_mem memobj); //! Post a memory object destruction event static void postMemObjectFree(cl_mem memobj); //! Post a memory transfer (acquired by device) event - static void postMemObjectAcquired(cl_mem memobj, cl_device_id device, cl_long elapsed_time); + static void postMemObjectAcquired(cl_mem memobj, cl_device_id device, int64_t elapsed_time); //! Post a sampler creation event static void postSamplerCreate(cl_sampler sampler); @@ -104,7 +104,7 @@ class Agent : public _vdi_agent { //! Post a kernel destruction event static void postKernelFree(cl_kernel kernel); //! Post a kernel set argument event - static void postKernelSetArg(cl_kernel kernel, cl_int arg_index, size_t size, + static void postKernelSetArg(cl_kernel kernel, int32_t arg_index, size_t size, const void* value_ptr); private: @@ -140,12 +140,12 @@ class Agent : public _vdi_agent { bool isReady() const { return ready_; } //! Set the callback vector for this agent - cl_int setCallbacks(const vdi_agent_callbacks* callbacks, size_t size); + int32_t setCallbacks(const vdi_agent_callbacks* callbacks, size_t size); //! Return the current capabilities. - cl_int getCapabilities(vdi_agent_capabilities* caps); + int32_t getCapabilities(vdi_agent_capabilities* caps); //! Set the current capabilities. - cl_int setCapabilities(const vdi_agent_capabilities* caps, bool install); + int32_t setCapabilities(const vdi_agent_capabilities* caps, bool install); //! Return the Agent instance from the given cl_agent inline static Agent* get(vdi_agent* agent) { diff --git a/rocclr/platform/command.cpp b/rocclr/platform/command.cpp index 406c0d52e4..b6f496930a 100644 --- a/rocclr/platform/command.cpp +++ b/rocclr/platform/command.cpp @@ -60,7 +60,7 @@ Event::~Event() { } } -uint64_t Event::recordProfilingInfo(cl_int status, uint64_t timeStamp) { +uint64_t Event::recordProfilingInfo(int32_t status, uint64_t timeStamp) { if (timeStamp == 0) { timeStamp = Os::timeNanos(); } @@ -85,10 +85,10 @@ uint64_t Event::recordProfilingInfo(cl_int status, uint64_t timeStamp) { return timeStamp; } -bool Event::setStatus(cl_int status, uint64_t timeStamp) { +bool Event::setStatus(int32_t status, uint64_t timeStamp) { assert(status <= CL_QUEUED && "invalid status"); - cl_int currentStatus = status_; + int32_t currentStatus = status_; if (currentStatus <= CL_COMPLETE || currentStatus <= status) { // We can only move forward in the execution status. return false; @@ -130,7 +130,7 @@ bool Event::setStatus(cl_int status, uint64_t timeStamp) { } -bool Event::setCallback(cl_int status, Event::CallBackFunction callback, void* data) { +bool Event::setCallback(int32_t status, Event::CallBackFunction callback, void* data) { assert(status >= CL_COMPLETE && status <= CL_QUEUED && "invalid status"); CallBackEntry* entry = new CallBackEntry(status, callback, data); @@ -153,9 +153,9 @@ bool Event::setCallback(cl_int status, Event::CallBackFunction callback, void* d } -void Event::processCallbacks(cl_int status) const { +void Event::processCallbacks(int32_t status) const { cl_event event = const_cast(as_cl(this)); - const cl_int mask = (status > CL_COMPLETE) ? status : CL_COMPLETE; + const int32_t mask = (status > CL_COMPLETE) ? status : CL_COMPLETE; // For_each callback: CallBackEntry* entry; @@ -307,7 +307,7 @@ NativeFnCommand::NativeFnCommand(HostQueue& queue, const EventWaitList& eventWai } } -cl_int NativeFnCommand::invoke() { +int32_t NativeFnCommand::invoke() { size_t numMemObjs = memObjects_.size(); for (size_t i = 0; i < numMemObjs; ++i) { void* hostMemPtr = memObjects_[i]->getHostMem(); @@ -431,15 +431,15 @@ bool MigrateMemObjectsCommand::validateMemory() { return true; } -cl_int NDRangeKernelCommand::captureAndValidate() { +int32_t NDRangeKernelCommand::captureAndValidate() { const amd::Device& device = queue()->device(); // Validate the kernel before submission if (!queue()->device().validateKernel(kernel(), queue()->vdev(), cooperativeGroups())) { return CL_OUT_OF_RESOURCES; } - cl_int error; - cl_ulong lclMemSize = kernel().getDeviceKernel(device)->workGroupInfo()->localMemSize_; + int32_t error; + uint64_t lclMemSize = kernel().getDeviceKernel(device)->workGroupInfo()->localMemSize_; parameters_ = kernel().parameters().capture(device, lclMemSize, &error); return error; } diff --git a/rocclr/platform/command.hpp b/rocclr/platform/command.hpp index 9412f95828..cf1bd59546 100644 --- a/rocclr/platform/command.hpp +++ b/rocclr/platform/command.hpp @@ -69,7 +69,7 @@ class HostQueue; * in a Context. */ class Event : public RuntimeObject { - typedef void(CL_CALLBACK* CallBackFunction)(cl_event event, cl_int command_exec_status, + typedef void(CL_CALLBACK* CallBackFunction)(cl_event event, int32_t command_exec_status, void* user_data); struct CallBackEntry : public HeapObject { @@ -77,9 +77,9 @@ class Event : public RuntimeObject { std::atomic callback_; //!< callback function pointer. void* data_; //!< user data passed to the callback function. - cl_int status_; //!< execution status triggering the callback. + int32_t status_; //!< execution status triggering the callback. - CallBackEntry(cl_int status, CallBackFunction callback, void* data) + CallBackEntry(int32_t status, CallBackFunction callback, void* data) : callback_(callback), data_(data), status_(status) {} }; @@ -90,7 +90,7 @@ class Event : public RuntimeObject { Monitor lock_; std::atomic callbacks_; //!< linked list of callback entries. - volatile cl_int status_; //!< current execution status. + volatile int32_t status_; //!< current execution status. std::atomic_flag notified_; //!< Command queue was notified protected: @@ -146,10 +146,10 @@ class Event : public RuntimeObject { //! Record the profiling info for the given change of \a status. // If the given \a timeStamp is 0 and profiling is enabled, // use the current host clock time instead. - uint64_t recordProfilingInfo(cl_int status, uint64_t timeStamp = 0); + uint64_t recordProfilingInfo(int32_t status, uint64_t timeStamp = 0); //! Process the callbacks for the given \a status change. - void processCallbacks(cl_int status) const; + void processCallbacks(int32_t status) const; public: //! Return the context for this event. @@ -163,10 +163,10 @@ class Event : public RuntimeObject { const ProfilingInfo& profilingInfo() const { return profilingInfo_; } //! Return this command's execution status. - cl_int status() const { return status_; } + int32_t status() const { return status_; } //! Insert the given \a callback into the callback stack. - bool setCallback(cl_int status, CallBackFunction callback, void* data); + bool setCallback(int32_t status, CallBackFunction callback, void* data); /*! \brief Set the event status. * @@ -176,7 +176,7 @@ class Event : public RuntimeObject { * * \see amd::Event::awaitCompletion */ - bool setStatus(cl_int status, uint64_t timeStamp = 0); + bool setStatus(int32_t status, uint64_t timeStamp = 0); //! Signal all threads waiting on this event. void signal() { @@ -214,7 +214,7 @@ class Command : public Event { Command* next_; const cl_command_type type_; //!< This command's OpenCL type. - volatile cl_int exception_; //!< The first raised exception. + volatile int32_t exception_; //!< The first raised exception. void* data_; protected: @@ -265,10 +265,10 @@ class Command : public Event { cl_command_type type() const { return type_; } //! Return the first raised exception or 0 if none. - cl_int exception() const { return exception_; } + int32_t exception() const { return exception_; } //! Set the exception for this command. - void setException(cl_int exception) { exception_ = exception; } + void setException(int32_t exception) { exception_ = exception; } //! Return the opaque, device specific data for this command. void* data() const { return data_; } @@ -551,7 +551,7 @@ class WriteMemoryCommand : public OneMemoryArgCommand { class FillMemoryCommand : public OneMemoryArgCommand { public: - const static size_t MaxFillPatterSize = sizeof(cl_double16); + const static size_t MaxFillPatterSize = sizeof(double[16]); private: Coord3D origin_; //!< Origin of the region to write to. @@ -768,7 +768,7 @@ class MigrateMemObjectsCommand : public Command { //! Returns the migration flags cl_mem_migration_flags migrationFlags() const { return migrationFlags_; } //! Returns the number of memory objects in the command - cl_uint numMemObjects() const { return (cl_uint)memObjects_.size(); } + uint32_t numMemObjects() const { return (uint32_t)memObjects_.size(); } //! Returns a pointer to the memory objects const std::vector& memObjects() const { return memObjects_; } @@ -845,7 +845,7 @@ class NDRangeKernelCommand : public Command { //! Set the local work size. void setLocalWorkSize(const NDRange& local) { sizes_.local() = local; } - cl_int captureAndValidate(); + int32_t captureAndValidate(); }; class NativeFnCommand : public Command { @@ -872,7 +872,7 @@ class NativeFnCommand : public Command { virtual void submit(device::VirtualDevice& device) { device.submitNativeFn(*this); } - cl_int invoke(); + int32_t invoke(); }; class Marker : public Command { @@ -902,7 +902,7 @@ class ExtObjectsCommand : public Command { public: //! Construct a new AcquireExtObjectsCommand - ExtObjectsCommand(HostQueue& queue, const EventWaitList& eventWaitList, cl_uint num_objects, + ExtObjectsCommand(HostQueue& queue, const EventWaitList& eventWaitList, uint32_t num_objects, const std::vector& memoryObjects, cl_command_type type) : Command(queue, type, eventWaitList) { for (const auto& it : memoryObjects) { @@ -920,7 +920,7 @@ class ExtObjectsCommand : public Command { } //! Get number of GL objects - cl_uint getNumObjects() { return (cl_uint)memObjects_.size(); } + uint32_t getNumObjects() { return (uint32_t)memObjects_.size(); } //! Get pointer to GL object list const std::vector& getMemList() const { return memObjects_; } bool validateMemory(); @@ -931,7 +931,7 @@ class AcquireExtObjectsCommand : public ExtObjectsCommand { public: //! Construct a new AcquireExtObjectsCommand AcquireExtObjectsCommand(HostQueue& queue, const EventWaitList& eventWaitList, - cl_uint num_objects, const std::vector& memoryObjects, + uint32_t num_objects, const std::vector& memoryObjects, cl_command_type type) : ExtObjectsCommand(queue, eventWaitList, num_objects, memoryObjects, type) {} @@ -944,7 +944,7 @@ class ReleaseExtObjectsCommand : public ExtObjectsCommand { public: //! Construct a new ReleaseExtObjectsCommand ReleaseExtObjectsCommand(HostQueue& queue, const EventWaitList& eventWaitList, - cl_uint num_objects, const std::vector& memoryObjects, + uint32_t num_objects, const std::vector& memoryObjects, cl_command_type type) : ExtObjectsCommand(queue, eventWaitList, num_objects, memoryObjects, type) {} @@ -1027,7 +1027,7 @@ class ThreadTraceMemObjectsCommand : public Command { } //! Get number of CL memory objects - cl_uint getNumObjects() { return (cl_uint)memObjects_.size(); } + uint32_t getNumObjects() { return (uint32_t)memObjects_.size(); } //! Get pointer to CL memory object list const std::vector& getMemList() const { return memObjects_; } @@ -1104,21 +1104,21 @@ class ThreadTraceCommand : public Command { class SignalCommand : public OneMemoryArgCommand { private: - cl_uint markerValue_; - cl_ulong markerOffset_; + uint32_t markerValue_; + uint64_t markerOffset_; public: SignalCommand(HostQueue& queue, cl_command_type cmdType, const EventWaitList& eventWaitList, - Memory& memory, cl_uint value, cl_ulong offset = 0) + Memory& memory, uint32_t value, uint64_t offset = 0) : OneMemoryArgCommand(queue, cmdType, eventWaitList, memory), markerValue_(value), markerOffset_(offset) {} virtual void submit(device::VirtualDevice& device) { device.submitSignal(*this); } - const cl_uint markerValue() { return markerValue_; } + const uint32_t markerValue() { return markerValue_; } Memory& memory() { return *memory_; } - const cl_ulong markerOffset() { return markerOffset_; } + const uint64_t markerOffset() { return markerOffset_; } }; class MakeBuffersResidentCommand : public Command { @@ -1155,7 +1155,7 @@ class MakeBuffersResidentCommand : public Command { //! A deallocation command used to free SVM or system pointers. class SvmFreeMemoryCommand : public Command { public: - typedef void(CL_CALLBACK* freeCallBack)(cl_command_queue, cl_uint, void**, void*); + typedef void(CL_CALLBACK* freeCallBack)(cl_command_queue, uint32_t, void**, void*); private: std::vector svmPointers_; //!< List of pointers to deallocate @@ -1163,7 +1163,7 @@ class SvmFreeMemoryCommand : public Command { void* userData_; //!< Data passed to user-defined callback public: - SvmFreeMemoryCommand(HostQueue& queue, const EventWaitList& eventWaitList, cl_uint numSvmPointers, + SvmFreeMemoryCommand(HostQueue& queue, const EventWaitList& eventWaitList, uint32_t numSvmPointers, void** svmPointers, freeCallBack pfnFreeFunc, void* userData) : Command(queue, CL_COMMAND_SVM_FREE, eventWaitList), //! We copy svmPointers since it can be reused/deallocated after diff --git a/rocclr/platform/kernel.cpp b/rocclr/platform/kernel.cpp index 299b92038c..9d9a187073 100644 --- a/rocclr/platform/kernel.cpp +++ b/rocclr/platform/kernel.cpp @@ -153,7 +153,7 @@ void KernelParameters::set(size_t index, size_t size, const void* value, bool sv desc.info_.defined_ = true; } -address KernelParameters::capture(const Device& device, cl_ulong lclMemSize, cl_int* error) { +address KernelParameters::capture(const Device& device, uint64_t lclMemSize, int32_t* error) { *error = CL_SUCCESS; //! Information about which arguments are SVM pointers is stored after // the actual parameters, but only if the device has any SVM capability diff --git a/rocclr/platform/kernel.hpp b/rocclr/platform/kernel.hpp index 653339a8ae..fdc58e14f6 100644 --- a/rocclr/platform/kernel.hpp +++ b/rocclr/platform/kernel.hpp @@ -210,7 +210,7 @@ class KernelParameters : protected HeapObject { size_t localMemSize(size_t minDataTypeAlignment) const; //! Capture the state of the parameters and return the stack base pointer. - address capture(const Device& device, cl_ulong lclMemSize, cl_int* error); + address capture(const Device& device, uint64_t lclMemSize, int32_t* error); //! Release the captured state of the parameters. void release(address parameters, const amd::Device& device) const; diff --git a/rocclr/platform/memory.cpp b/rocclr/platform/memory.cpp index 70e4c8a0a2..2ee2bb1986 100644 --- a/rocclr/platform/memory.cpp +++ b/rocclr/platform/memory.cpp @@ -932,12 +932,12 @@ cl_image_format Image::supportedFormats[] = { {CL_DEPTH, CL_FLOAT}, }; -const cl_uint NUM_CHANNEL_ORDER_OF_RGB = 1; // The number of channel orders of RGB at the end of +const uint32_t NUM_CHANNEL_ORDER_OF_RGB = 1; // The number of channel orders of RGB at the end of // the table supportedFormats above and before sRGB and // depth. -const cl_uint NUM_CHANNEL_ORDER_OF_sRGB = 1; // The number of channel orders of sRGB at the end of +const uint32_t NUM_CHANNEL_ORDER_OF_sRGB = 1; // The number of channel orders of sRGB at the end of // the table supportedFormats above and before depth. -const cl_uint NUM_CHANNEL_ORDER_OF_DEPTH = +const uint32_t NUM_CHANNEL_ORDER_OF_DEPTH = 2; // The number of channel orders of DEPTH at the end of the table supportedFormats above. // definition of list of supported RA formats @@ -953,7 +953,7 @@ cl_image_format Image::supportedDepthStencilFormats[] = { {CL_DEPTH_STENCIL, CL_FLOAT}, {CL_DEPTH_STENCIL, CL_UNORM_INT24}}; -cl_uint Image::numSupportedFormats(const Context& context, cl_mem_object_type image_type, +uint32_t Image::numSupportedFormats(const Context& context, cl_mem_object_type image_type, cl_mem_flags flags) { const std::vector& devices = context.devices(); uint numFormats = sizeof(supportedFormats) / sizeof(cl_image_format); @@ -1007,8 +1007,8 @@ cl_uint Image::numSupportedFormats(const Context& context, cl_mem_object_type im return numFormats; } -cl_uint Image::getSupportedFormats(const Context& context, cl_mem_object_type image_type, - const cl_uint num_entries, cl_image_format* image_formats, +uint32_t Image::getSupportedFormats(const Context& context, cl_mem_object_type image_type, + const uint32_t num_entries, cl_image_format* image_formats, cl_mem_flags flags) { const std::vector& devices = context.devices(); uint numFormats = 0; @@ -1203,9 +1203,9 @@ static int round_to_even(float v) { static uint16_t float2half_rtz(float f) { union { float f; - cl_uint u; + uint32_t u; } u = {f}; - cl_uint sign = (u.u >> 16) & 0x8000; + uint32_t sign = (u.u >> 16) & 0x8000; float x = fabsf(f); // Nan diff --git a/rocclr/platform/memory.hpp b/rocclr/platform/memory.hpp index 28a76534e6..0775639823 100644 --- a/rocclr/platform/memory.hpp +++ b/rocclr/platform/memory.hpp @@ -426,10 +426,10 @@ class Image : public Memory { static cl_image_format supportedFormats[]; static cl_image_format supportedFormatsRA[]; static cl_image_format supportedDepthStencilFormats[]; - static cl_uint numSupportedFormats(const Context& context, cl_mem_object_type image_type, + static uint32_t numSupportedFormats(const Context& context, cl_mem_object_type image_type, cl_mem_flags flags = 0); - static cl_uint getSupportedFormats(const Context& context, cl_mem_object_type image_type, - const cl_uint num_entries, cl_image_format* image_formats, + static uint32_t getSupportedFormats(const Context& context, cl_mem_object_type image_type, + const uint32_t num_entries, cl_image_format* image_formats, cl_mem_flags flags = 0); //! Helper struct to manipulate image formats. diff --git a/rocclr/platform/program.cpp b/rocclr/platform/program.cpp index fbd55f68c9..fe9940796b 100644 --- a/rocclr/platform/program.cpp +++ b/rocclr/platform/program.cpp @@ -81,7 +81,7 @@ const Symbol* Program::findSymbol(const char* kernelName) const { return (it == symbolTable_->cend()) ? NULL : &it->second; } -cl_int Program::addDeviceProgram(Device& device, const void* image, size_t length, +int32_t Program::addDeviceProgram(Device& device, const void* image, size_t length, amd::option::Options* options) { if (image != NULL && !amd::isElfMagic((const char*)image)) { if (device.settings().useLightning_) { @@ -185,14 +185,14 @@ device::Program* Program::getDeviceProgram(const Device& device) const { Monitor Program::buildLock_("OCL build program", true); -cl_int Program::compile(const std::vector& devices, size_t numHeaders, +int32_t Program::compile(const std::vector& devices, size_t numHeaders, const std::vector& headerPrograms, const char** headerIncludeNames, const char* options, void(CL_CALLBACK* notifyFptr)(cl_program, void*), void* data, bool optionChangable) { ScopedLock sl(buildLock_); - cl_int retval = CL_SUCCESS; + int32_t retval = CL_SUCCESS; // Clear the program object clear(); @@ -247,7 +247,7 @@ cl_int Program::compile(const std::vector& devices, size_t numHeaders, if (sourceCode_.empty()) { return CL_INVALID_OPERATION; } - cl_int result = + int32_t result = devProgram->compile(sourceCode_, headers, headerIncludeNames, options, &parsedOptions); // Check if the previous device failed a build @@ -267,12 +267,12 @@ cl_int Program::compile(const std::vector& devices, size_t numHeaders, return retval; } -cl_int Program::link(const std::vector& devices, size_t numInputs, +int32_t Program::link(const std::vector& devices, size_t numInputs, const std::vector& inputPrograms, const char* options, void(CL_CALLBACK* notifyFptr)(cl_program, void*), void* data, bool optionChangable) { ScopedLock sl(buildLock_); - cl_int retval = CL_SUCCESS; + int32_t retval = CL_SUCCESS; if (symbolTable_ == NULL) { symbolTable_ = new symbols_t; @@ -369,7 +369,7 @@ cl_int Program::link(const std::vector& devices, size_t numInputs, if (devProgram->buildStatus() != CL_BUILD_NONE) { continue; } - cl_int result = devProgram->link(inputDevPrograms, options, &parsedOptions); + int32_t result = devProgram->link(inputDevPrograms, options, &parsedOptions); // Check if the previous device failed a build if ((result != CL_SUCCESS) && (retval != CL_SUCCESS)) { @@ -457,11 +457,11 @@ void Program::StubProgramSource(const std::string& app_name) { program_counter++; } -cl_int Program::build(const std::vector& devices, const char* options, +int32_t Program::build(const std::vector& devices, const char* options, void(CL_CALLBACK* notifyFptr)(cl_program, void*), void* data, bool optionChangable) { ScopedLock sl(buildLock_); - cl_int retval = CL_SUCCESS; + int32_t retval = CL_SUCCESS; if (symbolTable_ == NULL) { symbolTable_ = new symbols_t; @@ -536,7 +536,7 @@ cl_int Program::build(const std::vector& devices, const char* options, if (devProgram->buildStatus() != CL_BUILD_NONE) { continue; } - cl_int result = devProgram->build(sourceCode_, options, &parsedOptions); + int32_t result = devProgram->build(sourceCode_, options, &parsedOptions); // Check if the previous device failed a build if ((result != CL_SUCCESS) && (retval != CL_SUCCESS)) { diff --git a/rocclr/platform/program.hpp b/rocclr/platform/program.hpp index 9a742df581..00a18c9838 100644 --- a/rocclr/platform/program.hpp +++ b/rocclr/platform/program.hpp @@ -166,7 +166,7 @@ class Program : public RuntimeObject { const std::string& programLog() const { return programLog_; } //! Add a new device program with or without binary image and options. - cl_int addDeviceProgram(Device&, const void* image = NULL, size_t len = 0, + int32_t addDeviceProgram(Device&, const void* image = NULL, size_t len = 0, amd::option::Options* options = NULL); //! Find the section for the given device. Return NULL if not found. @@ -182,20 +182,20 @@ class Program : public RuntimeObject { const std::string& kernelNames() const { return kernelNames_; } //! Compile the program for the given devices. - cl_int compile(const std::vector& devices, size_t numHeaders, + int32_t compile(const std::vector& devices, size_t numHeaders, const std::vector& headerPrograms, const char** headerIncludeNames, const char* options = NULL, void(CL_CALLBACK* notifyFptr)(cl_program, void*) = NULL, void* data = NULL, bool optionChangable = true); //! Link the programs for the given devices. - cl_int link(const std::vector& devices, size_t numInputs, + int32_t link(const std::vector& devices, size_t numInputs, const std::vector& inputPrograms, const char* options = NULL, void(CL_CALLBACK* notifyFptr)(cl_program, void*) = NULL, void* data = NULL, bool optionChangable = true); //! Build the program for the given devices. - cl_int build(const std::vector& devices, const char* options = NULL, + int32_t build(const std::vector& devices, const char* options = NULL, void(CL_CALLBACK* notifyFptr)(cl_program, void*) = NULL, void* data = NULL, bool optionChangable = true);