[bug] how to solve compile sycl-toolchain code occur error: '/tmp/libsycl-fallback-imf-d878ec.cubin': Permission denied

[wangzy0327]

Describe the bug

sometimes occur the following error. how to solve it ?
clang-16: warning: linked binaries do not contain expected 'nvptx64-nvidia-cuda' target; found targets: 'nvptx64-nvidia-cuda-sm_50' [-Wsycl-target]
/home/wzy/sycl_workspace/build-cuda-2022-09-debug/bin/clang-offload-bundler: error: '/tmp/libsycl-fallback-imf-d878ec.cubin': Permission denied
clang-16: error: clang-offload-bundler command failed with exit code 1 (use -v to see invocation)
CMakeFiles/hello-sycl.dir/build.make:102: recipe for target 'hello-sycl' failed
gmake[2]: *** [hello-sycl] Error 1
CMakeFiles/Makefile2:94: recipe for target 'CMakeFiles/hello-sycl.dir/all' failed
gmake[1]: *** [CMakeFiles/hello-sycl.dir/all] Error 2

To Reproduce

any sycl-example code occurs the error somtimes.

There provides a sycl example that occurs the following error when the code compile(cmake -B build && cmake --build build).

CMakeLists.txt

cmake_minimum_required(VERSION 2.8.12)

set(CMAKE_C_COMPILER "/home/wzy/sycl_workspace/build-cuda-2022-09-debug/bin/clang")
set(CMAKE_CXX_COMPILER "/home/wzy/sycl_workspace/build-cuda-2022-09-debug/bin/clang++")
set(CMAKE_CXX_STANDARD 17)

project(hello-sycl)

set(DPCPP_HOME "/home/wzy/sycl_workspace")
set(DPCPP_SYCL_HOME "${DPCPP_HOME}/build-cuda-2022-09-debug")

include_directories("${DPCPP_SYCL_HOME}/include/sycl")
include_directories("${DPCPP_SYCL_HOME}/include")

message(STATUS "dpcpp_home : ${DPCPP_HOME}")
message(STATUS "dpcpp_cuda_sycl_home : ${DPCPP_SYCL_HOME}")

message(STATUS "find library path : ${DPCPP_SYCL_HOME}/lib")
set(CMAKE_BUILD_RPATH "${DPCPP_SYCL_HOME}/lib;${CMAKE_BUILD_RPATH}")
message(STATUS "cmake build rpath : ${CMAKE_BUILD_RPATH}")

set(CMAKE_BUILD_TYPE "Debug")
set(CMAKE_CXX_FLAGS "-fsycl -fsycl-targets=nvptx64-nvidia-cuda")
set(CMAKE_CXX_FLAGS_DEBUG "$ENV{CXXFLAGS} -O0 -Wall -g -ggdb -std=c++17 -Wno-sycl-target -Wno-linker-warnings")

set(CMAKE_CXX_FLAGS_RELEASE "$ENV{CXXFLAGS} -O3 -Wall -std=c++17")

link_directories("${DPCPP_SYCL_HOME}/lib")

aux_source_directory(. DIR_SRCS)
add_executable(hello-sycl ${DIR_SRCS})
target_include_directories(hello-sycl PRIVATE "${DPCPP_SYCL_HOME}/include/sycl")
target_include_directories(hello-sycl PRIVATE "${DPCPP_SYCL_HOME}/include")
target_link_libraries(hello-sycl PRIVATE sycl)

my-example

#include <sycl/sycl.hpp>
#include <iostream>
#include <vector>
using namespace sycl;

int main(){
    std::cout<<"hello sycl"<<std::endl;
    const int N = 4;
    const int MAGIC_NUM = 42;
    std::vector<int> vec_num = {3,4,5,6,7,8,9};
    std::cout<<vec_num[0]<<std::endl;

    queue q;
    auto dev = q.get_device();
    std::cout<<"default queue select device info : "<<&dev<<std::endl;
    std::cout<<"default queue select device vendor id : "<<dev.get_info<sycl::info::device::vendor_id>()<<std::endl;
    std::cout<<"default queue select device vendor name : "<<dev.get_info<sycl::info::device::name>()<<std::endl;
    auto devices = dev.get_devices();
    for (auto &device : devices) {
            std::cout << "  queue get device: " << device.get_info<sycl::info::device::name>() << std::endl;
            std::cout << "  queue get device vendor id: " << device.get_info<sycl::info::device::vendor_id>()<< std::endl;
            std::cout << "  queue get device object address : " << &device << std::endl;                    
    }
    auto ctxt = q.get_context();
 
    auto platforms = sycl::platform::get_platforms();
    for (auto &platform : platforms) {
        std::cout << "Platform name : " << platform.get_info<sycl::info::platform::name>()<< std::endl;
        std::cout << "Platform profile : " << platform.get_info<sycl::info::platform::profile>()<<std::endl;
        std::cout << "Platform version : " << platform.get_info<sycl::info::platform::version>()<<std::endl;
        std::cout << "Platform vendor : " << platform.get_info<sycl::info::platform::vendor>()<<std::endl;
        std::cout << "Platform object address : " << &platform<< std::endl;            
        
        auto devices = platform.get_devices();
        for (auto &device : devices) {
            std::cout << "  Device: " << device.get_info<sycl::info::device::name>() << std::endl;
            std::cout << "  Device vendor id: " << device.get_info<sycl::info::device::vendor_id>()<< std::endl;
            std::cout << "  Device object address : " << &device << std::endl;                    
        }
    }

    std::vector<kernel_id> builtinKernelIds =
        dev.get_info<info::device::built_in_kernel_ids>();

    std::cout<<"built in kernel size : "<<builtinKernelIds.size()<<std::endl;

    auto myBundle = get_kernel_bundle<bundle_state::executable>(ctxt);

    return 0;
}

Environment (please complete the following information):

• OS: [Ubuntu 18.04]
• Target device and vendor: [Nvidia GPU]
• DPC++ version: sycl-2022-09
• Dependencies version: [cuda 11.2]

Additional context
Add any other context about the problem here.

[mmoadeli]

thanks @wangzy0327 for reporting this issue. We will look into this and get back to you shortly.

[ldrumm]

I've tried recreating this and I've been unsuccessful.
I recreating as much as I could of your environment in an 18.04 virtual machine with the cuda-11.2 toolkit:

mkdir 7777
cd 7777
git clone https://github.com/intel/llvm llvm-project
cat << 'EOD' > Vagrantfile
Vagrant.configure("2") do |config|
  config.vm.box = "ubuntu/bionic64"
  config.vm.synced_folder "llvm-project", "/vagrant/llvm-project"
    config.vm.provider "virtualbox" do |vb|
      vb.memory = "16384"
      vb.cpus = "8"
    end
    config.vm.provision "shell", inline: <<-SHELL
      apt-get update
      apt-get install -qy wget gnupg2 software-properties-common
      wget -q -O - https://apt.kitware.com/keys/kitware-archive-latest.asc | apt-key add - && apt-add-repository 'deb https://apt.kitware.com/ubuntu/ bionic main'
      apt-get update -qy && apt-get install -qy \
      wget \
      git \
      libncurses5-dev \
      libncursesw5-dev \
      vim \
      build-essential \
      ninja-build \
      cmake \
      python3-distutils \
      linux-headers-`uname -r`
      wget https://developer.download.nvidia.com/compute/cuda/11.2.0/local_installers/cuda_11.2.0_460.27.04_linux.run -q -O cuda.run
      sh cuda.run --silent --toolkit && rm cuda.run
    SHELL
end
EOD

cat << 'EOD' > example.cpp
#include <sycl/sycl.hpp>
#include <iostream>
#include <vector>
using namespace sycl;

int main(){
    std::cout<<"hello sycl"<<std::endl;
    const int N = 4;
    const int MAGIC_NUM = 42;
    std::vector<int> vec_num = {3,4,5,6,7,8,9};
    std::cout<<vec_num[0]<<std::endl;

    queue q;
    auto dev = q.get_device();
    std::cout<<"default queue select device info : "<<&dev<<std::endl;
    std::cout<<"default queue select device vendor id : "<<dev.get_info<sycl::info::device::vendor_id>()<<std::endl;
    std::cout<<"default queue select device vendor name : "<<dev.get_info<sycl::info::device::name>()<<std::endl;
    auto devices = dev.get_devices();
    for (auto &device : devices) {
            std::cout << "  queue get device: " << device.get_info<sycl::info::device::name>() << std::endl;
            std::cout << "  queue get device vendor id: " << device.get_info<sycl::info::device::vendor_id>()<< std::endl;
            std::cout << "  queue get device object address : " << &device << std::endl;
    }
    auto ctxt = q.get_context();

    auto platforms = sycl::platform::get_platforms();
    for (auto &platform : platforms) {
        std::cout << "Platform name : " << platform.get_info<sycl::info::platform::name>()<< std::endl;
        std::cout << "Platform profile : " << platform.get_info<sycl::info::platform::profile>()<<std::endl;
        std::cout << "Platform version : " << platform.get_info<sycl::info::platform::version>()<<std::endl;
        std::cout << "Platform vendor : " << platform.get_info<sycl::info::platform::vendor>()<<std::endl;
        std::cout << "Platform object address : " << &platform<< std::endl;

        auto devices = platform.get_devices();
        for (auto &device : devices) {
            std::cout << "  Device: " << device.get_info<sycl::info::device::name>() << std::endl;
            std::cout << "  Device vendor id: " << device.get_info<sycl::info::device::vendor_id>()<< std::endl;
            std::cout << "  Device object address : " << &device << std::endl;
        }
    }

    std::vector<kernel_id> builtinKernelIds =
        dev.get_info<info::device::built_in_kernel_ids>();

    std::cout<<"built in kernel size : "<<builtinKernelIds.size()<<std::endl;

    auto myBundle = get_kernel_bundle<bundle_state::executable>(ctxt);

    return 0;
}
EOD

cat <<'EOD' > CMakeLists.txt
cmake_minimum_required(VERSION 2.8.12)

set(DPCPP_HOME "/vagrant/llvm-project/")
set(CMAKE_C_COMPILER "${DPCPP_HOME}/build/bin/clang")
set(CMAKE_CXX_COMPILER "${DPCPP_HOME}/build/bin/clang++")
set(CMAKE_CXX_STANDARD 17)

project(hello-sycl)

set(DPCPP_SYCL_HOME "${DPCPP_HOME}/build")

include_directories("${DPCPP_SYCL_HOME}/include/sycl")
include_directories("${DPCPP_SYCL_HOME}/include")

message(STATUS "dpcpp_home : ${DPCPP_HOME}")
message(STATUS "dpcpp_cuda_sycl_home : ${DPCPP_SYCL_HOME}")

message(STATUS "find library path : ${DPCPP_SYCL_HOME}/lib")
set(CMAKE_BUILD_RPATH "${DPCPP_SYCL_HOME}/lib;${CMAKE_BUILD_RPATH}")
message(STATUS "cmake build rpath : ${CMAKE_BUILD_RPATH}")

set(CMAKE_BUILD_TYPE "Debug")
set(CMAKE_CXX_FLAGS "-fsycl -fsycl-targets=nvptx64-nvidia-cuda")
set(CMAKE_CXX_FLAGS_DEBUG "$ENV{CXXFLAGS} -O0 -Wall -g -ggdb -std=c++17 -Wno-sycl-target -Wno-linker-warnings")
# set(CMAKE_CXX_FLAGS_RELEASE "$ENV{CXXFLAGS} -O3 -Wall -std=c++17")

link_directories("${DPCPP_SYCL_HOME}/lib")

aux_source_directory(. DIR_SRCS)
add_executable(hello-sycl ${DIR_SRCS})
target_include_directories(hello-sycl PRIVATE "${DPCPP_SYCL_HOME}/include/sycl")
target_include_directories(hello-sycl PRIVATE "${DPCPP_SYCL_HOME}/include")
target_link_libraries(hello-sycl PRIVATE sycl)
EOD

cat << 'EOD' > recreate
#!/bin/bash

cd /vagrant/llvm-project
mkdir -p /vagrant/llvm-project/build
cd /vagrant/llvm-project/build
python3 ../buildbot/configure.py --cuda --cmake-opt=-DLLVM_USE_LINKER=gold --cmake-opt=-DLLVM_PARALLEL_LINK_JOBS=2 --cmake-opt=-DCUDA_CUDA_LIBRARY=/usr/local/cuda-11.2/targets/x86_64-linux/lib/stubs/libcuda.so
ninja check-sycl-deploy -k 0 # ignore some python command problems in the lit tests on 18.04
cd /vagrant
for _ in {1..100} ; do
 rm -rf build
 mkdir build
 cmake -B build
 if ! cmake --build build ; then echo FAILED ; break ; fi
done
EOD

vagrant up
vagrant ssh -c 'bash /vagrant/recreate'

I had no failures to link the example with the 100 tries looped in the last fragment listed.
My guess this permissions issue is unique to your system. I note with interest that you're using sudo to execute stat /tmp which should be unnecessary. Have you inadvertently executed the compiler with sudo at some point causing a change of ownership of something?
I'm going to close this as unreproducible, but if you can show (with data) it's really a bug in dpc++, then I'll happily investigate further

[wangzy0327]

@ldrumm Thank you for your effort. I have done a lot of network model testing with sycl programming model before by adding sycl backend to tvm. During testing, this problem did occur, but the probability is very low, and the same code is difficult to reproduce.

[wangzy0327]

@ldrumm I also met the problem in my scenary as following display. It quite exists the problem but is difficult to reproduce.

[ldrumm]

Could you please paste plain text rather than screenshots? It's hard to extract information from a screenshot

[wangzy0327]

@ldrumm @Pennycook Can you try to reproduce the problem in my same environment?

OS System environment

DISTRIB_ID=Ubuntu
DISTRIB_RELEASE=18.04
DISTRIB_CODENAME=bionic
DISTRIB_DESCRIPTION="Ubuntu 18.04.6 LTS"
NAME="Ubuntu"
VERSION="18.04.6 LTS (Bionic Beaver)"
ID=ubuntu
ID_LIKE=debian
PRETTY_NAME="Ubuntu 18.04.6 LTS"
VERSION_ID="18.04"
HOME_URL="https://www.ubuntu.com/"
SUPPORT_URL="https://help.ubuntu.com/"
BUG_REPORT_URL="https://bugs.launchpad.net/ubuntu/"
PRIVACY_POLICY_URL="https://www.ubuntu.com/legal/terms-and-policies/privacy-policy"
VERSION_CODENAME=bionic
UBUNTU_CODENAME=bionic

OS kernel version

Linux gxnzx1277 4.15.0-188-generic #199-Ubuntu SMP Wed Jun 15 20:42:56 UTC 2022 x86_64 x86_64 x86_64 GNU/Linux

cuda version

Tue Mar  7 02:13:46 2023       
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 460.27.04    Driver Version: 460.27.04    CUDA Version: 11.2     |

GPU 0: Tesla V100-PCIE-32GB (UUID: GPU-d4511668-7898-0d08-1441-4d81bbe88f68)
GPU 1: Tesla V100-PCIE-32GB (UUID: GPU-3d5ecc1e-153e-3012-1fa5-d46138f948bc)

Sycl version: 2022-12-release

reproduce example:

CMakeLists.txt

cmake_minimum_required(VERSION 2.8.12)

set(DPCPP_HOME "/home/wzy/sycl_workspace")
set(DPCPP_SYCL_HOME "${DPCPP_HOME}/build-cuda-2022-12")

set(CMAKE_C_COMPILER "${DPCPP_SYCL_HOME}/bin/clang")
set(CMAKE_CXX_COMPILER "${DPCPP_SYCL_HOME}/bin/clang++")
set(CMAKE_CXX_STANDARD 17)

project(complex-example)

include_directories("${DPCPP_SYCL_HOME}/include/sycl")
include_directories("${DPCPP_SYCL_HOME}/include")

message(STATUS "dpcpp_home : ${DPCPP_HOME}")
message(STATUS "dpcpp_cuda_sycl_home : ${DPCPP_SYCL_HOME}")


message(STATUS "find library path : ${DPCPP_SYCL_HOME}/lib")
set(CMAKE_BUILD_RPATH "${DPCPP_SYCL_HOME}/lib;${CMAKE_BUILD_RPATH}")
message(STATUS "cmake build rpath : ${CMAKE_BUILD_RPATH}")


set(CMAKE_BUILD_TYPE "Debug")
set(CMAKE_CXX_FLAGS "-fsycl -fsycl-targets=nvptx64-nvidia-cuda")
set(CMAKE_CXX_FLAGS_DEBUG "$ENV{CXXFLAGS} -O0 -Wall -g -ggdb -std=c++17")
# set(CMAKE_CXX_FLAGS_RELEASE "$ENV{CXXFLAGS} -O3 -Wall -std=c++17")


link_directories("${DPCPP_SYCL_HOME}/lib")

aux_source_directory(. DIR_SRCS)
add_executable(complex-example ${DIR_SRCS})
target_include_directories(complex-example PRIVATE "${DPCPP_SYCL_HOME}/include/sycl")
target_include_directories(complex-example PRIVATE "${DPCPP_SYCL_HOME}/include")
target_link_libraries(complex-example PRIVATE sycl)

complex.cc

#include <iostream>
#include <sycl/sycl.hpp>
using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names

// Size of the matrices
constexpr size_t N = 2000;
constexpr size_t M = 3000;

class MyKernelA;
class MyKernelB;
class MyKernelC;

int main() {
  // Create a queue to work on
  queue myQueue;

  // Create some 2D buffers of float for our matrices
  buffer<float, 2> a { range<2> { N, M } };
  buffer<float, 2> b { range<2> { N, M } };
  buffer<float, 2> c { range<2> { N, M } };

  // Launch an asynchronous kernel to initialize a
  myQueue.submit([&](handler& cgh) {
    // The kernel writes a, so get a write accessor on it
    accessor A { a, cgh, write_only };

    // Enqueue a parallel kernel iterating on a N*M 2D iteration space
    cgh.parallel_for<MyKernelA>(range<2> { N, M },
                     [=](id<2> index) { A[index] = index[0] * 2 + index[1]; });
  });

  // Launch an asynchronous kernel to initialize b
  myQueue.submit([&](handler& cgh) {
    // The kernel writes b, so get a write accessor on it
    accessor B { b, cgh, write_only };

    // From the access pattern above, the SYCL runtime detects that this
    // command_group is independent from the first one and can be
    // scheduled independently

    // Enqueue a parallel kernel iterating on a N*M 2D iteration space
    cgh.parallel_for<MyKernelB>(range<2> { N, M }, [=](id<2> index) {
      B[index] = index[0] * 2014 + index[1] * 42;
    });
  });

  // Launch an asynchronous kernel to compute matrix addition c = a + b
  myQueue.submit([&](handler& cgh) {
    // In the kernel a and b are read, but c is written
    accessor A { a, cgh, read_only };
    accessor B { b, cgh, read_only };
    accessor C { c, cgh, write_only };

    // From these accessors, the SYCL runtime will ensure that when
    // this kernel is run, the kernels computing a and b have completed

    // Enqueue a parallel kernel iterating on a N*M 2D iteration space
    cgh.parallel_for<MyKernelC>(range<2> { N, M },
                     [=](id<2> index) { C[index] = A[index] + B[index]; });
  });

  auto dev = myQueue.get_device();
  const std::vector<kernel_id> builtinKernelIds =
        dev.get_info<info::device::built_in_kernel_ids>();     

  std::cout<<"built in kernel size : "<<builtinKernelIds.size()<<std::endl;

  // Ask for an accessor to read c from application scope.  The SYCL runtime
  // waits for c to be ready before returning from the constructor
  host_accessor C { c, read_only };
  std::cout << std::endl << "Result:" << std::endl;
  for (size_t i = 0; i < N; i++) {
    for (size_t j = 0; j < M; j++) {
      // Compare the result to the analytic value
      if (C[i][j] != i * (2 + 2014) + j * (1 + 42)) {
        std::cout << "Wrong value " << C[i][j] << " on element " << i << " "
                  << j << std::endl;
        exit(-1);
      }
    }
  }

  std::cout << "Good computation!" << std::endl;
  return 0;
}

test.sh

for _ in {1..100} ; do
    rm -rf build
    mkdir build
    cmake -B build
    if ! cmake --build build ; then echo FAILED ; break ; fi
done
EOD    

I get the result which display /home/wzy/sycl_workspace/build-cuda-2022-12/bin/clang-offload-bundler: error: '/tmp/libsycl-cmath-7e8fa3.cubin': Permission denied

-- The C compiler identification is Clang 16.0.0
-- The CXX compiler identification is Clang 16.0.0
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Check for working C compiler: /home/wzy/sycl_workspace/build-cuda-2022-12/bin/clang - skipped
-- Detecting C compile features
-- Detecting C compile features - done
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Check for working CXX compiler: /home/wzy/sycl_workspace/build-cuda-2022-12/bin/clang++ - skipped
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- dpcpp_home : /home/wzy/sycl_workspace
-- dpcpp_cuda_sycl_home : /home/wzy/sycl_workspace/build-cuda-2022-12
-- find library path : /home/wzy/sycl_workspace/build-cuda-2022-12/lib
-- cmake build rpath : /home/wzy/sycl_workspace/build-cuda-2022-12/lib;
-- Configuring done
-- Generating done
-- Build files have been written to: /home/wzy/sycl_workspace/dpcpp-application/complex-example/build
Scanning dependencies of target complex-example
[ 50%] Building CXX object CMakeFiles/complex-example.dir/complex.cc.o
[100%] Linking CXX executable complex-example
/home/wzy/sycl_workspace/build-cuda-2022-12/bin/clang-offload-bundler: error: '/tmp/libsycl-cmath-7e8fa3.cubin': Permission denied
clang-16: error: clang-offload-bundler command failed with exit code 1 (use -v to see invocation)
CMakeFiles/complex-example.dir/build.make:102: recipe for target 'complex-example' failed
gmake[2]: *** [complex-example] Error 1
CMakeFiles/Makefile2:94: recipe for target 'CMakeFiles/complex-example.dir/all' failed
gmake[1]: *** [CMakeFiles/complex-example.dir/all] Error 2
Makefile:102: recipe for target 'all' failed
gmake: *** [all] Error 2
FAILED
test.sh: line 7: EOD: command not found

[ldrumm]

The only difference is my kernel version is

Linux ubuntu-bionic 4.15.0-206-generic #217-Ubuntu SMP Fri Feb 3 19:10:13 UTC 2023 x86_64 x86_64 x86_64 GNU/Linux

I reran using the binary release you posted (rather than building from source as before) and still can't recreate the bug
I had a thought that space might be an issue so artificially limited the amount of space available in /tmp.

[100%] Linking CXX executable complex-example
LLVM ERROR: IO failure on output stream: No space left on device
clang++: error: sycl-link command failed with exit code 1 (use -v to see invocation)
CMakeFiles/complex-example.dir/build.make:96: recipe for target 'complex-example' failed
make[2]: *** [complex-example] Error 1
CMakeFiles/Makefile2:82: recipe for target 'CMakeFiles/complex-example.dir/all' failed
make[1]: *** [CMakeFiles/complex-example.dir/all] Error 2
Makefile:90: recipe for target 'all' failed
make: *** [all] Error 2
FAILED

However, that's still not a permissions issue. I don't think there's anything I can do to help further, but ask the following, in order of simplicity

• Is some other process modifying permissions in /tmp (anitvirus, maybe)?
• How reproducible (how often) is this?
• can you wrap the compiler command in strace -f -e %file so we can get an idea of how the syscalls are failing?
• Does it happen with other release versions?
• Can you build from source and reproduce in a debugger?

Without being able to reproduce locally, I'm unable to help. I suggest you follow the above steps to determine where the issue lies.

[wangzy0327]

@ldrumm

(base) wzy@gxnzx1277:~/sycl_workspace$ df -lh /tmp
Filesystem      Size  Used Avail Use% Mounted on
/dev/sda5      1007G   46G  911G   5% /

There is enough amount of space avaliable in /tmp.
how to use the command in strace -f -e %file ? Is there a case study?
It also happen with sycl-2022-09 release version besides sycl-2022-12-release