Kernel within a kernel: out of bounds

franck.richard · July 31, 2024, 7:06pm

Hello,

I am currently working on a genetic sequencing program but I lack CUDA experience.
I have a kernel that contains the following loop:

#pragma unroll 1
                                    for (score_type k = lo; k <= hi; ++k)
                                    {
                                        score_type I1, I2, I, D1, D2, D, H1, H;

                                        I1 = (E_Band_gap_extend1_ptr && k - 1 >= E_lo && k - 1 <= E_hi) ? E_Band_gap_extend1_ptr[k - 1] : WFA_MIN1;
                                        I2 = (H_Band_gap_open1_ptr && k - 1 >= H_lo_o && k - 1 <= H_hi_o) ? H_Band_gap_open1_ptr[k - 1] : WFA_MIN1;
                                        I = nvbio::max(I1, I2) + 1u;
                                        E_Band_ptr[k] = I;

                                        D1 = (F_Band_gap_extend1_ptr && k + 1 >= F_lo && k + 1 <= F_hi) ? F_Band_gap_extend1_ptr[k + 1] : WFA_MIN1;
                                        D2 = (H_Band_gap_open1_ptr && k + 1 >= H_lo_o && k + 1 <= H_hi_o) ? H_Band_gap_open1_ptr[k + 1] : WFA_MIN1;
                                        D = nvbio::max(D1, D2);
                                        F_Band_ptr[k] = D;

                                        H1 = (H_Band_gap_mismatch_ptr && k >= H_lo_m && k <= H_hi_m) ? H_Band_gap_mismatch_ptr[k] : WFA_MIN1;
                                        H = nvbio::max3(H1 + 1, I, D);

                                        if ((uint16)H > M || (uint16)(H - k) > N)
                                            H = WAVEFRONT_OFFSET_NULL;

                                        H_Band_ptr[k] = H;
                                    }

I’d like to be able to launch a second kernel (from the first) which would execute this loop in parallel for all values of k between lo and hi.

So I defined a second kernel this way



template <typename score_type>
        __global__ void compute_band_kernel(score_type M, score_type N, score_type lo, score_type hi, score_type E_lo, score_type E_hi, score_type H_lo_o, score_type H_hi_o, score_type F_lo, score_type F_hi, score_type H_lo_m, score_type H_hi_m,
                                            const int16 *E_Band_gap_extend1_ptr, const int16 *H_Band_gap_open1_ptr, const int16 *F_Band_gap_extend1_ptr, const int16 *H_Band_gap_mismatch_ptr,
                                            int16 *E_Band_ptr, int16 *F_Band_ptr, int16 *H_Band_ptr)
        {
            int32 k = lo + blockIdx.x * blockDim.x + threadIdx.x;

            if (k >= lo && k <= hi)
            {
                score_type I1, I2, I, D1, D2, D, H1, H;

                I1 = (E_Band_gap_extend1_ptr && k - 1 >= E_lo && k - 1 <= E_hi) ? E_Band_gap_extend1_ptr[k - 1] : WFA_MIN1;
                I2 = (H_Band_gap_open1_ptr && k - 1 >= H_lo_o && k - 1 <= H_hi_o) ? H_Band_gap_open1_ptr[k - 1] : WFA_MIN1;
                I = nvbio::max(I1, I2) + 1;
                E_Band_ptr[k] = I;

                D1 = (F_Band_gap_extend1_ptr && k + 1 >= F_lo && k + 1 <= F_hi) ? F_Band_gap_extend1_ptr[k + 1] : WFA_MIN1;
                D2 = (H_Band_gap_open1_ptr && k + 1 >= H_lo_o && k + 1 <= H_hi_o) ? H_Band_gap_open1_ptr[k + 1] : WFA_MIN1;
                D = nvbio::max(D1, D2);
                F_Band_ptr[k] = D;

                H1 = (H_Band_gap_mismatch_ptr && k >= H_lo_m && k <= H_hi_m) ? H_Band_gap_mismatch_ptr[k] : WFA_MIN1;
                H = nvbio::max3(H1 + 1, I, D);

                if ((unsigned short)H > M || (unsigned short)(H - k) > N)
                    H = WAVEFRONT_OFFSET_NULL;

                H_Band_ptr[k] = H;
            }
        }

        // Fonction pour vérifier les erreurs CUDA
        inline void checkCudaErrors(cudaError_t err, const char *msg)
        {
            if (err != cudaSuccess)
            {
                std::cerr << msg << ": " << cudaGetErrorString(err) << std::endl;
                throw std::runtime_error(msg);
            }
        }

        template <typename score_type>
        void compute_band_host(score_type M, score_type N, score_type lo, score_type hi, score_type E_lo, score_type E_hi, score_type H_lo_o, score_type H_hi_o, score_type F_lo, score_type F_hi, score_type H_lo_m, score_type H_hi_m,
                               const int16 *E_Band_gap_extend1_ptr, const int16 *H_Band_gap_open1_ptr, const int16 *F_Band_gap_extend1_ptr, const int16 *H_Band_gap_mismatch_ptr,
                               int16 *E_Band_ptr, int16 *F_Band_ptr, int16 *H_Band_ptr)
        {
            score_type num_threads = hi - lo + 1;
            score_type block_size = 256;
            score_type num_blocks = (num_threads + block_size - 1) / block_size;

            compute_band_kernel<<<num_blocks, block_size>>>(M, N, lo, hi, E_lo, E_hi, H_lo_o, H_hi_o, F_lo, F_hi, H_lo_m, H_hi_m,
                                                            E_Band_gap_extend1_ptr, H_Band_gap_open1_ptr, F_Band_gap_extend1_ptr, H_Band_gap_mismatch_ptr,
                                                            E_Band_ptr, F_Band_ptr, H_Band_ptr);

            checkCudaErrors(cudaDeviceSynchronize(), "Erreur lors de cudaDeviceSynchronize");
        }

        template <typename score_type>
        NVBIO_FORCEINLINE NVBIO_HOST_DEVICE void compute_band(score_type M, score_type N, score_type lo, score_type hi, score_type E_lo, score_type E_hi, score_type H_lo_o, score_type H_hi_o, score_type F_lo, score_type F_hi, score_type H_lo_m, score_type H_hi_m,
                                                              const int16 *E_Band_gap_extend1_ptr, const int16 *H_Band_gap_open1_ptr, const int16 *F_Band_gap_extend1_ptr, const int16 *H_Band_gap_mismatch_ptr,
                                                              int16 *E_Band_ptr, int16 *F_Band_ptr, int16 *H_Band_ptr)
        {
#ifndef __CUDA_ARCH__
            // Contexte hôte
            compute_band_host(M, N, lo, hi, E_lo, E_hi, H_lo_o, H_hi_o, F_lo, F_hi, H_lo_m, H_hi_m, E_Band_gap_extend1_ptr, H_Band_gap_open1_ptr, F_Band_gap_extend1_ptr, H_Band_gap_mismatch_ptr, E_Band_ptr, F_Band_ptr, H_Band_ptr);
#else
            dim3 threadsPerBlock(16, 16);
            dim3 numBlocks((1 + threadsPerBlock.x - 1) / threadsPerBlock.x,
                           (1 + threadsPerBlock.y - 1) / threadsPerBlock.y);

            score_type num_threads = hi - lo + 1;
            score_type block_size = 5;
            score_type num_blocks = (num_threads + block_size - 1) / block_size;

            // Contexte device : appels de kernel ici
            // compute_band_kernel<<<(hi - lo + 256) / 256, 256>>>(M, N, lo, hi, E_lo, E_hi, H_lo_o, H_hi_o, F_lo, F_hi, H_lo_m, H_hi_m, E_Band_gap_extend1_ptr, H_Band_gap_open1_ptr, F_Band_gap_extend1_ptr, H_Band_gap_mismatch_ptr, E_Band_ptr, F_Band_ptr, H_Band_ptr);
            // compute_band_kernel<<<numBlocks, threadsPerBlock>>>(M, N, lo, hi, E_lo, E_hi, H_lo_o, H_hi_o, F_lo, F_hi, H_lo_m, H_hi_m, E_Band_gap_extend1_ptr, H_Band_gap_open1_ptr, F_Band_gap_extend1_ptr, H_Band_gap_mismatch_ptr, E_Band_ptr, F_Band_ptr, H_Band_ptr);
            compute_band_kernel<<<num_blocks, block_size>>>(M, N, lo, hi, E_lo, E_hi, H_lo_o, H_hi_o, F_lo, F_hi, H_lo_m, H_hi_m, E_Band_gap_extend1_ptr, H_Band_gap_open1_ptr, F_Band_gap_extend1_ptr, H_Band_gap_mismatch_ptr, E_Band_ptr, F_Band_ptr, H_Band_ptr);
#endif
        }

Problem: I have a crash at runtime when modifying variables E_Band_ptr, F_Band_ptr and H_Band_ptr in the second kernel…

Memcheck gives me :



========= Invalid __global__ write of size 2 bytes
=========     at void nvbio::aln::compute_band_kernel<int, nvbio::aln::wfa_type<int>>(T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T1, T2)+0x1d80 in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio/alignment/wfa.h:59
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x57f149e65ee2 is out of bounds
=========     and is 40.843.114.815.774 bytes before the nearest allocation at 0x7d16d1400000 of size 150 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x2c914f]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:libcudart_static_4d8b33a106dceb3c07a56e26de61f2d53bb62a68 [0x2696cd]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:cudaLaunchKernel [0x2cd13d]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) in /usr/local/cuda-12.5/include/cuda_runtime.h:216 [0x494af]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__device_stub__ZN5nvbio3aln19compute_band_kernelIiNS0_8wfa_typeIiEEEEvT_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_T0_(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) in /tmp/tmpxft_0009c696_00000000-6_alignment_test.cudafe1.stub.c:287 [0x4338f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::__wrapper__device_stub_compute_band_kernel<int, nvbio::aln::wfa_type<int> >(int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, nvbio::aln::wfa_type<int>&) in /tmp/tmpxft_0009c696_00000000-6_alignment_test.cudafe1.stub.c:290 [0x4346e]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band_kernel<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>) [0x5796b]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band_host<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) [0x8eb19]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) [0x88d63]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:int nvbio::aln::priv::wfah_alignment_score_dispatch<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag, unsigned char>::update<true, nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>, uchar2*, int, unsigned char*, nvbio::aln::BestSink<int>, nvbio::aln::SimpleWfahScheme, nvbio::aln::wfa_type<int> >(nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>&, int, int, int, int, unsigned char*, uchar2*, int&, int*, int*, nvbio::aln::BestSink<int>&, int, int&, int, int, int, int, int, nvbio::aln::SimpleWfahScheme const&, nvbio::aln::wfa_type<int>&) [0x7ed3f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::priv::wfah_alignment_score_dispatch<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag, unsigned char>::run<nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::SimpleWfahScheme, nvbio::aln::BestSink<int>, short2*, nvbio::aln::wfa_type<int> >(nvbio::aln::SimpleWfahScheme const&, nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>&, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, int, int, short2*, nvbio::aln::wfa_type<int>&) [0x7711e]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::priv::alignment_score_dispatch<nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, short2*>::dispatch<nvbio::aln::BestSink<int>, nvbio::aln::wfa_type<int> >(nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, short2*, nvbio::aln::wfa_type<int>&) [0x71612]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::alignment_score<nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::BestSink<int>, short2*, nvbio::aln::wfa_type<int> >(nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, short2*, nvbio::aln::wfa_type<int>&) [0x5c73f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::SingleTest::full_wfa<128u, 181u, 150u, nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag> >(char const*, nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, char const*) [0x4ef60]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::test(int, char**) in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/alignment_test.cu:1827 [0x3e7bc]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:main in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/nvbio-test.cpp:207 [0x3031d]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x2a1c9]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main in ../csu/libc-start.c:360 [0x2a28a]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x13a54]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
========= 
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaLaunchKernel.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x404b45]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaLaunchKernel [0x2cd177]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) in /usr/local/cuda-12.5/include/cuda_runtime.h:216 [0x494af]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__device_stub__ZN5nvbio3aln19compute_band_kernelIiNS0_8wfa_typeIiEEEEvT_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_S4_T0_(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) in /tmp/tmpxft_0009c696_00000000-6_alignment_test.cudafe1.stub.c:287 [0x4338f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::__wrapper__device_stub_compute_band_kernel<int, nvbio::aln::wfa_type<int> >(int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, int&, nvbio::aln::wfa_type<int>&) in /tmp/tmpxft_0009c696_00000000-6_alignment_test.cudafe1.stub.c:290 [0x4346e]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band_kernel<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>) [0x5796b]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band_host<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) [0x8eb19]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) [0x88d63]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:int nvbio::aln::priv::wfah_alignment_score_dispatch<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag, unsigned char>::update<true, nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>, uchar2*, int, unsigned char*, nvbio::aln::BestSink<int>, nvbio::aln::SimpleWfahScheme, nvbio::aln::wfa_type<int> >(nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>&, int, int, int, int, unsigned char*, uchar2*, int&, int*, int*, nvbio::aln::BestSink<int>&, int, int&, int, int, int, int, int, nvbio::aln::SimpleWfahScheme const&, nvbio::aln::wfa_type<int>&) [0x7ed3f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::priv::wfah_alignment_score_dispatch<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag, unsigned char>::run<nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::SimpleWfahScheme, nvbio::aln::BestSink<int>, short2*, nvbio::aln::wfa_type<int> >(nvbio::aln::SimpleWfahScheme const&, nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>&, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, int, int, short2*, nvbio::aln::wfa_type<int>&) [0x7711e]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::priv::alignment_score_dispatch<nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, short2*>::dispatch<nvbio::aln::BestSink<int>, nvbio::aln::wfa_type<int> >(nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, short2*, nvbio::aln::wfa_type<int>&) [0x71612]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::alignment_score<nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::BestSink<int>, short2*, nvbio::aln::wfa_type<int> >(nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, short2*, nvbio::aln::wfa_type<int>&) [0x5c73f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::SingleTest::full_wfa<128u, 181u, 150u, nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag> >(char const*, nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, char const*) [0x4ef60]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::test(int, char**) in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/alignment_test.cu:1827 [0x3e7bc]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:main in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/nvbio-test.cpp:207 [0x3031d]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x2a1c9]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main in ../csu/libc-start.c:360 [0x2a28a]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x13a54]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
========= 
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x404b45]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaDeviceSynchronize [0x2a2f73]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band_host<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) [0x8eb25]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::compute_band<int, nvbio::aln::wfa_type<int> >(int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, nvbio::aln::wfa_type<int>&) [0x88d63]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:int nvbio::aln::priv::wfah_alignment_score_dispatch<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag, unsigned char>::update<true, nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>, uchar2*, int, unsigned char*, nvbio::aln::BestSink<int>, nvbio::aln::SimpleWfahScheme, nvbio::aln::wfa_type<int> >(nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>&, int, int, int, int, unsigned char*, uchar2*, int&, int*, int*, nvbio::aln::BestSink<int>&, int, int&, int, int, int, int, int, nvbio::aln::SimpleWfahScheme const&, nvbio::aln::wfa_type<int>&) [0x7ed3f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::priv::wfah_alignment_score_dispatch<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag, unsigned char>::run<nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::SimpleWfahScheme, nvbio::aln::BestSink<int>, short2*, nvbio::aln::wfa_type<int> >(nvbio::aln::SimpleWfahScheme const&, nvbio::aln::priv::WfahScoringContext<1u, (nvbio::aln::AlignmentType)2, nvbio::aln::PatternBlockingTag>&, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, int, int, short2*, nvbio::aln::wfa_type<int>&) [0x7711e]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::priv::alignment_score_dispatch<nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, short2*>::dispatch<nvbio::aln::BestSink<int>, nvbio::aln::wfa_type<int> >(nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, short2*, nvbio::aln::wfa_type<int>&) [0x71612]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:bool nvbio::aln::alignment_score<nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::BestSink<int>, short2*, nvbio::aln::wfa_type<int> >(nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, nvbio::vector_view<unsigned char const*, unsigned int>, nvbio::aln::trivial_quality_string, nvbio::vector_view<unsigned char const*, unsigned int>, int, nvbio::aln::BestSink<int>&, short2*, nvbio::aln::wfa_type<int>&) [0x5c73f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:void nvbio::aln::SingleTest::full_wfa<128u, 181u, 150u, nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag> >(char const*, nvbio::aln::WfahAligner<(nvbio::aln::AlignmentType)2, nvbio::aln::SimpleWfahScheme, nvbio::aln::PatternBlockingTag>, char const*) [0x4ef60]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::test(int, char**) in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/alignment_test.cu:1827 [0x3e7bc]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:main in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/nvbio-test.cpp:207 [0x3031d]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x2a1c9]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main in ../csu/libc-start.c:360 [0x2a28a]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x13a54]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
========= 
Erreur lors de cudaDeviceSynchronize: unspecified launch failure
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaFree.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x404b45]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaFree [0x2af957]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::system::cuda::detail::cuda_memory_resource<&cudaMalloc, &cudaFree, thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default> >::do_deallocate(thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default>, unsigned long, unsigned long) [0x8dd2f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::device_ptr_memory_resource<thrust::THRUST_200400_860_NS::system::cuda::detail::cuda_memory_resource<&cudaMalloc, &cudaFree, thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default> > >::do_deallocate(thrust::THRUST_200400_860_NS::device_ptr<void>, unsigned long, unsigned long) [0x8795c]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::mr::allocator<unsigned char, thrust::THRUST_200400_860_NS::device_ptr_memory_resource<thrust::THRUST_200400_860_NS::system::cuda::detail::cuda_memory_resource<&cudaMalloc, &cudaFree, thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default> > > >::deallocate(thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long) [0x87bbd]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::allocator_traits<thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::deallocate(thrust::THRUST_200400_860_NS::device_allocator<unsigned char>&, thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long)::workaround_warnings::deallocate(thrust::THRUST_200400_860_NS::device_allocator<unsigned char>&, thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long) [0x7d3ca]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::allocator_traits<thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::deallocate(thrust::THRUST_200400_860_NS::device_allocator<unsigned char>&, thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long) [0x7d3fc]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::contiguous_storage<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::deallocate() [0x760a0]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::contiguous_storage<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::~contiguous_storage() [0x702ad]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::vector_base<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::~vector_base() [0x58b7c]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::device_vector<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::~device_vector() [0x4d7e7]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::SingleTest::~SingleTest() [0x4cc05]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::test(int, char**) in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/alignment_test.cu:1834 [0x408a8]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:main in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/nvbio-test.cpp:207 [0x3031d]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x2a1c9]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main in ../csu/libc-start.c:360 [0x2a28a]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x13a54]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
========= 
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaGetLastError.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x404b45]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaGetLastError [0x2a5cd0]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::cuda_cub::throw_on_error(cudaError, char const*) [0x4c35f]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::system::cuda::detail::cuda_memory_resource<&cudaMalloc, &cudaFree, thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default> >::do_deallocate(thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default>, unsigned long, unsigned long) [0x8dd4c]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::device_ptr_memory_resource<thrust::THRUST_200400_860_NS::system::cuda::detail::cuda_memory_resource<&cudaMalloc, &cudaFree, thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default> > >::do_deallocate(thrust::THRUST_200400_860_NS::device_ptr<void>, unsigned long, unsigned long) [0x8795c]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::mr::allocator<unsigned char, thrust::THRUST_200400_860_NS::device_ptr_memory_resource<thrust::THRUST_200400_860_NS::system::cuda::detail::cuda_memory_resource<&cudaMalloc, &cudaFree, thrust::THRUST_200400_860_NS::pointer<void, thrust::THRUST_200400_860_NS::cuda_cub::tag, thrust::THRUST_200400_860_NS::tagged_reference<void, thrust::THRUST_200400_860_NS::cuda_cub::tag>, thrust::THRUST_200400_860_NS::use_default> > > >::deallocate(thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long) [0x87bbd]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::allocator_traits<thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::deallocate(thrust::THRUST_200400_860_NS::device_allocator<unsigned char>&, thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long)::workaround_warnings::deallocate(thrust::THRUST_200400_860_NS::device_allocator<unsigned char>&, thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long) [0x7d3ca]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::allocator_traits<thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::deallocate(thrust::THRUST_200400_860_NS::device_allocator<unsigned char>&, thrust::THRUST_200400_860_NS::device_ptr<unsigned char>, unsigned long) [0x7d3fc]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::contiguous_storage<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::deallocate() [0x760a0]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::contiguous_storage<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::~contiguous_storage() [0x702ad]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::detail::vector_base<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::~vector_base() [0x58b7c]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:thrust::THRUST_200400_860_NS::device_vector<unsigned char, thrust::THRUST_200400_860_NS::device_allocator<unsigned char> >::~device_vector() [0x4d7e7]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::SingleTest::~SingleTest() [0x4cc05]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:nvbio::aln::test(int, char**) in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/alignment_test.cu:1834 [0x408a8]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:main in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/nvbio-test/nvbio-test.cpp:207 [0x3031d]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
=========     Host Frame:__libc_start_call_main in ../sysdeps/nptl/libc_start_call_main.h:58 [0x2a1c9]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main in ../csu/libc-start.c:360 [0x2a28a]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x13a54]
=========                in /home/karlhal/bioinformatic2/code/RbowtieCuda/src/nvbio/build_deb/./nvbio-test/nvbio-test
========= 
terminate called after throwing an instance of 'thrust::THRUST_200400_860_NS::system::system_error'
  what():  CUDA free failed: cudaErrorLaunchFailure: unspecified launch failure
========= Error: process didn't terminate successfully
========= Target application returned an error
========= ERROR SUMMARY: 5 errors

Why is there such a difference in the addresses (out of bounds) defining the three variables mentioned above?

With my thanks,

Franck

Robert_Crovella · July 31, 2024, 8:39pm

If E_band_ptr refers to a location in local or shared memory, it cannot be passed to a child kernel for use in the child kernel. I don’t know for sure that is the case since you’ve provided incomplete code. Also, cuda-memcheck if that is what you are referring to as “Memcheck” can identify the specific offending line of code for you, if you compile with -lineinfo.

In that case, once you know the offending line, you can add printf statements to print out pointer values and indexes/offsets, and spot the issue yourself.

franck.richard · August 1, 2024, 10:14am

Thank you Robert for your clarification about memory. I use Thrust to reserve memory space for variables. I will study this.

Topic		Replies	Views
Unknow error when calling device subroutine Legacy PGI Compilers (archived)	18	10641	September 1, 2015
Using cudaMemCheck Legacy PGI Compilers (archived)	3	6580	November 21, 2013
Memcheck CUDA Programming and Performance	2	614	July 20, 2017
Weird pointer arithmetic bug ? CUDA Programming and Performance	18	2651	July 12, 2011
Invalid __global__ write of size 4 Error on kernel Launch CUDA Programming and Performance	7	6271	March 18, 2013
copyin 2D failed Legacy PGI Compilers (archived)	24	15243	August 18, 2015
strange error in summation memory problems CUDA Programming and Performance	14	16843	July 21, 2010
How to deal with "invalid argument" errors when launching kernels with arguments pointing to large data? CUDA Programming and Performance	8	25171	March 1, 2019
Problem with read access violation for large arrays in unified memory CUDA Programming and Performance	16	3425	July 16, 2018
kernel works on Gtx280/295/480 but not on C2050 unspecified launch failure CUDA Programming and Performance	38	3336	September 23, 2010

Kernel within a kernel: out of bounds

Related topics