Danylo Lytovchenko
421 řádky
15 KiB
C++
421 řádky
15 KiB
C++
/* Copyright (c) 2010 - 2021 Advanced Micro Devices, Inc.
|
|
|
|
Permission is hereby granted, free of charge, to any person obtaining a copy
|
|
of this software and associated documentation files (the "Software"), to deal
|
|
in the Software without restriction, including without limitation the rights
|
|
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
|
copies of the Software, and to permit persons to whom the Software is
|
|
furnished to do so, subject to the following conditions:
|
|
|
|
The above copyright notice and this permission notice shall be included in
|
|
all copies or substantial portions of the Software.
|
|
|
|
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
|
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
|
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
|
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
|
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
|
THE SOFTWARE. */
|
|
|
|
#include "OCLPerfDoubleDMA.h"
|
|
|
|
#include <Timer.h>
|
|
#include <assert.h>
|
|
#include <stdio.h>
|
|
|
|
#include <cmath>
|
|
#include <sstream>
|
|
#include <string>
|
|
|
|
#include "CL/cl.h"
|
|
#include "CL/cl_ext.h"
|
|
|
|
const size_t blockX = 256;
|
|
const size_t blockY = 256;
|
|
const size_t blockZ = 512;
|
|
|
|
const size_t chunk = 16;
|
|
const size_t size_S = blockX * blockY * blockZ * sizeof(cl_float4);
|
|
const size_t size_s = blockX * blockY * chunk * sizeof(cl_float4);
|
|
static const int WindowWidth = 80;
|
|
|
|
const size_t MaxQueues = 3;
|
|
bool profEnable = false;
|
|
|
|
static const char* strKernel =
|
|
"__kernel void dummy(__global float4* out) \n"
|
|
"{ \n"
|
|
" uint id = get_global_id(0); \n"
|
|
" float4 value = (float4)(1.0f, 2.0f, 3.0f, 4.0f); \n"
|
|
" uint factorial = 1; \n"
|
|
" for (uint i = 1; i < (id / 0x400); ++i)\n"
|
|
" { \n"
|
|
" factorial *= i; \n"
|
|
" } \n"
|
|
" out[id] = value * factorial; \n"
|
|
"} \n";
|
|
|
|
class ProfileQueue {
|
|
public:
|
|
enum Operation { Write = 0, Execute, Read, Total };
|
|
|
|
static const char* OperationName[Total];
|
|
static const char StartCommand[Total];
|
|
static const char ExecCommand[Total];
|
|
|
|
ProfileQueue() {}
|
|
~ProfileQueue() {
|
|
for (size_t op = 0; op < Total; ++op) {
|
|
for (size_t idx = 0; idx < events_[op].size(); ++idx) {
|
|
clReleaseEvent(events_[op][idx]);
|
|
}
|
|
}
|
|
}
|
|
|
|
void addEvent(Operation op, cl_event event) { events_[op].push_back(event); }
|
|
|
|
void findMinMax(cl_long* min, cl_long* max) {
|
|
// Find time min/max ranges for the frame scaling
|
|
for (size_t op = 0; (op < ProfileQueue::Total); ++op) {
|
|
cl_long time;
|
|
if (events_[op].size() == 0) continue;
|
|
clGetEventProfilingInfo(events_[op][0], CL_PROFILING_COMMAND_START, sizeof(cl_long), &time,
|
|
NULL);
|
|
if (0 == *min) {
|
|
*min = time;
|
|
} else {
|
|
*min = std::min(*min, time);
|
|
}
|
|
clGetEventProfilingInfo(events_[op][events_[op].size() - 1], CL_PROFILING_COMMAND_END,
|
|
sizeof(cl_long), &time, NULL);
|
|
if (0 == *max) {
|
|
*max = time;
|
|
} else {
|
|
*max = std::max(*max, time);
|
|
}
|
|
}
|
|
}
|
|
|
|
void display(cl_long start, cl_long finish) {
|
|
std::string graph;
|
|
graph.resize(WindowWidth + 1);
|
|
graph[WindowWidth] = '\x0';
|
|
cl_long timeFrame = finish - start;
|
|
cl_long interval = timeFrame / WindowWidth;
|
|
|
|
// Find time min/max ranges for the frame scaling
|
|
for (size_t op = 0; (op < Total); ++op) {
|
|
if (events_[op].size() == 0) continue;
|
|
cl_long timeStart, timeEnd;
|
|
int begin = 0, end = 0;
|
|
for (size_t idx = 0; idx < events_[op].size(); ++idx) {
|
|
bool cutStart = false;
|
|
clGetEventProfilingInfo(events_[op][idx], CL_PROFILING_COMMAND_START, sizeof(cl_long),
|
|
&timeStart, NULL);
|
|
clGetEventProfilingInfo(events_[op][idx], CL_PROFILING_COMMAND_END, sizeof(cl_long),
|
|
&timeEnd, NULL);
|
|
|
|
// Continue if out of the frame scope
|
|
if (timeStart >= finish) continue;
|
|
if (timeEnd <= start) continue;
|
|
|
|
if (timeStart <= start) {
|
|
timeStart = start;
|
|
cutStart = true;
|
|
}
|
|
|
|
if (timeEnd >= finish) {
|
|
timeEnd = finish;
|
|
}
|
|
|
|
// Readjust time to the frame
|
|
timeStart -= start;
|
|
timeEnd -= start;
|
|
timeStart = static_cast<cl_long>(floor(static_cast<float>(timeStart) / interval + 0.5f));
|
|
timeEnd = static_cast<cl_long>(floor(static_cast<float>(timeEnd) / interval + 0.5f));
|
|
begin = static_cast<int>(timeStart);
|
|
// Idle from end to begin
|
|
for (int c = end; c < begin; ++c) {
|
|
graph[c] = '-';
|
|
}
|
|
end = static_cast<int>(timeEnd);
|
|
for (int c = begin; c < end; ++c) {
|
|
if ((c == begin) && !cutStart) {
|
|
graph[c] = StartCommand[op];
|
|
} else {
|
|
graph[c] = ExecCommand[op];
|
|
}
|
|
}
|
|
if ((begin == end) && (end < WindowWidth)) {
|
|
graph[begin] = '+';
|
|
}
|
|
}
|
|
if (end < WindowWidth) {
|
|
for (int c = end; c < WindowWidth; ++c) {
|
|
graph[c] = '-';
|
|
}
|
|
}
|
|
printf("%s\n", graph.c_str());
|
|
}
|
|
}
|
|
|
|
private:
|
|
// Profiling events
|
|
std::vector<cl_event> events_[Total];
|
|
};
|
|
|
|
const char* ProfileQueue::OperationName[Total] = {"BufferWrite", "KernelExecution", "BufferRead"};
|
|
const char ProfileQueue::StartCommand[Total] = {'W', 'X', 'R'};
|
|
const char ProfileQueue::ExecCommand[Total] = {'>', '#', '<'};
|
|
|
|
class Profile {
|
|
public:
|
|
Profile(bool profEna, int numQueues)
|
|
: profileEna_(profEna), numQueues_(numQueues), min_(0), max_(0), execTime_(0) {}
|
|
|
|
~Profile() {}
|
|
|
|
void addEvent(int queue, ProfileQueue::Operation op, cl_event event) {
|
|
if (profileEna_) {
|
|
profQueue[queue].addEvent(op, event);
|
|
}
|
|
}
|
|
|
|
cl_long findExecTime() {
|
|
if (execTime_ != 0) return execTime_;
|
|
for (int q = 0; q < numQueues_; ++q) {
|
|
profQueue[q].findMinMax(&min_, &max_);
|
|
}
|
|
execTime_ = max_ - min_;
|
|
return execTime_;
|
|
}
|
|
|
|
void display(cl_long start, cl_long finish) {
|
|
if (!profileEna_) return;
|
|
printf("\n ----------- Time frame %.3f (us), scale 1:%.0f\n", (float)(finish - start) / 1000,
|
|
(float)(finish - start) / (1000 * WindowWidth));
|
|
for (size_t op = 0; (op < ProfileQueue::Total); ++op) {
|
|
printf("%s - %c%c; ", ProfileQueue::OperationName[op], ProfileQueue::StartCommand[op],
|
|
ProfileQueue::ExecCommand[op]);
|
|
}
|
|
printf("\n");
|
|
for (int q = 0; q < numQueues_; ++q) {
|
|
printf("CommandQueue #%d\n", q);
|
|
profQueue[q].display(min_ + start, min_ + finish);
|
|
}
|
|
}
|
|
|
|
private:
|
|
bool profileEna_;
|
|
int numQueues_; //!< Total number of queues
|
|
cl_long min_; //!< Min HW timestamp
|
|
cl_long max_; //!< Max HW timestamp
|
|
cl_long execTime_; //!< Profile time
|
|
ProfileQueue profQueue[MaxQueues];
|
|
};
|
|
|
|
OCLPerfDoubleDMA::OCLPerfDoubleDMA() {
|
|
_numSubTests = 2 * MaxQueues * 2;
|
|
failed_ = false;
|
|
}
|
|
|
|
OCLPerfDoubleDMA::~OCLPerfDoubleDMA() {}
|
|
|
|
void OCLPerfDoubleDMA::open(unsigned int test, char* units, double& conversion,
|
|
unsigned int deviceId) {
|
|
_deviceId = deviceId;
|
|
OCLTestImp::open(test, units, conversion, deviceId);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "Error opening test");
|
|
test_ = test;
|
|
cl_device_type deviceType;
|
|
error_ = _wrapper->clGetDeviceInfo(devices_[deviceId], CL_DEVICE_TYPE, sizeof(deviceType),
|
|
&deviceType, NULL);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "CL_DEVICE_TYPE failed");
|
|
|
|
if (!(deviceType & CL_DEVICE_TYPE_GPU)) {
|
|
printf("GPU device is required for this test!\n");
|
|
failed_ = true;
|
|
return;
|
|
}
|
|
program_ = _wrapper->clCreateProgramWithSource(context_, 1, &strKernel, NULL, &error_);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "clCreateProgramWithSource() failed");
|
|
error_ = _wrapper->clBuildProgram(program_, 1, &devices_[deviceId], NULL, NULL, NULL);
|
|
if (error_ != CL_SUCCESS) {
|
|
char programLog[1024];
|
|
_wrapper->clGetProgramBuildInfo(program_, devices_[deviceId], CL_PROGRAM_BUILD_LOG, 1024,
|
|
programLog, 0);
|
|
printf("\n%s\n", programLog);
|
|
fflush(stdout);
|
|
}
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "clBuildProgram() failed");
|
|
kernel_ = _wrapper->clCreateKernel(program_, "dummy", &error_);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "clCreateKernel() failed");
|
|
|
|
size_t bufSize = size_s;
|
|
cl_mem buffer;
|
|
if (test_ >= (2 * MaxQueues)) {
|
|
profEnable = true;
|
|
}
|
|
test_ %= 2 * MaxQueues;
|
|
size_t numBufs = (test_ % MaxQueues) + 1;
|
|
for (size_t b = 0; b < numBufs; ++b) {
|
|
buffer = _wrapper->clCreateBuffer(context_, CL_MEM_READ_WRITE, bufSize, NULL, &error_);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "clCreateBuffer() failed");
|
|
buffers_.push_back(buffer);
|
|
}
|
|
|
|
buffer = _wrapper->clCreateBuffer(context_, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, size_S,
|
|
NULL, &error_);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "clCreateBuffer() failed");
|
|
buffers_.push_back(buffer);
|
|
}
|
|
|
|
static void CL_CALLBACK notify_callback(const char* errinfo, const void* private_info, size_t cb,
|
|
void* user_data) {}
|
|
|
|
void OCLPerfDoubleDMA::run(void) {
|
|
if (failed_) {
|
|
return;
|
|
}
|
|
CPerfCounter timer;
|
|
const int numQueues = (test_ % MaxQueues) + 1;
|
|
const bool useKernel = ((test_ / MaxQueues) > 0);
|
|
const int numBufs = numQueues;
|
|
Profile profile(profEnable, numQueues);
|
|
|
|
std::vector<cl_command_queue> cmdQueues(numQueues);
|
|
int q;
|
|
cl_command_queue_properties qProp = (profEnable) ? CL_QUEUE_PROFILING_ENABLE : 0;
|
|
for (q = 0; q < numQueues; ++q) {
|
|
cl_command_queue cmdQueue =
|
|
_wrapper->clCreateCommandQueue(context_, devices_[_deviceId], qProp, &error_);
|
|
CHECK_RESULT((error_ != CL_SUCCESS), "clCreateCommandQueue() failed");
|
|
cmdQueues[q] = cmdQueue;
|
|
}
|
|
|
|
float* Data_s = (float*)_wrapper->clEnqueueMapBuffer(cmdQueues[0], buffers_[numBufs], CL_TRUE,
|
|
CL_MAP_READ | CL_MAP_WRITE, 0, size_S, 0,
|
|
NULL, NULL, &error_);
|
|
|
|
size_t gws[1] = {size_s / (4 * sizeof(float))};
|
|
size_t lws[1] = {256};
|
|
|
|
// Warm-up
|
|
for (q = 0; q < numQueues; ++q) {
|
|
error_ |= _wrapper->clEnqueueWriteBuffer(cmdQueues[q], buffers_[q], CL_FALSE, 0, size_s,
|
|
(char*)Data_s, 0, NULL, NULL);
|
|
error_ |= _wrapper->clSetKernelArg(kernel_, 0, sizeof(cl_mem), (void*)&buffers_[q]);
|
|
error_ |=
|
|
_wrapper->clEnqueueNDRangeKernel(cmdQueues[q], kernel_, 1, NULL, gws, lws, 0, NULL, NULL);
|
|
error_ |= _wrapper->clEnqueueReadBuffer(cmdQueues[q], buffers_[q], CL_FALSE, 0, size_s,
|
|
(char*)Data_s, 0, NULL, NULL);
|
|
error_ |= _wrapper->clFinish(cmdQueues[q]);
|
|
}
|
|
|
|
CHECK_RESULT_NO_RETURN((error_ != CL_SUCCESS), "Execution failed");
|
|
|
|
size_t s_done = 0;
|
|
cl_event r[MaxQueues] = {0}, w[MaxQueues] = {0}, x[MaxQueues] = {0};
|
|
|
|
/*---------- pass2: copy Data_s to and from GPU Buffers ----------*/
|
|
s_done = 0;
|
|
timer.Reset();
|
|
timer.Start();
|
|
int idx = numBufs - 1;
|
|
// Start from the last so read/write won't go to the same DMA when kernel is
|
|
// executed
|
|
q = numQueues - 1;
|
|
size_t iter = 0;
|
|
while (1) {
|
|
if (0 == r[idx]) {
|
|
error_ |= _wrapper->clEnqueueWriteBuffer(cmdQueues[q], buffers_[idx], CL_FALSE, 0, size_s,
|
|
(char*)Data_s + s_done, 0, NULL, &w[idx]);
|
|
} else {
|
|
error_ |= _wrapper->clEnqueueWriteBuffer(cmdQueues[q], buffers_[idx], CL_FALSE, 0, size_s,
|
|
(char*)Data_s + s_done, 1, &r[idx], &w[idx]);
|
|
if (!profEnable) {
|
|
error_ |= _wrapper->clReleaseEvent(r[idx]);
|
|
}
|
|
}
|
|
_wrapper->clFlush(cmdQueues[q]);
|
|
profile.addEvent(q, ProfileQueue::Write, w[idx]);
|
|
|
|
if (useKernel) {
|
|
// Change the queue
|
|
++q %= numQueues;
|
|
// Implicit flush of DMA engine on kernel start, because memory dependency
|
|
error_ |= _wrapper->clSetKernelArg(kernel_, 0, sizeof(cl_mem), (void*)&buffers_[idx]);
|
|
error_ |= _wrapper->clEnqueueNDRangeKernel(cmdQueues[q], kernel_, 1, NULL, gws, lws, 1,
|
|
&w[idx], &x[idx]);
|
|
if (!profEnable) {
|
|
error_ |= _wrapper->clReleaseEvent(w[idx]);
|
|
}
|
|
profile.addEvent(q, ProfileQueue::Execute, x[idx]);
|
|
}
|
|
_wrapper->clFlush(cmdQueues[q]);
|
|
|
|
// Change the queue
|
|
++q %= numQueues;
|
|
error_ |= _wrapper->clEnqueueReadBuffer(cmdQueues[q], buffers_[idx], CL_FALSE, 0, size_s,
|
|
(char*)Data_s + s_done, 1,
|
|
(useKernel) ? &x[idx] : &w[idx], &r[idx]);
|
|
if (!profEnable) {
|
|
error_ |= _wrapper->clReleaseEvent((useKernel) ? x[idx] : w[idx]);
|
|
}
|
|
profile.addEvent(q, ProfileQueue::Read, r[idx]);
|
|
_wrapper->clFlush(cmdQueues[q]);
|
|
|
|
if ((s_done += size_s) >= size_S) {
|
|
if (!profEnable) {
|
|
error_ |= _wrapper->clReleaseEvent(r[idx]);
|
|
}
|
|
break;
|
|
}
|
|
++iter;
|
|
++idx %= numBufs;
|
|
++q %= numQueues;
|
|
}
|
|
|
|
for (q = 0; q < numQueues; ++q) {
|
|
error_ |= _wrapper->clFinish(cmdQueues[q]);
|
|
}
|
|
timer.Stop();
|
|
|
|
error_ =
|
|
_wrapper->clEnqueueUnmapMemObject(cmdQueues[0], buffers_[numBufs], Data_s, 0, NULL, NULL);
|
|
|
|
error_ |= _wrapper->clFinish(cmdQueues[0]);
|
|
CHECK_RESULT_NO_RETURN((error_ != CL_SUCCESS), "Execution failed");
|
|
|
|
cl_long gpuTimeFrame = profile.findExecTime();
|
|
cl_long oneIter = gpuTimeFrame / iter;
|
|
|
|
// Display 4 iterations in the middle
|
|
cl_long startFrame = oneIter * (iter / 2 - 2);
|
|
cl_long finishFrame = oneIter * (iter / 2 + 2);
|
|
profile.display(startFrame, finishFrame);
|
|
|
|
for (q = 0; q < numQueues; ++q) {
|
|
error_ = _wrapper->clReleaseCommandQueue(cmdQueues[q]);
|
|
CHECK_RESULT_NO_RETURN((error_ != CL_SUCCESS), "clReleaseCommandQueue() failed");
|
|
}
|
|
|
|
double GBytes = (double)(2 * size_S) / (double)(1000 * 1000 * 1000);
|
|
_perfInfo = static_cast<float>(GBytes / timer.GetElapsedTime());
|
|
|
|
std::stringstream stream;
|
|
if (useKernel) {
|
|
stream << "Write/Kernel/Read operation ";
|
|
} else {
|
|
stream << "Write/Read operation ";
|
|
}
|
|
stream << numQueues << " queues; profiling " << ((profEnable) ? "enabled" : "disabled")
|
|
<< " [GB/s]";
|
|
|
|
stream.flags(std::ios::right | std::ios::showbase);
|
|
testDescString = stream.str();
|
|
}
|
|
|
|
unsigned int OCLPerfDoubleDMA::close(void) { return OCLTestImp::close(); }
|