Skip to content

Commit dfc5abe

Browse files
veselypetaveselypeta
committed
[HIP] Fix Kernel Compilation on AMD
1 parent 3653e58 commit dfc5abe

File tree

9 files changed

+208
-17
lines changed

9 files changed

+208
-17
lines changed

‎.github/workflows/cmake.yml

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -164,7 +164,7 @@ jobs:
164164
matrix:
165165
adapter: [
166166
{name: CUDA, triplet: nvptx64-nvidia-cuda},
167-
{name: HIP, triplet: spir64}, # should be amdgcn-amdhsa, but build scripts for device binaries are currently broken for this target.
167+
{name: HIP, triplet: amdgcn-amd-amdhsa},
168168
{name: L0, triplet: spir64}
169169
]
170170
build_type: [Debug, Release]
@@ -197,6 +197,8 @@ jobs:
197197
-DUR_BUILD_ADAPTER_${{matrix.adapter.name}}=ON
198198
-DUR_DPCXX=${{github.workspace}}/dpcpp_compiler/bin/clang++
199199
-DUR_CONFORMANCE_TARGET_TRIPLES=${{matrix.adapter.triplet}}
200+
${{ matrix.adapter.name == 'HIP' && '-DAMD_ARCH=gfx1030' || '' }}
201+
${{ matrix.adapter.name == 'HIP' && '-DUR_HIP_PLATFORM=AMD' || '' }}
200202
201203
- name: Build
202204
# This is so that device binaries can find the sycl runtime library

‎source/adapters/hip/kernel.cpp

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,12 @@ urKernelCreate(ur_program_handle_t hProgram, const char *pKernelName,
2222
ScopedContext Active(hProgram->getContext()->getDevice());
2323

2424
hipFunction_t HIPFunc;
25-
UR_CHECK_ERROR(
26-
hipModuleGetFunction(&HIPFunc, hProgram->get(), pKernelName));
25+
hipError_t KernelError =
26+
hipModuleGetFunction(&HIPFunc, hProgram->get(), pKernelName);
27+
if (KernelError == hipErrorNotFound) {
28+
return UR_RESULT_ERROR_INVALID_KERNEL_NAME;
29+
}
30+
UR_CHECK_ERROR(KernelError);
2731

2832
std::string KernelNameWoffset = std::string(pKernelName) + "_with_offset";
2933
hipFunction_t HIPFuncWithOffsetParam;

‎test/conformance/device_code/CMakeLists.txt

Lines changed: 12 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,19 @@ macro(add_device_binary SOURCE_FILE)
99
file(MAKE_DIRECTORY ${DEVICE_BINARY_DIR})
1010
foreach(TRIPLE ${TARGET_TRIPLES})
1111
set(EXE_PATH "${DEVICE_BINARY_DIR}/${KERNEL_NAME}_${TRIPLE}")
12+
if(${TRIPLE} MATCHES "amd")
13+
set(AMD_TARGET_BACKEND -Xsycl-target-backend=${TRIPLE})
14+
set(AMD_OFFLOAD_ARCH --offload-arch=${AMD_ARCH})
15+
endif()
16+
# images are not yet supported in sycl on AMD
17+
if(${TRIPLE} MATCHES "amd" AND ${KERNEL_NAME} MATCHES "image_copy")
18+
continue()
19+
endif()
1220
add_custom_command(OUTPUT ${EXE_PATH}
13-
COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off
14-
${SOURCE_FILE} -o ${EXE_PATH}
21+
COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off
22+
${AMD_TARGET_BACKEND} ${AMD_OFFLOAD_ARCH} ${SOURCE_FILE}
23+
-o ${EXE_PATH}
24+
1525
COMMAND ${CMAKE_COMMAND} -E env SYCL_DUMP_IMAGES=true
1626
${EXE_PATH} || (exit 0)
1727
WORKING_DIRECTORY "${DEVICE_BINARY_DIR}"

‎test/conformance/enqueue/enqueue_adapter_hip.match

Lines changed: 87 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1,87 @@
1-
Segmentation fault
1+
{{OPT}}Segmentation Fault
2+
{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.Success/AMD_HIP_BACKEND___{{.*}}_
3+
{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.InvalidEventWaitInvalidEvent/AMD_HIP_BACKEND___{{.*}}_
4+
{{OPT}}urEnqueueDeviceGetGlobalVariableWriteTest.InvalidEventWaitInvalidEvent/AMD_HIP_BACKEND___{{.*}}_
5+
{{OPT}}urEnqueueMemBufferCopyRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___copy_row_2D
6+
{{OPT}}urEnqueueMemBufferCopyRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___copy_3d_2d
7+
{{OPT}}urEnqueueMemBufferCopyRectTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_
8+
{{OPT}}urEnqueueMemBufferFillTest.Success/AMD_HIP_BACKEND___{{.*}}___size__256__patternSize__256
9+
{{OPT}}urEnqueueMemBufferFillTest.Success/AMD_HIP_BACKEND___{{.*}}___size__1024__patternSize__256
10+
{{OPT}}urEnqueueMemBufferMapTest.SuccessMultiMaps/AMD_HIP_BACKEND___{{.*}}_
11+
{{OPT}}urEnqueueMemBufferReadTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_
12+
{{OPT}}urEnqueueMemBufferReadRectTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_
13+
{{OPT}}urEnqueueMemBufferWriteTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_
14+
{{OPT}}urEnqueueMemBufferWriteRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___write_row_2D
15+
{{OPT}}urEnqueueMemBufferWriteRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___write_3d_2d
16+
{{OPT}}urEnqueueMemBufferWriteRectTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_
17+
{{OPT}}urEnqueueMemImageCopyTest.Success/AMD_HIP_BACKEND___{{.*}}___1D
18+
{{OPT}}urEnqueueMemImageCopyTest.Success/AMD_HIP_BACKEND___{{.*}}___2D
19+
{{OPT}}urEnqueueMemImageCopyTest.Success/AMD_HIP_BACKEND___{{.*}}___3D
20+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopy/AMD_HIP_BACKEND___{{.*}}___1D
21+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopy/AMD_HIP_BACKEND___{{.*}}___2D
22+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopy/AMD_HIP_BACKEND___{{.*}}___3D
23+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithSrcOffset/AMD_HIP_BACKEND___{{.*}}___1D
24+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithSrcOffset/AMD_HIP_BACKEND___{{.*}}___2D
25+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithSrcOffset/AMD_HIP_BACKEND___{{.*}}___3D
26+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithDstOffset/AMD_HIP_BACKEND___{{.*}}___1D
27+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithDstOffset/AMD_HIP_BACKEND___{{.*}}___2D
28+
{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithDstOffset/AMD_HIP_BACKEND___{{.*}}___3D
29+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleQueue/AMD_HIP_BACKEND___{{.*}}___1D
30+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleQueue/AMD_HIP_BACKEND___{{.*}}___3D
31+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageSrc/AMD_HIP_BACKEND___{{.*}}___1D
32+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageSrc/AMD_HIP_BACKEND___{{.*}}___3D
33+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageDst/AMD_HIP_BACKEND___{{.*}}___1D
34+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageDst/AMD_HIP_BACKEND___{{.*}}___3D
35+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullPtrEventWaitList/AMD_HIP_BACKEND___{{.*}}___1D
36+
{{OPT}}urEnqueueMemImageCopyTest.InvalidNullPtrEventWaitList/AMD_HIP_BACKEND___{{.*}}___3D
37+
{{OPT}}urEnqueueMemImageCopyTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___1D
38+
{{OPT}}urEnqueueMemImageCopyTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___2D
39+
{{OPT}}urEnqueueMemImageCopyTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___3D
40+
{{OPT}}urEnqueueMemImageReadTest.Success1D/AMD_HIP_BACKEND___{{.*}}_
41+
{{OPT}}urEnqueueMemImageReadTest.Success3D/AMD_HIP_BACKEND___{{.*}}_
42+
{{OPT}}urEnqueueMemImageReadTest.InvalidOrigin1D/AMD_HIP_BACKEND___{{.*}}_
43+
{{OPT}}urEnqueueMemImageReadTest.InvalidOrigin2D/AMD_HIP_BACKEND___{{.*}}_
44+
{{OPT}}urEnqueueMemImageReadTest.InvalidOrigin3D/AMD_HIP_BACKEND___{{.*}}_
45+
{{OPT}}urEnqueueMemImageReadTest.InvalidRegion1D/AMD_HIP_BACKEND___{{.*}}_
46+
{{OPT}}urEnqueueMemImageReadTest.InvalidRegion2D/AMD_HIP_BACKEND___{{.*}}_
47+
{{OPT}}urEnqueueMemImageReadTest.InvalidRegion3D/AMD_HIP_BACKEND___{{.*}}_
48+
{{OPT}}urEnqueueMemImageWriteTest.Success1D/AMD_HIP_BACKEND___{{.*}}_
49+
{{OPT}}urEnqueueMemImageWriteTest.Success3D/AMD_HIP_BACKEND___{{.*}}_
50+
{{OPT}}urEnqueueMemImageWriteTest.InvalidOrigin1D/AMD_HIP_BACKEND___{{.*}}_
51+
{{OPT}}urEnqueueMemImageWriteTest.InvalidOrigin2D/AMD_HIP_BACKEND___{{.*}}_
52+
{{OPT}}urEnqueueMemImageWriteTest.InvalidOrigin3D/AMD_HIP_BACKEND___{{.*}}_
53+
{{OPT}}urEnqueueMemImageWriteTest.InvalidRegion1D/AMD_HIP_BACKEND___{{.*}}_
54+
{{OPT}}urEnqueueMemImageWriteTest.InvalidRegion2D/AMD_HIP_BACKEND___{{.*}}_
55+
{{OPT}}urEnqueueMemImageWriteTest.InvalidRegion3D/AMD_HIP_BACKEND___{{.*}}_
56+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1__patternSize__1
57+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1__patternSize__256
58+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1__patternSize__4
59+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__57__height__1__patternSize__1
60+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1__patternSize__256
61+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1__patternSize__1024
62+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256__patternSize__1
63+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256__patternSize__256
64+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256__patternSize__65536
65+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__1__patternSize__1
66+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__35__patternSize__1
67+
{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__35__patternSize__128
68+
{{OPT}}urEnqueueUSMFill2DNegativeTest.OutOfBounds/AMD_HIP_BACKEND___{{.*}}_
69+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1
70+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1
71+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1
72+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256
73+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__23
74+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__1
75+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1
76+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1
77+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1
78+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256
79+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__23
80+
{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__1
81+
{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidNullHandleQueue/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1
82+
{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidNullPointer/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1
83+
{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1
84+
{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidEventWaitList/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1
85+
{{OPT}}urEnqueueUSMPrefetchWithParamTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_USM_MIGRATION_FLAG_DEFAULT
86+
{{OPT}}urEnqueueUSMPrefetchWithParamTest.CheckWaitEvent/AMD_HIP_BACKEND___{{.*}}___UR_USM_MIGRATION_FLAG_DEFAULT
87+
{{OPT}}urEnqueueUSMPrefetchTest.InvalidSizeTooLarge/AMD_HIP_BACKEND___{{.*}}_

‎test/conformance/kernel/kernel_adapter_hip.match

Lines changed: 25 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1,25 @@
1-
Segmentation fault
1+
{{OPT}}Segmentation Fault
2+
{{OPT}}urKernelGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_NUM_REGS
3+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_FUNCTION_NAME
4+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_NUM_ARGS
5+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_REFERENCE_COUNT
6+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_CONTEXT
7+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_PROGRAM
8+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES
9+
{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_NUM_REGS
10+
{{OPT}}urKernelSetArgLocalTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_
11+
{{OPT}}urKernelSetArgMemObjTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_
12+
{{OPT}}urKernelSetArgPointerTest.SuccessShared/AMD_HIP_BACKEND___{{.*}}_
13+
{{OPT}}urKernelSetArgPointerNegativeTest.InvalidNullHandleKernel/AMD_HIP_BACKEND___{{.*}}_
14+
{{OPT}}urKernelSetArgPointerNegativeTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_
15+
{{OPT}}urKernelSetArgSamplerTest.Success/AMD_HIP_BACKEND___{{.*}}_
16+
{{OPT}}urKernelSetArgSamplerTest.InvalidNullHandleKernel/AMD_HIP_BACKEND___{{.*}}_
17+
{{OPT}}urKernelSetArgSamplerTest.InvalidNullHandleArgValue/AMD_HIP_BACKEND___{{.*}}_
18+
{{OPT}}urKernelSetArgSamplerTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_
19+
{{OPT}}urKernelSetArgValueTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_
20+
{{OPT}}urKernelSetArgValueTest.InvalidKernelArgumentSize/AMD_HIP_BACKEND___{{.*}}_
21+
{{OPT}}urKernelSetExecInfoUSMPointersTest.SuccessShared/AMD_HIP_BACKEND___{{.*}}_
22+
{{OPT}}urKernelSetSpecializationConstantsTest.Success/AMD_HIP_BACKEND___{{.*}}_
23+
{{OPT}}urKernelSetSpecializationConstantsTest.InvalidNullHandleKernel/AMD_HIP_BACKEND___{{.*}}_
24+
{{OPT}}urKernelSetSpecializationConstantsTest.InvalidNullPointerSpecConstants/AMD_HIP_BACKEND___{{.*}}_
25+
{{OPT}}urKernelSetSpecializationConstantsTest.InvalidSizeCount/AMD_HIP_BACKEND___{{.*}}_

‎test/conformance/kernel/urKernelSetArgSampler.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,14 @@
77

88
struct urKernelSetArgSamplerTest : uur::urKernelTest {
99
void SetUp() {
10+
// Images and samplers are not available on AMD
11+
ur_platform_backend_t backend;
12+
ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND,
13+
sizeof(backend), &backend, nullptr));
14+
if (backend == UR_PLATFORM_BACKEND_HIP) {
15+
GTEST_SKIP() << "Sampler are not supported on hip.";
16+
}
17+
1018
program_name = "image_copy";
1119
UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp());
1220
ur_sampler_desc_t sampler_desc = {

‎test/conformance/program/program_adapter_hip.match

Lines changed: 25 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1 +1,25 @@
1-
Segmentation fault
1+
{{OPT}}Segmentation Fault
2+
{{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_
3+
{{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/AMD_HIP_BACKEND___{{.*}}_
4+
{{OPT}}urProgramGetBuildInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE
5+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_STATUS
6+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS
7+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_LOG
8+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE
9+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_STATUS
10+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS
11+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_LOG
12+
{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE
13+
{{OPT}}urProgramGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS
14+
{{OPT}}urProgramGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES
15+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_REFERENCE_COUNT
16+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_CONTEXT
17+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_DEVICES
18+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_DEVICES
19+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_SOURCE
20+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_BINARY_SIZES
21+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_BINARIES
22+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS
23+
{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES
24+
{{OPT}}urProgramLinkTest.Success/AMD_HIP_BACKEND___{{.*}}_
25+
{{OPT}}urProgramSetSpecializationConstantsTest.Success/AMD_HIP_BACKEND___{{.*}}_

‎test/conformance/source/environment.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -266,6 +266,17 @@ std::string KernelsEnvironment::getSupportedILPostfix(uint32_t device_index) {
266266
return {};
267267
}
268268

269+
// special case for AMD as it doesn't support IL.
270+
ur_platform_backend_t backend;
271+
if (urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, sizeof(backend),
272+
&backend, nullptr)) {
273+
error = "failed to get backend from platform.";
274+
return {};
275+
}
276+
if (backend == UR_PLATFORM_BACKEND_HIP) {
277+
return ".bin";
278+
}
279+
269280
auto device = instance->GetDevices()[device_index];
270281
std::string IL_version;
271282
if (uur::GetDeviceILVersion(device, IL_version)) {

‎test/conformance/testing/include/uur/fixtures.h

Lines changed: 31 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1032,15 +1032,37 @@ struct urKernelExecutionTest : urKernelTest {
10321032
ASSERT_SUCCESS(urKernelSetArgMemObj(kernel, current_arg_index, nullptr,
10331033
mem_handle));
10341034

1035-
// This emulates the offset struct sycl adds for a 1D buffer accessor.
1036-
struct {
1037-
size_t offsets[1] = {0};
1038-
} accessor;
1039-
ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 1,
1040-
sizeof(accessor), nullptr,
1041-
&accessor));
1042-
1043-
current_arg_index += 2;
1035+
// SYCL device kernels have different interfaces depending on the
1036+
// backend being used. Typically a kernel which takes a buffer argument
1037+
// will take a pointer to the start of the buffer and a sycl::id param
1038+
// which is a struct that encodes the accessor to the buffer. However
1039+
// the AMD backend handles this differently and uses three separate
1040+
// arguments for each of the three dimensions of the accessor.
1041+
1042+
ur_platform_backend_t backend;
1043+
ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND,
1044+
sizeof(backend), &backend, nullptr));
1045+
if (backend == UR_PLATFORM_BACKEND_HIP) {
1046+
// this emulates the three offset params for buffer accessor on AMD.
1047+
size_t val = 0;
1048+
ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 1,
1049+
sizeof(size_t), nullptr, &val));
1050+
ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 2,
1051+
sizeof(size_t), nullptr, &val));
1052+
ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 3,
1053+
sizeof(size_t), nullptr, &val));
1054+
current_arg_index += 4;
1055+
} else {
1056+
// This emulates the offset struct sycl adds for a 1D buffer accessor.
1057+
struct {
1058+
size_t offsets[1] = {0};
1059+
} accessor;
1060+
ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 1,
1061+
sizeof(accessor), nullptr,
1062+
&accessor));
1063+
current_arg_index += 2;
1064+
}
1065+
10441066
buffer_args.push_back(mem_handle);
10451067
*out_buffer = mem_handle;
10461068
}

0 commit comments

Comments
 (0)