Skip to content

Commit 88f22c5

Browse files
bashbaugbashbaug
committed
add a tester to determine if SVM allocs and frees are blocking
1 parent 47c50ac commit 88f22c5

File tree

3 files changed

+176
-0
lines changed

3 files changed

+176
-0
lines changed

‎samples/99_svmfree/CMakeLists.txt‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
# Copyright (c) 2025 Ben Ashbaugh
2+
#
3+
# SPDX-License-Identifier: MIT
4+
5+
add_opencl_sample(
6+
TEST
7+
NUMBER 99
8+
TARGET svmfree
9+
VERSION 200
10+
SOURCES main.cpp)

‎samples/99_svmfree/main.cpp‎

Lines changed: 164 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,164 @@
1+
/*
2+
// Copyright (c) 2025 Ben Ashbaugh
3+
//
4+
// SPDX-License-Identifier: MIT
5+
*/
6+
7+
#include <popl/popl.hpp>
8+
9+
#include <CL/opencl.hpp>
10+
11+
#include <chrono>
12+
13+
using test_clock = std::chrono::steady_clock;
14+
15+
static const char kernelString[] = R"CLC(
16+
kernel void eat_time(global int* ptr, int kernelOperationsCount )
17+
{
18+
volatile int value = kernelOperationsCount;
19+
while (--value) {}
20+
if (get_global_id(0) > 10000000) {
21+
ptr[0] = 0;
22+
}
23+
}
24+
)CLC";
25+
26+
int kernelOperationsCount = 10000000;
27+
28+
static void test_malloc(cl::Context& context, cl::CommandQueue& queue, cl::Kernel& kernel, size_t count)
29+
{
30+
std::cout << "Testing SVM alloc for " << count * sizeof(int) << " bytes:\n";
31+
32+
auto ptr = (int*)clSVMAlloc(
33+
context(),
34+
CL_MEM_READ_WRITE,
35+
count * sizeof(int),
36+
0);
37+
38+
kernel.setArg(0, ptr);
39+
kernel.setArg(1, kernelOperationsCount);
40+
41+
auto start = test_clock::now();
42+
queue.enqueueNDRangeKernel(
43+
kernel,
44+
cl::NullRange,
45+
cl::NDRange{1024},
46+
cl::NDRange{1});
47+
auto submit = test_clock::now();
48+
auto another = (int*)clSVMAlloc(
49+
context(),
50+
CL_MEM_READ_WRITE,
51+
count * sizeof(int),
52+
0);
53+
auto alloc = test_clock::now();
54+
queue.finish();
55+
auto wait = test_clock::now();
56+
57+
std::cout << "\tAlloc time: " << std::chrono::duration_cast<std::chrono::microseconds>(alloc - submit).count() << " us" << std::endl;
58+
std::cout << "\tWait time: " << std::chrono::duration_cast<std::chrono::microseconds>(wait - alloc).count() << " us" << std::endl;
59+
60+
clSVMFree(context(), ptr);
61+
clSVMFree(context(), another);
62+
}
63+
64+
static void test_free(cl::Context& context, cl::CommandQueue& queue, cl::Kernel& kernel, size_t count, bool inUse)
65+
{
66+
std::cout << "Testing SVM free while pointer is" << (inUse ? "" : " NOT") << " in use for " << count * sizeof(int) << " bytes:\n";
67+
68+
auto ptr = (int*)clSVMAlloc(
69+
context(),
70+
CL_MEM_READ_WRITE,
71+
count * sizeof(int),
72+
0);
73+
auto another = (int*)clSVMAlloc(
74+
context(),
75+
CL_MEM_READ_WRITE,
76+
count * sizeof(int),
77+
0);
78+
79+
int kernelOperationsCount = 10000000;
80+
81+
kernel.setArg(0, inUse ? ptr : another);
82+
kernel.setArg(1, kernelOperationsCount);
83+
84+
auto start = test_clock::now();
85+
queue.enqueueNDRangeKernel(
86+
kernel,
87+
cl::NullRange,
88+
cl::NDRange{1024},
89+
cl::NDRange{1});
90+
auto submit = test_clock::now();
91+
clSVMFree(context(), ptr);
92+
auto free = test_clock::now();
93+
queue.finish();
94+
auto wait = test_clock::now();
95+
96+
std::cout << "\tFree time: " << std::chrono::duration_cast<std::chrono::microseconds>(free - submit).count() << " us" << std::endl;
97+
std::cout << "\tWait time: " << std::chrono::duration_cast<std::chrono::microseconds>(wait - free).count() << " us" << std::endl;
98+
99+
clSVMFree(context(), another);
100+
}
101+
102+
int main(int argc, char** argv)
103+
{
104+
int platformIndex = 0;
105+
int deviceIndex = 0;
106+
107+
size_t smallCount = 1024;
108+
size_t largeCount = 32 * 1024 * 1024;
109+
110+
{
111+
popl::OptionParser op("Supported Options");
112+
op.add<popl::Value<int>>("p", "platform", "Platform Index", platformIndex, &platformIndex);
113+
op.add<popl::Value<int>>("d", "device", "Device Index", deviceIndex, &deviceIndex);
114+
op.add<popl::Value<int>>("t", "timer", "Timer Value", kernelOperationsCount, &kernelOperationsCount);
115+
op.add<popl::Value<size_t>>("s", "small", "Small Count", smallCount, &smallCount);
116+
op.add<popl::Value<size_t>>("l", "large", "Large Count", largeCount, &largeCount);
117+
118+
bool printUsage = false;
119+
try {
120+
op.parse(argc, argv);
121+
} catch (std::exception& e) {
122+
fprintf(stderr, "Error: %s\n\n", e.what());
123+
printUsage = true;
124+
}
125+
126+
if (printUsage || !op.unknown_options().empty() || !op.non_option_args().empty()) {
127+
fprintf(stderr,
128+
"Usage: svmfree [options]\n"
129+
"%s", op.help().c_str());
130+
return -1;
131+
}
132+
}
133+
134+
std::vector<cl::Platform> platforms;
135+
cl::Platform::get(&platforms);
136+
137+
printf("Running on platform: %s\n",
138+
platforms[platformIndex].getInfo<CL_PLATFORM_NAME>().c_str() );
139+
140+
std::vector<cl::Device> devices;
141+
platforms[platformIndex].getDevices(CL_DEVICE_TYPE_ALL, &devices);
142+
143+
printf("Running on device: %s\n",
144+
devices[deviceIndex].getInfo<CL_DEVICE_NAME>().c_str() );
145+
146+
cl::Context context{devices[deviceIndex]};
147+
cl::CommandQueue commandQueue{context, devices[deviceIndex]};
148+
149+
cl::Program program{ context, kernelString };
150+
program.build();
151+
cl::Kernel kernel = cl::Kernel{ program, "eat_time" };
152+
153+
std::cout << "Testing small SVM allocations...\n";
154+
test_malloc(context, commandQueue, kernel, smallCount);
155+
test_free(context, commandQueue, kernel, smallCount, false);
156+
test_free(context, commandQueue, kernel, smallCount, true);
157+
158+
std::cout << "Testing large SVM allocations...\n";
159+
test_malloc(context, commandQueue, kernel, largeCount);
160+
test_free(context, commandQueue, kernel, largeCount, false);
161+
test_free(context, commandQueue, kernel, largeCount, true);
162+
163+
return 0;
164+
}

‎samples/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -91,3 +91,5 @@ if(BUILD_EXTENSION_SAMPLES)
9191
add_subdirectory( 14_ooqcommandbuffers )
9292
add_subdirectory( 15_mutablecommandbufferasserts )
9393
endif()
94+
95+
add_subdirectory( 99_svmfree )

0 commit comments

Comments
 (0)