UR CUDA ERROR Runtime error: CUDA_ERROR_INVALID_VALUE, unified runtime error

[vidyalatha-badde]

Hi,

UR CUDA ERROR:
        Value:           1
        Name:            CUDA_ERROR_INVALID_VALUE
        Description:     invalid argument
        Function:        urEnqueueKernelLaunch
        Source Location: sycl_workspace/llvm/sycl/plugins/unified_runtime/ur/adapters/cuda/enqueue.cpp:418

terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -30 (PI_ERROR_INVALID_VALUE) -30 (PI_ERROR_INVALID_VALUE)
Aborted (core dumped) 

Has anyone encountered a similar runtime error during the execution of migrated sycl code on NVIDIA as the backend? what could be the cause of the error?

[vidyalatha-badde]

Could someone help me out here?

[mmoadeli]

Hi @vidyalatha-badde
would you please be able to provide more details on how to produce the error?

[vidyalatha-badde]

Hi @mmoadeli

Thanks for looking into the issue.
I have a cuda code that got migrated and when trying to run on NVIDIA A100 GPU it is giving the above runtime error. Below are the code snippets

([&]() { dpct::get_default_queue().submit([&](sycl::handler &cgh) {
        sycl::local_accessor<half, 2> dpct_local_acc_ct1(
            sycl::range<2>(SHMEM_SZ, CHUNK_K * K + SKEW_HALF) , cgh);
        cgh.parallel_for(
            sycl::nd_range<3>(
                sycl::range<3>(1, 1, deviceProp.get_max_compute_units()) *
                    sycl::range<3>(1, 1, THREADS_PER_BLOCK),
                sycl::range<3>(1, 1, THREADS_PER_BLOCK)),
            [=](sycl::nd_item<3> item_ct1) {
            sycl::sub_group sg = item_ct1.get_sub_group();
             func1( A, B, C, D, alpha, beta, item_ct1, dpct_local_acc_ct1, sg);
            });

But in the function definition, the migrated code contains pointer to dpct_local accessor variable

void func1( const sycl::half *A, const sycl::half *B, const float *C,
                  float *D, float alpha, float beta,
                  const sycl::nd_item<3> &item_ct1, sycl::half* dpct_local, sycl::sub_group sg) {

auto shmem = (sycl::half (*)[CHUNK_K * K + SKEW_HALF])dpct_local;

Compiling the above code gives an error as

 no matching function for call to 'func1'
  263 |              func1( A, B, C, D, alpha, beta, item_ct1, dpct_local_acc_ct1, sg);
      |              ^~~~~~~~~~~~
local_acc_repro.cpp:106:6: note: candidate function not viable: no known conversion from 'const sycl::local_accessor<half, 2>' (aka 'const local_accessor<sycl::detail::half_impl::half, 2>') to 'sycl::half *' (aka 'sycl::detail::half_impl::half *') for 8th argument
  106 | void func1( const sycl::half *A, const sycl::half *B, const float *C,
      |      ^

Hence I added .get_pointer() to dpct_local_acc_ct1 as dpct_local_acc_ct1.get_pointer() during the function call and that removed the compile time error and now I'm getting the CUDA_ERROR_INVALID_VALUE.

Could you please let me know if I'm following the correct syntax for dpct_local accessor pointer here? Please let me know if you require any additional details.

Thanks,
Vidya.

[vidyalathabadde]

Hi @mmoadeli,
Any idea so far?

[mmoadeli]

Apologies @vidyalathabadde for late reply.
I haven't been able to find out the reason for the above error using the provided information. Having the source code, if possible, should help to further investigate.

[vidyalathabadde]

Hi @mmoadeli,

Thanks for getting back to me.
Please find the attached reproducer code (test-case.zip) which helps in reproducing the runtime error that I'm getting on NVIDIA hardware (A100).

To compile:
clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda test.cpp -o repro -I include/ -I Common/

To Run:
./repro

Output:

Preparing data for GPU...

UR CUDA ERROR:
        Value:           1
        Name:            CUDA_ERROR_INVALID_VALUE
        Description:     invalid argument
        Function:        urEnqueueKernelLaunch
        Source Location: sycl_workspace/llvm/sycl/plugins/unified_runtime/ur/adapters/cuda/enqueue.cpp:418

terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -30 (PI_ERROR_INVALID_VALUE) -30 (PI_ERROR_INVALID_VALUE)
Aborted (core dumped)

Clang version:

clang version 17.0.0 (https://github.com/intel/llvm c1d3d7ce2e4a3b05f70b0b2d8f6d7621ea0595aa)
Target: x86_64-unknown-linux-gnu
Thread model: posix

Could you please take a look at this issue and let me know what is the cause of the runtime error?

test-case.zip

[vidyalathabadde]

Hi @mmoadeli ,

Did you get a chance to reproduce the issue? Please let me know if it is an issue with the code or if is it a bug.

Thanks,
Vidya.

[mmoadeli]

Hi @vidyalathabadde
Sorry for late reply.
I have been able to re-produce the issue. It is going to be investigated shortly.

[abagusetty]

Apparently the issue is related to the size of local memory. The requested local memory is for this 2D array (73728, 144). So 73728 * 144 * sizeof(sycl::half) certainly exceeds q.get_device().get_info<sycl::info::device::local_mem_size>()

The error code seems not be the right one.

[vidyalathabadde]

Thanks for looking into this issue. Could you please let me know how can i resolve it?

[abagusetty]

Thanks for looking into this issue. Could you please let me know how can i resolve it?

There is nothing from the runtime/compilers that could resolve the issue since you are clearly trying to allocate more local memory than the hardware supports. Please try to look into to your original cuda code and run it on A100, it should also have similar problems unless the DPCT port is doing something wrong.

[vidyalathabadde]

The original Cuda source code is working fine and could run on A100 with the same size. So it seems to be an issue during the porting itself. Let me check with the concerned team then. Thanks for your time.

[vidyalathabadde]

The error code seems not be the right one.

Do you think the error code should be changed then? I could not interpret the error message initially hence I posted the issue here. Also, I'll check with a smaller size and get back to you if I see similar error. Will that be fine?

[abagusetty]

The error code seems not be the right one.

Do you think the error code should be changed then? I could not interpret the error message initially hence I posted the issue here. Also, I'll check with a smaller size and get back to you if I see similar error. Will that be fine?

If you set the following env variables: export UR_CUDA_MAX_LOCAL_MEM_SIZE=49000 & export SYCL_RT_WARNING_LEVEL=2

llvm/sycl/doc/EnvironmentVariables.md

Line 165 in 2744b41

| `SYCL_PI_CUDA_MAX_LOCAL_MEM_SIZE` | Integer | Specifies the maximum size of a local memory allocation in bytes. If the value exceeds the device's capabilities then a `sycl::runtime_error` is thrown. In order for the full error message to be printed, `SYCL_RT_WARNING_LEVEL=2` must be set. The default value for `SYCL_PI_CUDA_MAX_LOCAL_MEM_SIZE` is determined by the hardware. |

and run the executable again, hopefully you should see an error like the one below:

The program was built for 1 devices
Build program log for 'NVIDIA A100 80GB PCIe':

Too much local memory allocated for device
terminate called after throwing an instance of 'sycl::_V1::runtime_error'
  what():  Native API failed. Native API returns: -996 (The plugin has emitted a backend specific error)
Too much local memory allocated for device
 -996 (The plugin has emitted a backend specific error)
Aborted (core dumped)