[HIP] Fix Kernel Compilation on AMD

.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]
@@ -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

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;

test/conformance/device_code/CMakeLists.txt

@@ -9,9 +9,15 @@ 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}"

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___{{.*}}_

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___{{.*}}_

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 = {

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___{{.*}}_

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)) {

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;
  }