diff --git a/.github/workflows/cmake.yml b/.github/workflows/cmake.yml index 9edbd459d7..2b38d9ffc7 100644 --- a/.github/workflows/cmake.yml +++ b/.github/workflows/cmake.yml @@ -164,7 +164,7 @@ jobs: matrix: adapter: [ {name: CUDA, triplet: nvptx64-nvidia-cuda}, - {name: HIP, triplet: spir64}, # should be amdgcn-amdhsa, but build scripts for device binaries are currently broken for this target. + {name: HIP, triplet: amdgcn-amd-amdhsa}, {name: L0, triplet: spir64} ] build_type: [Debug, Release] @@ -180,7 +180,7 @@ jobs: - name: Download DPC++ run: | - wget -O ${{github.workspace}}/dpcpp_compiler.tar.gz https://github.com/intel/llvm/releases/download/nightly-2023-09-21/sycl_linux.tar.gz + wget -O ${{github.workspace}}/dpcpp_compiler.tar.gz https://github.com/intel/llvm/releases/download/nightly-2023-10-18/sycl_linux.tar.gz mkdir dpcpp_compiler tar -xvf ${{github.workspace}}/dpcpp_compiler.tar.gz -C dpcpp_compiler @@ -197,6 +197,8 @@ jobs: -DUR_BUILD_ADAPTER_${{matrix.adapter.name}}=ON -DUR_DPCXX=${{github.workspace}}/dpcpp_compiler/bin/clang++ -DUR_CONFORMANCE_TARGET_TRIPLES=${{matrix.adapter.triplet}} + ${{ matrix.adapter.name == 'HIP' && '-DAMD_ARCH=gfx1030' || '' }} + ${{ matrix.adapter.name == 'HIP' && '-DUR_HIP_PLATFORM=AMD' || '' }} - name: Build # This is so that device binaries can find the sycl runtime library diff --git a/source/adapters/hip/device.cpp b/source/adapters/hip/device.cpp index 7cec6def8b..b2bbce7817 100644 --- a/source/adapters/hip/device.cpp +++ b/source/adapters/hip/device.cpp @@ -30,6 +30,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, static constexpr uint32_t MaxWorkItemDimensions = 3u; switch ((uint32_t)propName) { + case UR_DEVICE_INFO_VIRTUAL_MEMORY_SUPPORT: { + return ReturnValue(false); + } case UR_DEVICE_INFO_TYPE: { return ReturnValue(UR_DEVICE_TYPE_GPU); } diff --git a/source/adapters/hip/kernel.cpp b/source/adapters/hip/kernel.cpp index b433d3a3b4..724a255cad 100644 --- a/source/adapters/hip/kernel.cpp +++ b/source/adapters/hip/kernel.cpp @@ -22,8 +22,12 @@ urKernelCreate(ur_program_handle_t hProgram, const char *pKernelName, ScopedContext Active(hProgram->getContext()->getDevice()); hipFunction_t HIPFunc; - UR_CHECK_ERROR( - hipModuleGetFunction(&HIPFunc, hProgram->get(), pKernelName)); + hipError_t KernelError = + hipModuleGetFunction(&HIPFunc, hProgram->get(), pKernelName); + if (KernelError == hipErrorNotFound) { + return UR_RESULT_ERROR_INVALID_KERNEL_NAME; + } + UR_CHECK_ERROR(KernelError); std::string KernelNameWoffset = std::string(pKernelName) + "_with_offset"; hipFunction_t HIPFuncWithOffsetParam; diff --git a/test/conformance/device_code/CMakeLists.txt b/test/conformance/device_code/CMakeLists.txt index 1d3f28df7f..646da15241 100644 --- a/test/conformance/device_code/CMakeLists.txt +++ b/test/conformance/device_code/CMakeLists.txt @@ -9,9 +9,13 @@ macro(add_device_binary SOURCE_FILE) file(MAKE_DIRECTORY ${DEVICE_BINARY_DIR}) foreach(TRIPLE ${TARGET_TRIPLES}) set(EXE_PATH "${DEVICE_BINARY_DIR}/${KERNEL_NAME}_${TRIPLE}") + if(${TRIPLE} MATCHES "amd") + set(AMD_TARGET_BACKEND -Xsycl-target-backend=${TRIPLE}) + set(AMD_OFFLOAD_ARCH --offload-arch=${AMD_ARCH}) + endif() add_custom_command(OUTPUT ${EXE_PATH} - COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off - ${SOURCE_FILE} -o ${EXE_PATH} + COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off ${AMD_TARGET_BACKEND} ${AMD_OFFLOAD_ARCH} + ${SOURCE_FILE} -o ${EXE_PATH} || (exit 0) COMMAND ${CMAKE_COMMAND} -E env SYCL_DUMP_IMAGES=true ${EXE_PATH} || (exit 0) WORKING_DIRECTORY "${DEVICE_BINARY_DIR}" diff --git a/test/conformance/enqueue/enqueue_adapter_hip.match b/test/conformance/enqueue/enqueue_adapter_hip.match index 7a1c0d5b8e..9d48681c1a 100644 --- a/test/conformance/enqueue/enqueue_adapter_hip.match +++ b/test/conformance/enqueue/enqueue_adapter_hip.match @@ -1 +1,87 @@ -Segmentation fault +{{OPT}}Segmentation Fault +{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.Success/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueDeviceGetGlobalVariableReadTest.InvalidEventWaitInvalidEvent/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueDeviceGetGlobalVariableWriteTest.InvalidEventWaitInvalidEvent/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemBufferCopyRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___copy_row_2D +{{OPT}}urEnqueueMemBufferCopyRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___copy_3d_2d +{{OPT}}urEnqueueMemBufferCopyRectTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemBufferFillTest.Success/AMD_HIP_BACKEND___{{.*}}___size__256__patternSize__256 +{{OPT}}urEnqueueMemBufferFillTest.Success/AMD_HIP_BACKEND___{{.*}}___size__1024__patternSize__256 +{{OPT}}urEnqueueMemBufferMapTest.SuccessMultiMaps/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemBufferReadTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemBufferReadRectTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemBufferWriteTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemBufferWriteRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___write_row_2D +{{OPT}}urEnqueueMemBufferWriteRectTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___write_3d_2d +{{OPT}}urEnqueueMemBufferWriteRectTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageCopyTest.Success/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.Success/AMD_HIP_BACKEND___{{.*}}___2D +{{OPT}}urEnqueueMemImageCopyTest.Success/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopy/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopy/AMD_HIP_BACKEND___{{.*}}___2D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopy/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithSrcOffset/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithSrcOffset/AMD_HIP_BACKEND___{{.*}}___2D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithSrcOffset/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithDstOffset/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithDstOffset/AMD_HIP_BACKEND___{{.*}}___2D +{{OPT}}urEnqueueMemImageCopyTest.SuccessPartialCopyWithDstOffset/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleQueue/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleQueue/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageSrc/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageSrc/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageDst/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullHandleImageDst/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullPtrEventWaitList/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.InvalidNullPtrEventWaitList/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageCopyTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___1D +{{OPT}}urEnqueueMemImageCopyTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___2D +{{OPT}}urEnqueueMemImageCopyTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___3D +{{OPT}}urEnqueueMemImageReadTest.Success1D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.Success3D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.InvalidOrigin1D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.InvalidOrigin2D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.InvalidOrigin3D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.InvalidRegion1D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.InvalidRegion2D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageReadTest.InvalidRegion3D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.Success1D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.Success3D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.InvalidOrigin1D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.InvalidOrigin2D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.InvalidOrigin3D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.InvalidRegion1D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.InvalidRegion2D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueMemImageWriteTest.InvalidRegion3D/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1__patternSize__1 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1__patternSize__256 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1__patternSize__4 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__57__height__1__patternSize__1 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1__patternSize__256 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1__patternSize__1024 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256__patternSize__1 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256__patternSize__256 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256__patternSize__65536 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__1__patternSize__1 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__35__patternSize__1 +{{OPT}}urEnqueueUSMFill2DTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__35__patternSize__128 +{{OPT}}urEnqueueUSMFill2DNegativeTest.OutOfBounds/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__23 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__1024__height__1 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__1024__width__256__height__256 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__23 +{{OPT}}urEnqueueUSMMemcpy2DTestWithParam.SuccessNonBlocking/AMD_HIP_BACKEND___{{.*}}___pitch__234__width__233__height__1 +{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidNullHandleQueue/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1 +{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidNullPointer/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1 +{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1 +{{OPT}}urEnqueueUSMMemcpy2DNegativeTest.InvalidEventWaitList/AMD_HIP_BACKEND___{{.*}}___pitch__1__width__1__height__1 +{{OPT}}urEnqueueUSMPrefetchWithParamTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_USM_MIGRATION_FLAG_DEFAULT +{{OPT}}urEnqueueUSMPrefetchWithParamTest.CheckWaitEvent/AMD_HIP_BACKEND___{{.*}}___UR_USM_MIGRATION_FLAG_DEFAULT +{{OPT}}urEnqueueUSMPrefetchTest.InvalidSizeTooLarge/AMD_HIP_BACKEND___{{.*}}_ diff --git a/test/conformance/kernel/kernel_adapter_hip.match b/test/conformance/kernel/kernel_adapter_hip.match index 7a1c0d5b8e..96d579f088 100644 --- a/test/conformance/kernel/kernel_adapter_hip.match +++ b/test/conformance/kernel/kernel_adapter_hip.match @@ -1 +1,25 @@ -Segmentation fault +{{OPT}}Segmentation Fault +{{OPT}}urKernelGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_NUM_REGS +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_FUNCTION_NAME +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_NUM_ARGS +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_REFERENCE_COUNT +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_CONTEXT +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_PROGRAM +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_ATTRIBUTES +{{OPT}}urKernelGetInfoTest.InvalidSizeSmall/AMD_HIP_BACKEND___{{.*}}___UR_KERNEL_INFO_NUM_REGS +{{OPT}}urKernelSetArgLocalTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgMemObjTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgPointerTest.SuccessShared/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgPointerNegativeTest.InvalidNullHandleKernel/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgPointerNegativeTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgSamplerTest.Success/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgSamplerTest.InvalidNullHandleKernel/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgSamplerTest.InvalidNullHandleArgValue/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgSamplerTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgValueTest.InvalidKernelArgumentIndex/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetArgValueTest.InvalidKernelArgumentSize/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetExecInfoUSMPointersTest.SuccessShared/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetSpecializationConstantsTest.Success/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetSpecializationConstantsTest.InvalidNullHandleKernel/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetSpecializationConstantsTest.InvalidNullPointerSpecConstants/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urKernelSetSpecializationConstantsTest.InvalidSizeCount/AMD_HIP_BACKEND___{{.*}}_ diff --git a/test/conformance/kernel/urKernelCreate.cpp b/test/conformance/kernel/urKernelCreate.cpp index f89eb2a72a..ac4f3fdf9f 100644 --- a/test/conformance/kernel/urKernelCreate.cpp +++ b/test/conformance/kernel/urKernelCreate.cpp @@ -28,12 +28,12 @@ struct urKernelCreateTest : uur::urProgramTest { UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urKernelCreateTest); TEST_P(urKernelCreateTest, Success) { - ASSERT_SUCCESS(urKernelCreate(program, kernel_name.data(), &kernel)); + ASSERT_SUCCESS(urKernelCreate(program, kernel_name.c_str(), &kernel)); } TEST_P(urKernelCreateTest, InvalidNullHandleProgram) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_HANDLE, - urKernelCreate(nullptr, kernel_name.data(), &kernel)); + urKernelCreate(nullptr, kernel_name.c_str(), &kernel)); } TEST_P(urKernelCreateTest, InvalidNullPointerName) { @@ -43,11 +43,11 @@ TEST_P(urKernelCreateTest, InvalidNullPointerName) { TEST_P(urKernelCreateTest, InvalidNullPointerKernel) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_POINTER, - urKernelCreate(program, kernel_name.data(), nullptr)); + urKernelCreate(program, kernel_name.c_str(), nullptr)); } TEST_P(urKernelCreateTest, InvalidKernelName) { std::string invalid_name = "incorrect_kernel_name"; ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_KERNEL_NAME, - urKernelCreate(program, invalid_name.data(), &kernel)); + urKernelCreate(program, invalid_name.c_str(), &kernel)); } diff --git a/test/conformance/kernel/urKernelSetArgSampler.cpp b/test/conformance/kernel/urKernelSetArgSampler.cpp index 4a044383ff..814b79a153 100644 --- a/test/conformance/kernel/urKernelSetArgSampler.cpp +++ b/test/conformance/kernel/urKernelSetArgSampler.cpp @@ -7,6 +7,14 @@ struct urKernelSetArgSamplerTest : uur::urKernelTest { void SetUp() { + // Images and samplers are not available on AMD + ur_platform_backend_t backend; + ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, + sizeof(backend), &backend, nullptr)); + if (backend == UR_PLATFORM_BACKEND_HIP) { + GTEST_SKIP() << "Sampler are not supported on hip."; + } + program_name = "image_copy"; UUR_RETURN_ON_FATAL_FAILURE(urKernelTest::SetUp()); ur_sampler_desc_t sampler_desc = { diff --git a/test/conformance/program/program_adapter_hip.match b/test/conformance/program/program_adapter_hip.match index 7a1c0d5b8e..67f98ec2f7 100644 --- a/test/conformance/program/program_adapter_hip.match +++ b/test/conformance/program/program_adapter_hip.match @@ -1 +1,25 @@ -Segmentation fault +{{OPT}}Segmentation Fault +{{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urProgramCreateWithNativeHandleTest.InvalidNullPointerProgram/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urProgramGetBuildInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_STATUS +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_LOG +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_STATUS +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_OPTIONS +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_LOG +{{OPT}}urProgramGetBuildInfoTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_BUILD_INFO_BINARY_TYPE +{{OPT}}urProgramGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS +{{OPT}}urProgramGetInfoTest.Success/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_REFERENCE_COUNT +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_CONTEXT +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_DEVICES +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_DEVICES +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_SOURCE +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_BINARY_SIZES +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_BINARIES +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_NUM_KERNELS +{{OPT}}urProgramGetInfoTest.InvalidNullHandleProgram/AMD_HIP_BACKEND___{{.*}}___UR_PROGRAM_INFO_KERNEL_NAMES +{{OPT}}urProgramLinkTest.Success/AMD_HIP_BACKEND___{{.*}}_ +{{OPT}}urProgramSetSpecializationConstantsTest.Success/AMD_HIP_BACKEND___{{.*}}_ diff --git a/test/conformance/source/environment.cpp b/test/conformance/source/environment.cpp index 287310f679..875ceb63ef 100644 --- a/test/conformance/source/environment.cpp +++ b/test/conformance/source/environment.cpp @@ -266,6 +266,17 @@ std::string KernelsEnvironment::getSupportedILPostfix(uint32_t device_index) { return {}; } + // special case for AMD as it doesn't support IL. + ur_platform_backend_t backend; + if (urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, sizeof(backend), + &backend, nullptr)) { + error = "failed to get backend from platform."; + return {}; + } + if (backend == UR_PLATFORM_BACKEND_HIP) { + return ".bin"; + } + auto device = instance->GetDevices()[device_index]; std::string IL_version; if (uur::GetDeviceILVersion(device, IL_version)) { diff --git a/test/conformance/testing/include/uur/fixtures.h b/test/conformance/testing/include/uur/fixtures.h index fbb8a48fb1..2c6cc1dde9 100644 --- a/test/conformance/testing/include/uur/fixtures.h +++ b/test/conformance/testing/include/uur/fixtures.h @@ -1032,15 +1032,37 @@ struct urKernelExecutionTest : urKernelTest { ASSERT_SUCCESS(urKernelSetArgMemObj(kernel, current_arg_index, nullptr, mem_handle)); - // This emulates the offset struct sycl adds for a 1D buffer accessor. - struct { - size_t offsets[1] = {0}; - } accessor; - ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 1, - sizeof(accessor), nullptr, - &accessor)); - - current_arg_index += 2; + // SYCL device kernels have different interfaces depending on the + // backend being used. Typically a kernel which takes a buffer argument + // will take a pointer to the start of the buffer and a sycl::id param + // which is a struct that encodes the accessor to the buffer. However + // the AMD backend handles this differently and uses three separate + // arguments for each of the three dimensions of the accessor. + + ur_platform_backend_t backend; + ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, + sizeof(backend), &backend, nullptr)); + if (backend == UR_PLATFORM_BACKEND_HIP) { + // this emulates the three offset params for buffer accessor on AMD. + size_t val = 0; + ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 1, + sizeof(size_t), nullptr, &val)); + ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 2, + sizeof(size_t), nullptr, &val)); + ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 3, + sizeof(size_t), nullptr, &val)); + current_arg_index += 4; + } else { + // This emulates the offset struct sycl adds for a 1D buffer accessor. + struct { + size_t offsets[1] = {0}; + } accessor; + ASSERT_SUCCESS(urKernelSetArgValue(kernel, current_arg_index + 1, + sizeof(accessor), nullptr, + &accessor)); + current_arg_index += 2; + } + buffer_args.push_back(mem_handle); *out_buffer = mem_handle; } diff --git a/test/conformance/virtual_memory/virtual_memory_adapter_hip.match b/test/conformance/virtual_memory/virtual_memory_adapter_hip.match index 2e26995f76..7df750cfdb 100644 --- a/test/conformance/virtual_memory/virtual_memory_adapter_hip.match +++ b/test/conformance/virtual_memory/virtual_memory_adapter_hip.match @@ -1,55 +1,3 @@ -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___1 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___2 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___3 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___7 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___12 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___44 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___1024 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___4000 -urPhysicalMemCreateTest.Success/AMD_HIP_BACKEND___{{.*}}___12345 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___1 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___2 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___3 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___7 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___12 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___44 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___1024 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___4000 -urPhysicalMemCreateTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}___12345 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___1 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___2 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___3 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___7 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___12 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___44 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___1024 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___4000 -urPhysicalMemCreateTest.InvalidNullHandleDevice/AMD_HIP_BACKEND___{{.*}}___12345 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___1 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___2 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___3 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___7 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___12 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___44 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___1024 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___4000 -urPhysicalMemCreateTest.InvalidNullPointerPhysicalMem/AMD_HIP_BACKEND___{{.*}}___12345 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___1 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___2 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___3 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___7 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___12 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___44 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___1024 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___4000 -urPhysicalMemCreateTest.InvalidSize/AMD_HIP_BACKEND___{{.*}}___12345 -urPhysicalMemReleaseTest.Success/AMD_HIP_BACKEND___{{.*}}_ -urPhysicalMemReleaseTest.InvalidNullHandlePhysicalMem/AMD_HIP_BACKEND___{{.*}}_ -urPhysicalMemRetainTest.Success/AMD_HIP_BACKEND___{{.*}}_ -urPhysicalMemRetainTest.InvalidNullHandlePhysicalMem/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemFreeTest.Success/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemFreeTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemFreeTest.InvalidNullPointerStart/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemGetInfoTestWithParam.Success/AMD_HIP_BACKEND___{{.*}}___UR_VIRTUAL_MEM_INFO_ACCESS_MODE urVirtualMemGetInfoTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemGetInfoTest.InvalidNullPointerStart/AMD_HIP_BACKEND___{{.*}}_ @@ -62,42 +10,6 @@ urVirtualMemGranularityGetInfoNegativeTest.InvalidNullPointerPropSizeRet/AMD_HIP urVirtualMemGranularityGetInfoNegativeTest.InvalidNullPointerPropValue/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemGranularityGetInfoNegativeTest.InvalidPropSizeZero/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemGranularityGetInfoNegativeTest.InvalidSizePropSizeSmall/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemMapTest.Success/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemMapTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemMapTest.InvalidNullHandlePhysicalMem/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemMapTest.InvalidNullPointerStart/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemMapTest.InvalidEnumerationFlags/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___2 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___4 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___8 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___16 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___32 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___64 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___128 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___256 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___512 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___1024 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___2048 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___5000 -urVirtualMemReserveTestWithParam.SuccessNoStartPointer/AMD_HIP_BACKEND___{{.*}}___100000 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___2 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___4 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___8 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___16 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___32 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___64 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___128 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___256 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___512 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___1024 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___2048 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___5000 -urVirtualMemReserveTestWithParam.SuccessWithStartPointer/AMD_HIP_BACKEND___{{.*}}___100000 -urVirtualMemReserveTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemReserveTest.InvalidNullPointer/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemSetAccessTest.Success/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemSetAccessTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ urVirtualMemSetAccessTest.InvalidNullPointerStart/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemUnmapTest.Success/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemUnmapTest.InvalidNullHandleContext/AMD_HIP_BACKEND___{{.*}}_ -urVirtualMemUnmapTest.InvalidNullPointerStart/AMD_HIP_BACKEND___{{.*}}_