Skip to content

WorkGroupScratchMemory/copy_dynamic_size.cpp Fails on cuda when built for multiple triples #16968

@ayylol

Description

@ayylol

Describe the bug

WorkGroupScratchMemory/copy_dynamic_size.cpp will fail executing on cuda devices when it is built for more than just the nvptx triple. The test passes fine when it is built only for nvptx.

To reproduce

To reproduce first build the test for spir and nvptx:

llvm-lit WorkGroupScratchMemory/copy_dynamic_size.cpp --param test-mode=build-only --param sycl_build_targets="spir;nvidia" -a

Then run for cuda:gpu

llvm-lit WorkGroupScratchMemory/copy_dynamic_size.cpp --param test-mode=run-only --param sycl_devices=cuda:gpu -a

I got the following error:

-- Testing: 1 tests, 1 workers --
FAIL: SYCL :: WorkGroupScratchMemory/copy_dynamic_size.cpp (1 of 1)
******************** TEST 'SYCL :: WorkGroupScratchMemory/copy_dynamic_size.cpp' FAILED ********************
Exit Code: 134

Command Output (stderr):
--
RUN: at line 2: env UR_CUDA_ENABLE_IMAGE_SUPPORT=1 ONEAPI_DEVICE_SELECTOR=cuda:gpu  /iusers/dgarciao/llvm/build/tools/sycl/test-e2e/WorkGroupScratchMemory/Output/copy_dynamic_size.cpp.tmp.out
+ env UR_CUDA_ENABLE_IMAGE_SUPPORT=1 ONEAPI_DEVICE_SELECTOR=cuda:gpu /iusers/dgarciao/llvm/build/tools/sycl/test-e2e/WorkGroupScratchMemory/Output/copy_dynamic_size.cpp.tmp.out
<CUDA>[ERROR]:
UR CUDA ERROR:
        Value:           700
        Name:            CUDA_ERROR_ILLEGAL_ADDRESS
        Description:     an illegal memory access was encountered
        Function:        wait
        Source Location: /iusers/dgarciao/llvm/build/_deps/unified-runtime-src/source/adapters/cuda/event.cpp:134

terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  Native API failed. Native API returns: 2147483646 (UR_RESULT_ERROR_UNKNOWN)
/iusers/dgarciao/llvm/build/tools/sycl/test-e2e/WorkGroupScratchMemory/Output/copy_dynamic_size.cpp.script: line 1: 4046523 Aborted                 (core dumped) env UR_CUDA_ENABLE_IMAGE_SUPPORT=1 ONEAPI_DEVICE_SELECTOR=cuda:gpu /iusers/dgarciao/llvm/build/tools/sycl/test-e2e/WorkGroupScratchMemory/Output/copy_dynamic_size.cpp.tmp.out

--

If instead we only build for nvptx in the first step, the test will pass

Environment

  • OS: Linux
  • Target device and vendor: cuda:gpu
  • DPC++ version: b8a2ab7

Additional context

No response

Metadata

Metadata

Assignees

No one assigned

    Labels

    bugSomething isn't workingcudaCUDA back-end

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions