AI/rocm-systems
2
0
Tükrözés 0
Kód Hibajegyek Egyesítési kérések Actions Packages Projektek Kiadások Wiki Tevékenység

Replace cl_* integral types with standard types.

Forráskód böngészése 
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


[ROCm/clr commit: d9d9c69399]
This commit is contained in:
Laurent Morichetti
2020-02-12 13:16:06 -08:00
szülő aaf154b30b
commit b3297f189d
42 fájl változott, egészen pontosan 419 új sor hozzáadva és 419 régi sor törölve
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/blit.cpp
+8 -8
Fájl megtekintése
@@ -596,17 +596,17 @@ bool HostBlitManager::fillImage(device::Memory& memory, const void* pattern,
size_t devSlicePitch;
void* newpattern = const_cast<void*>(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<float*>(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<void*>(&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<void*>(&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
projects/clr/rocclr/device/blit.hpp
+1 -1
Fájl megtekintése
@@ -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
projects/clr/rocclr/device/device.cpp
+6 -6
Fájl megtekintése
@@ -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);
// }
projects/clr/rocclr/device/device.hpp
+86 -86
Fájl megtekintése
@@ -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;
}
projects/clr/rocclr/device/devkernel.hpp
+2 -2
Fájl megtekintése
@@ -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
projects/clr/rocclr/device/devprogram.cpp
+3 -3
Fájl megtekintése
@@ -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<const std::string*>& 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<Program*>& inputPrograms, const char* origLinkOptions,
int32_t Program::link(const std::vector<Program*>& inputPrograms, const char* origLinkOptions,
amd::option::Options* linkOptions) {
lastBuildOptionsArg_ = origLinkOptions ? origLinkOptions : "";
if (linkOptions) {
@@ -1643,7 +1643,7 @@ cl_int Program::link(const std::vector<Program*>& 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) {
projects/clr/rocclr/device/devprogram.hpp
+6 -6
Fájl megtekintése
@@ -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<const std::string*>& headers,
int32_t compile(const std::string& sourceCode, const std::vector<const std::string*>& headers,
const char** headerIncludeNames, const char* origOptions,
amd::option::Options* options);
//! Builds the device program.
cl_int link(const std::vector<Program*>& inputPrograms, const char* origOptions,
int32_t link(const std::vector<Program*>& 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_; }
projects/clr/rocclr/device/gpu/gpublit.cpp
+36 -36
Fájl megtekintése
@@ -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), &copySize);
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<void*>(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<void*>(&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);
projects/clr/rocclr/device/gpu/gpucounters.hpp
+3 -3
Fájl megtekintése
@@ -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),
projects/clr/rocclr/device/gpu/gpudebugmanager.cpp
+3 -3
Fájl megtekintése
@@ -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;
projects/clr/rocclr/device/gpu/gpudebugmanager.hpp
+3 -3
Fájl megtekintése
@@ -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_; }
projects/clr/rocclr/device/gpu/gpudevice.cpp
+17 -17
Fájl megtekintése
@@ -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<cl_ulong>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
(static_cast<uint64_t>(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<cl_ulong>(memInfo.cardMemAvailableBytes + memInfo.cardExtMemAvailableBytes) /
static_cast<uint64_t>(memInfo.cardMemAvailableBytes + memInfo.cardExtMemAvailableBytes) /
100u);
#else
info_.globalMemSize_ = (static_cast<cl_ulong>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
static_cast<cl_ulong>(calAttr.localRAM) / 100u) *
info_.globalMemSize_ = (static_cast<uint64_t>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
static_cast<uint64_t>(calAttr.localRAM) / 100u) *
Mi;
#endif
int uswcPercentAvailable = (calAttr.uncachedRemoteRAM > 1536 && IS_WINDOWS) ? 75 : 50;
if (settings().apuSystem_) {
info_.globalMemSize_ +=
(static_cast<cl_ulong>(calAttr.uncachedRemoteRAM) * Mi * uswcPercentAvailable) / 100;
(static_cast<uint64_t>(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<cl_ulong>(calAttr.uncachedRemoteRAM) * Mi * uswcPercentAvailable) / 100,
(static_cast<uint64_t>(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<cl_ulong>(settings().maxAllocSize_));
std::min(info_.maxMemAllocSize_, static_cast<uint64_t>(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<const Memory*>& 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_;
projects/clr/rocclr/device/gpu/gpudevice.hpp
+1 -1
Fájl megtekintése
@@ -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
projects/clr/rocclr/device/gpu/gpukernel.cpp
+4 -4
Fájl megtekintése
@@ -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_;
projects/clr/rocclr/device/gpu/gpukernel.hpp
+2 -2
Fájl megtekintése
@@ -425,7 +425,7 @@ class NullKernel : public device::Kernel {
const NullProgram& nullProg() const { return reinterpret_cast<const NullProgram&>(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
projects/clr/rocclr/device/gpu/gpuprogram.cpp
+2 -2
Fájl megtekintése
@@ -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<device::Program*>& 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());
projects/clr/rocclr/device/gpu/gpuvirtual.cpp
+4 -4
Fájl megtekintése
@@ -1424,7 +1424,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& vcmd) {
std::vector<void*>& 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<amd::Memory*> 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];
}
projects/clr/rocclr/device/hwdebug.hpp
+2 -2
Fájl megtekintése
@@ -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;
projects/clr/rocclr/device/pal/palblit.cpp
+36 -36
Fájl megtekintése
@@ -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), &copySize);
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<void*>(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<void*>(&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;
projects/clr/rocclr/device/pal/palconstbuf.cpp
+1 -1
Fájl megtekintése
@@ -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);
projects/clr/rocclr/device/pal/palcounters.hpp
+3 -3
Fájl megtekintése
@@ -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;
projects/clr/rocclr/device/pal/paldebugmanager.cpp
+3 -3
Fájl megtekintése
@@ -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;
projects/clr/rocclr/device/pal/paldebugmanager.hpp
+3 -3
Fájl megtekintése
@@ -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<const pal::Device*>(device_); }
projects/clr/rocclr/device/pal/paldevice.cpp
+17 -17
Fájl megtekintése
@@ -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<cl_ulong>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
static_cast<cl_ulong>(localRAM) / 100u);
info_.globalMemSize_ = (static_cast<uint64_t>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
static_cast<uint64_t>(localRAM) / 100u);
uint uswcPercentAvailable =
((static_cast<cl_ulong>(heaps[Pal::GpuHeapGartUswc].heapSize) / Mi) > 1536 && IS_WINDOWS)
((static_cast<uint64_t>(heaps[Pal::GpuHeapGartUswc].heapSize) / Mi) > 1536 && IS_WINDOWS)
? 75
: 50;
if (settings().apuSystem_) {
info_.globalMemSize_ +=
(static_cast<cl_ulong>(heaps[Pal::GpuHeapGartUswc].heapSize) * uswcPercentAvailable) / 100;
(static_cast<uint64_t>(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<cl_ulong>(heaps[Pal::GpuHeapGartUswc].heapSize) * uswcPercentAvailable) / 100,
(static_cast<uint64_t>(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<cl_ulong>(settings().maxAllocSize_));
std::min(info_.maxMemAllocSize_, static_cast<uint64_t>(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_;
projects/clr/rocclr/device/pal/paldevice.hpp
+1 -1
Fájl megtekintése
@@ -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_; }
projects/clr/rocclr/device/pal/palvirtual.cpp
+1 -1
Fájl megtekintése
@@ -2075,7 +2075,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& vcmd) {
std::vector<void*>& 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 {
projects/clr/rocclr/device/rocm/rocblit.cpp
+40 -40
Fájl megtekintése
@@ -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<amd::Image*>(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<amd::Memory>(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<amd::Memory>(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<amd::Memory>(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), &copySize);
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<void*>(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<float*>(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<void*>(&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<void*>(&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<amd::Memory>(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);
projects/clr/rocclr/device/rocm/roccounters.cpp
+3 -3
Fájl megtekintése
@@ -419,9 +419,9 @@ static const std::array<std::pair<hsa_ven_amd_aqlprofile_block_name_t, int>, 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) {
projects/clr/rocclr/device/rocm/roccounters.hpp
+3 -3
Fájl megtekintése
@@ -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();
projects/clr/rocclr/device/rocm/rocdevice.cpp
+8 -8
Fájl megtekintése
@@ -1129,13 +1129,13 @@ bool Device::populateOCLDeviceConstants() {
}
assert(global_segment_size > 0);
info_.globalMemSize_ = static_cast<cl_ulong>(global_segment_size);
info_.globalMemSize_ = static_cast<uint64_t>(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<cl_ulong>(gpuvm_segment_max_alloc_);
info_.maxMemAllocSize_ = static_cast<uint64_t>(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;
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+1 -1
Fájl megtekintése
@@ -1137,7 +1137,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) {
const std::vector<void*>& 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 {
projects/clr/rocclr/include/vdi_agent_amd.h
+22 -22
Fájl megtekintése
@@ -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 */);
projects/clr/rocclr/platform/agent.cpp
+16 -16
Fájl megtekintése
@@ -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()) {
projects/clr/rocclr/platform/agent.hpp
+7 -7
Fájl megtekintése
@@ -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) {
projects/clr/rocclr/platform/command.cpp
+10 -10
Fájl megtekintése
@@ -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<cl_event>(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;
}
projects/clr/rocclr/platform/command.hpp
+28 -28
Fájl megtekintése
@@ -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<CallBackFunction> 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<CallBackEntry*> 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<amd::Memory*>& 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<amd::Memory*>& 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<amd::Memory*>& 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<amd::Memory*>& memoryObjects,
uint32_t num_objects, const std::vector<amd::Memory*>& 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<amd::Memory*>& memoryObjects,
uint32_t num_objects, const std::vector<amd::Memory*>& 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<amd::Memory*>& 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<void*> 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
projects/clr/rocclr/platform/kernel.cpp
+1 -1
Fájl megtekintése
@@ -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
projects/clr/rocclr/platform/kernel.hpp
+1 -1
Fájl megtekintése
@@ -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;
projects/clr/rocclr/platform/memory.cpp
+8 -8
Fájl megtekintése
@@ -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<amd::Device*>& 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<amd::Device*>& 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
projects/clr/rocclr/platform/memory.hpp
+3 -3
Fájl megtekintése
@@ -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.
projects/clr/rocclr/platform/program.cpp
+10 -10
Fájl megtekintése
@@ -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<Device*>& devices, size_t numHeaders,
int32_t Program::compile(const std::vector<Device*>& devices, size_t numHeaders,
const std::vector<const Program*>& 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<Device*>& 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<Device*>& devices, size_t numHeaders,
return retval;
}
cl_int Program::link(const std::vector<Device*>& devices, size_t numInputs,
int32_t Program::link(const std::vector<Device*>& devices, size_t numInputs,
const std::vector<Program*>& 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<Device*>& 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<Device*>& devices, const char* options,
int32_t Program::build(const std::vector<Device*>& 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<Device*>& 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)) {
projects/clr/rocclr/platform/program.hpp
+4 -4
Fájl megtekintése
@@ -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<Device*>& devices, size_t numHeaders,
int32_t compile(const std::vector<Device*>& devices, size_t numHeaders,
const std::vector<const Program*>& 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<Device*>& devices, size_t numInputs,
int32_t link(const std::vector<Device*>& devices, size_t numInputs,
const std::vector<Program*>& 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<Device*>& devices, const char* options = NULL,
int32_t build(const std::vector<Device*>& devices, const char* options = NULL,
void(CL_CALLBACK* notifyFptr)(cl_program, void*) = NULL, void* data = NULL,
bool optionChangable = true);