[HIP] Correctly set HIP Kernel Buffer Arguments

.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,20 @@ 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})
+            set(AMD_NOGPULIB -nogpulib)
+        endif()
+        # images are not yet supported in sycl on AMD
+        if(${TRIPLE} MATCHES "amd" AND ${KERNEL_NAME} MATCHES "image_copy")
+            continue()
+        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} ${AMD_NOGPULIB} ${SOURCE_FILE} 
+            -o ${EXE_PATH}
+
             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;
     }