Skip to content

Commit c43356c

Browse files
callumfarecallumfare
committed
WIP program and kernel launch testing
1 parent 3b3af0e commit c43356c

18 files changed

+498
-2
lines changed

‎offload/liboffload/src/OffloadImpl.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -530,7 +530,7 @@ ol_impl_result_t olCreateProgram_impl(ol_device_handle_t Device, void *ProgData,
530530
StringRef(reinterpret_cast<char *>(ProgData), ProgDataSize));
531531
__tgt_device_image DeviceImage{
532532
const_cast<char *>(ImageData->getBuffer().data()),
533-
const_cast<char *>(ImageData->getBuffer().data()) + ProgDataSize - 1,
533+
const_cast<char *>(ImageData->getBuffer().data()) + ProgDataSize,
534534
nullptr, nullptr};
535535

536536
ol_program_handle_t Prog = new ol_program_impl_t();

‎offload/unittests/OffloadAPI/CMakeLists.txt

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,9 @@
11
set(PLUGINS_TEST_COMMON LLVMOffload)
22
set(PLUGINS_TEST_INCLUDE ${LIBOMPTARGET_INCLUDE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/common)
33

4+
add_subdirectory(device_code)
5+
message(${OFFLOAD_TEST_DEVICE_CODE_PATH})
6+
47
add_libompt_unittest("offload.unittests"
58
${CMAKE_CURRENT_SOURCE_DIR}/common/Environment.cpp
69
${CMAKE_CURRENT_SOURCE_DIR}/platform/olGetPlatform.cpp
@@ -21,7 +24,17 @@ add_libompt_unittest("offload.unittests"
2124
${CMAKE_CURRENT_SOURCE_DIR}/enqueue/olEnqueueDataRead.cpp
2225
${CMAKE_CURRENT_SOURCE_DIR}/enqueue/olEnqueueDataCopy.cpp
2326
${CMAKE_CURRENT_SOURCE_DIR}/enqueue/olEnqueueMemcpy.cpp
27+
${CMAKE_CURRENT_SOURCE_DIR}/enqueue/olEnqueueKernelLaunch.cpp
28+
${CMAKE_CURRENT_SOURCE_DIR}/program/olCreateProgram.cpp
29+
${CMAKE_CURRENT_SOURCE_DIR}/program/olRetainProgram.cpp
30+
${CMAKE_CURRENT_SOURCE_DIR}/program/olReleaseProgram.cpp
31+
${CMAKE_CURRENT_SOURCE_DIR}/kernel/olCreateKernel.cpp
32+
${CMAKE_CURRENT_SOURCE_DIR}/kernel/olReleaseKernel.cpp
33+
${CMAKE_CURRENT_SOURCE_DIR}/kernel/olRetainKernel.cpp
34+
${CMAKE_CURRENT_SOURCE_DIR}/kernel/olSetKernelArgValue.cpp
35+
${CMAKE_CURRENT_SOURCE_DIR}/kernel/olSetKernelArgsData.cpp
2436
)
25-
add_dependencies("offload.unittests" ${PLUGINS_TEST_COMMON})
37+
add_dependencies("offload.unittests" ${PLUGINS_TEST_COMMON} LibompUnitTestsDeviceBins)
38+
target_compile_definitions("offload.unittests" PRIVATE DEVICE_CODE_PATH="${OFFLOAD_TEST_DEVICE_CODE_PATH}")
2639
target_link_libraries("offload.unittests" PRIVATE ${PLUGINS_TEST_COMMON})
2740
target_include_directories("offload.unittests" PRIVATE ${PLUGINS_TEST_INCLUDE})

‎offload/unittests/OffloadAPI/common/Environment.cpp

Lines changed: 51 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010
#include "Fixtures.hpp"
1111
#include "llvm/Support/CommandLine.h"
1212
#include <OffloadAPI.h>
13+
#include <fstream>
1314

1415
using namespace llvm;
1516

@@ -94,3 +95,53 @@ ol_platform_handle_t TestEnvironment::getPlatform() {
9495

9596
return Platform;
9697
}
98+
99+
// TODO: Define via cmake, also override via cmd line arg
100+
const std::string DeviceBinsDirectory = DEVICE_CODE_PATH;
101+
102+
bool TestEnvironment::loadDeviceBinary(
103+
const std::string &BinaryName, ol_platform_handle_t Platform,
104+
std::shared_ptr<std::vector<char>> &BinaryOut) {
105+
106+
// Get the platform type
107+
ol_platform_backend_t Backend = OL_PLATFORM_BACKEND_UNKNOWN;
108+
olGetPlatformInfo(Platform, OL_PLATFORM_INFO_BACKEND, sizeof(Backend),
109+
&Backend);
110+
std::string FileExtension;
111+
if (Backend == OL_PLATFORM_BACKEND_AMDGPU) {
112+
FileExtension = ".amdgpu.bin";
113+
} else if (Backend == OL_PLATFORM_BACKEND_CUDA) {
114+
FileExtension = ".nvptx64.bin";
115+
} else {
116+
errs() << "Unsupported platform type for a device binary test.\n";
117+
return false;
118+
}
119+
120+
std::string SourcePath =
121+
DeviceBinsDirectory + "/" + BinaryName + FileExtension;
122+
123+
std::ifstream SourceFile;
124+
SourceFile.open(SourcePath, std::ios::binary | std::ios::in | std::ios::ate);
125+
126+
if (!SourceFile.is_open()) {
127+
errs() << "failed opening device binary path: " + SourcePath;
128+
return false;
129+
}
130+
131+
size_t SourceSize = static_cast<size_t>(SourceFile.tellg());
132+
SourceFile.seekg(0, std::ios::beg);
133+
134+
std::vector<char> DeviceBinary(SourceSize);
135+
SourceFile.read(DeviceBinary.data(), SourceSize);
136+
if (!SourceFile) {
137+
SourceFile.close();
138+
errs() << "failed reading device binary data from file: " + SourcePath;
139+
return false;
140+
}
141+
SourceFile.close();
142+
143+
auto BinaryPtr = std::make_shared<std::vector<char>>(std::move(DeviceBinary));
144+
145+
BinaryOut = BinaryPtr;
146+
return true;
147+
}

‎offload/unittests/OffloadAPI/common/Environment.hpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,4 +14,7 @@
1414
namespace TestEnvironment {
1515
const std::vector<ol_platform_handle_t> &getPlatforms();
1616
ol_platform_handle_t getPlatform();
17+
bool loadDeviceBinary(const std::string &BinaryName,
18+
ol_platform_handle_t Platform,
19+
std::shared_ptr<std::vector<char>> &BinaryOut);
1720
} // namespace TestEnvironment

‎offload/unittests/OffloadAPI/common/Fixtures.hpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,14 @@
2727
} while (0)
2828
#endif
2929

30+
#ifndef ASSERT_ANY_ERROR
31+
#define ASSERT_ANY_ERROR(ACTUAL) \
32+
do { \
33+
ol_result_t Res = ACTUAL; \
34+
ASSERT_TRUE(Res); \
35+
} while (0)
36+
#endif
37+
3038
#define RETURN_ON_FATAL_FAILURE(...) \
3139
__VA_ARGS__; \
3240
if (this->HasFatalFailure() || this->IsSkipped()) { \
@@ -63,6 +71,44 @@ struct offloadDeviceTest : offloadPlatformTest {
6371
ol_device_handle_t Device = nullptr;
6472
};
6573

74+
// Fixture for a generic program test. If you want a different program, use
75+
// offloadQueueTest and create your own program handle with the binary you want.
76+
struct offloadProgramTest : offloadDeviceTest {
77+
void SetUp() override {
78+
RETURN_ON_FATAL_FAILURE(offloadDeviceTest::SetUp());
79+
ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Platform, DeviceBin));
80+
ASSERT_GE(DeviceBin->size(), 0lu);
81+
ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->data(), DeviceBin->size(),
82+
&Program));
83+
}
84+
85+
void TearDown() override {
86+
if (Program) {
87+
olReleaseProgram(Program);
88+
}
89+
RETURN_ON_FATAL_FAILURE(offloadDeviceTest::TearDown());
90+
}
91+
92+
ol_program_handle_t Program = nullptr;
93+
std::shared_ptr<std::vector<char>> DeviceBin;
94+
};
95+
96+
struct offloadKernelTest : offloadProgramTest {
97+
void SetUp() override {
98+
RETURN_ON_FATAL_FAILURE(offloadProgramTest::SetUp());
99+
ASSERT_SUCCESS(olCreateKernel(Program, "foo", &Kernel));
100+
}
101+
102+
void TearDown() override {
103+
if (Kernel) {
104+
olReleaseKernel(Kernel);
105+
}
106+
RETURN_ON_FATAL_FAILURE(offloadProgramTest::TearDown());
107+
}
108+
109+
ol_kernel_handle_t Kernel = nullptr;
110+
};
111+
66112
struct offloadQueueTest : offloadDeviceTest {
67113
void SetUp() override {
68114
RETURN_ON_FATAL_FAILURE(offloadDeviceTest::SetUp());

‎offload/unittests/OffloadAPI/device_code/CMakeLists.txt

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
macro(add_offload_test_device_code test_filename test_name)
2+
message("Building Offload API device code for test '${test_name}'")
3+
set(SRC_PATH ${CMAKE_CURRENT_SOURCE_DIR}/${test_filename})
4+
set(UTIL_PATH ${CMAKE_CURRENT_SOURCE_DIR}/util.h)
5+
6+
# Build for NVPTX
7+
set(BIN_PATH ${CMAKE_CURRENT_BINARY_DIR}/${test_name}.nvptx64.bin)
8+
add_custom_command(OUTPUT ${BIN_PATH}
9+
COMMAND
10+
${CMAKE_C_COMPILER} --target=nvptx64-nvidia-cuda -march=native
11+
--cuda-path=/usr/local/cuda
12+
${SRC_PATH} -o ${BIN_PATH}
13+
DEPENDS ${SRC_PATH} ${UTIL_PATH}
14+
)
15+
list(APPEND BIN_PATHS ${BIN_PATH})
16+
17+
# Build for AMDGPU
18+
set(BIN_PATH ${CMAKE_CURRENT_BINARY_DIR}/${test_name}.amdgpu.bin)
19+
add_custom_command(OUTPUT ${BIN_PATH}
20+
COMMAND
21+
${CMAKE_C_COMPILER} --target=amdgcn-amd-amdhsa -nogpulib
22+
${SRC_PATH} -o ${BIN_PATH}
23+
DEPENDS ${SRC_PATH} ${UTIL_PATH}
24+
)
25+
list(APPEND BIN_PATHS ${BIN_PATH})
26+
27+
# TODO: Build for host CPU
28+
endmacro()
29+
30+
31+
add_offload_test_device_code(foo.c foo)
32+
add_offload_test_device_code(bar.c bar)
33+
34+
add_custom_target(LibompUnitTestsDeviceBins DEPENDS ${BIN_PATHS})
35+
36+
set(OFFLOAD_TEST_DEVICE_CODE_PATH ${CMAKE_CURRENT_BINARY_DIR} PARENT_SCOPE)

‎offload/unittests/OffloadAPI/device_code/bar.c

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
#include "util.h"
2+
3+
KERNEL void bar(int *out) {
4+
out[get_thread_id_x()] = get_thread_id_x() + 1;
5+
}

‎offload/unittests/OffloadAPI/device_code/foo.c

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
#include "util.h"
2+
3+
KERNEL void foo(int *out) {
4+
out[get_thread_id_x()] = get_thread_id_x();
5+
}

‎offload/unittests/OffloadAPI/device_code/util.h

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
//===------- Offload API tests - helper for test device code --------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#if defined(__AMDGPU__)
10+
#define KERNEL [[clang::amdgpu_kernel]]
11+
#define get_thread_id_x() __builtin_amdgcn_workitem_id_x()
12+
#elif defined(__NVPTX__)
13+
#define KERNEL [[clang::nvptx_kernel]]
14+
#define get_thread_id_x() __nvvm_read_ptx_sreg_tid_x()
15+
#else
16+
#error "Unsupported target"
17+
#endif

‎offload/unittests/OffloadAPI/enqueue/olEnqueueKernelLaunch.cpp

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
//===------- Offload API tests - olEnqueueKernelLaunch --------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include "../common/Fixtures.hpp"
10+
#include <OffloadAPI.h>
11+
#include <gtest/gtest.h>
12+
13+
struct olEnqueueKernelLaunchTest : offloadQueueTest {
14+
void SetUp() override {
15+
RETURN_ON_FATAL_FAILURE(offloadQueueTest::SetUp());
16+
ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Platform, DeviceBin));
17+
ASSERT_GE(DeviceBin->size(), 0lu);
18+
ASSERT_SUCCESS(olCreateProgram(Device, DeviceBin->data(), DeviceBin->size(),
19+
&Program));
20+
ASSERT_SUCCESS(olCreateKernel(Program, "foo", &Kernel));
21+
}
22+
23+
void TearDown() override {
24+
if (Kernel) {
25+
olReleaseKernel(Kernel);
26+
}
27+
if (Program) {
28+
olReleaseProgram(Program);
29+
}
30+
RETURN_ON_FATAL_FAILURE(offloadQueueTest::TearDown());
31+
}
32+
33+
std::shared_ptr<std::vector<char>> DeviceBin;
34+
ol_program_handle_t Program = nullptr;
35+
ol_kernel_handle_t Kernel = nullptr;
36+
};
37+
38+
TEST_F(olEnqueueKernelLaunchTest, Success) {
39+
void *Mem;
40+
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_SHARED, 1024, &Mem));
41+
ol_kernel_launch_size_args_t LaunchArgs{};
42+
LaunchArgs.Dimensions = 1;
43+
LaunchArgs.GroupSizeX = 64;
44+
LaunchArgs.GroupSizeY = 1;
45+
LaunchArgs.GroupSizeZ = 1;
46+
47+
LaunchArgs.NumGroupsX = 1;
48+
LaunchArgs.NumGroupsY = 1;
49+
LaunchArgs.NumGroupsZ = 1;
50+
51+
ASSERT_SUCCESS(olSetKernelArgValue(Kernel, 0, sizeof(Mem), &Mem));
52+
ASSERT_SUCCESS(olEnqueueKernelLaunch(Queue, Kernel, &LaunchArgs, nullptr));
53+
54+
ASSERT_SUCCESS(olFinishQueue(Queue));
55+
56+
int *Data = (int *)Mem;
57+
for (int i = 0; i < 64; i++) {
58+
ASSERT_EQ(Data[i], i);
59+
}
60+
61+
ASSERT_SUCCESS(olMemFree(Device, OL_ALLOC_TYPE_SHARED, Mem));
62+
}

‎offload/unittests/OffloadAPI/kernel/olCreateKernel.cpp

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
//===------- Offload API tests - olCreateKernel ---------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include "../common/Fixtures.hpp"
10+
#include <OffloadAPI.h>
11+
#include <gtest/gtest.h>
12+
13+
using olCreateKernelTest = offloadProgramTest;
14+
15+
TEST_F(olCreateKernelTest, Success) {
16+
// std::shared_ptr<std::vector<char>> DeviceBin2;
17+
// ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Platform, DeviceBin2));
18+
19+
ol_kernel_handle_t Kernel = nullptr;
20+
ASSERT_SUCCESS(olCreateKernel(Program, "foo", &Kernel));
21+
ASSERT_NE(Kernel, nullptr);
22+
ASSERT_SUCCESS(olReleaseKernel(Kernel));
23+
}
24+
25+
TEST_F(olCreateKernelTest, InvalidNullProgram) {
26+
ol_kernel_handle_t Kernel = nullptr;
27+
ASSERT_ERROR(OL_ERRC_INVALID_NULL_HANDLE,
28+
olCreateKernel(nullptr, "foo", &Kernel));
29+
}
30+
31+
TEST_F(olCreateKernelTest, InvalidNullKernelPointer) {
32+
ASSERT_ERROR(OL_ERRC_INVALID_NULL_POINTER,
33+
olCreateKernel(Program, "foo", nullptr));
34+
}
35+
36+
// TEST_F(olCreateKernelTest, InvalidKernelName) {
37+
// ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Platform, DeviceBin));
38+
// ol_kernel_handle_t Kernel = nullptr;
39+
// ASSERT_ANY_ERROR(olCreateKernel(Program, "bad_kernel_name", &Kernel));
40+
// ASSERT_EQ(Kernel, nullptr);
41+
// }

‎offload/unittests/OffloadAPI/kernel/olReleaseKernel.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===------- Offload API tests - olReleaseKernel --------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include "../common/Fixtures.hpp"
10+
#include <OffloadAPI.h>
11+
#include <gtest/gtest.h>
12+
13+
using olReleaseKernelTest = offloadKernelTest;
14+
15+
TEST_F(olReleaseKernelTest, Success) {
16+
ASSERT_SUCCESS(olRetainKernel(Kernel));
17+
ASSERT_SUCCESS(olReleaseKernel(Kernel));
18+
}
19+
20+
TEST_F(olReleaseKernelTest, InvalidNullHandle) {
21+
ASSERT_ERROR(OL_ERRC_INVALID_NULL_HANDLE, olReleaseKernel(nullptr));
22+
}

‎offload/unittests/OffloadAPI/kernel/olRetainKernel.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
//===------- Offload API tests - olRetainKernel ---------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#include "../common/Fixtures.hpp"
10+
#include <OffloadAPI.h>
11+
#include <gtest/gtest.h>
12+
13+
using olRetainKernelTest = offloadKernelTest;
14+
15+
TEST_F(olRetainKernelTest, Success) {
16+
ASSERT_SUCCESS(olRetainKernel(Kernel));
17+
}
18+
19+
TEST_F(olRetainKernelTest, InvalidNullHandle) {
20+
ASSERT_ERROR(OL_ERRC_INVALID_NULL_HANDLE, olRetainKernel(nullptr));
21+
}

0 commit comments

Comments
 (0)