Previously posted this in HPC section incorrectly(?) Think it belongs here instead.
I have created a small MWE to demonstrate that dynamic shared memory violations seem to not trigger any errors in compute-sanitizer.
// clang-format off
// nvcc -std=c++17 -I../include -o shared_memory_violations ../proj_slipups/shared_memory_violations.cu
// compute-sanitizer emits no errors??
#include "sharedmem.cuh"
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/sequence.h>
template <typename T> __global__ void kernel1(const T *a, T *b) {
SharedMemory<T> smem;
T *ptr = smem.getPointer();
ptr[threadIdx.x] = a[blockIdx.x * blockDim.x + threadIdx.x];
printf("ptr[%d] = %d\n", threadIdx.x, ptr[threadIdx.x]);
__syncthreads();
// Some non-trivial work to make sure not everything is optimized away?
b[blockIdx.x * blockDim.x + threadIdx.x] =
ptr[(threadIdx.x + 1) % blockDim.x];
}
__global__ void fixedkernel(const int *a, int *b){
extern __shared__ int smem[];
int *ptr = &smem[0];
ptr[threadIdx.x] = a[blockIdx.x * blockDim.x + threadIdx.x];
printf("fixedkernel: ptr[%d], %p = %d\n", threadIdx.x, &ptr[threadIdx.x], ptr[threadIdx.x]);
__syncthreads();
// Some non-trivial work to make sure not everything is optimized away?
b[blockIdx.x * blockDim.x + threadIdx.x] =
ptr[(threadIdx.x + 1) % blockDim.x];
}
int main() {
dim3 tpb(32);
dim3 bpg(1);
thrust::device_vector<int> d_a(32);
thrust::sequence(d_a.begin(), d_a.end());
thrust::device_vector<int> d_b(32);
// Don't dynamically allocate any shared mem
// kernel1<<<bpg, tpb>>>(thrust::raw_pointer_cast(d_a.data()),
// thrust::raw_pointer_cast(d_b.data()));
fixedkernel<<<bpg, tpb>>>(thrust::raw_pointer_cast(d_a.data()),
thrust::raw_pointer_cast(d_b.data()));
cudaDeviceSynchronize();
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("Error: %s\n", cudaGetErrorString(err));
}
thrust::host_vector<int> h_b = d_b;
for (int i = 0; i < 32; ++i) {
printf("d_b[%d] = %d\n", i, h_b[i]);
}
return 0;
}
At first, I thought this was a template problem with the new shared memory helper container, so I wrote fixedkernel to remove all templates.
This still does not trigger any errors from compute-sanitizer nor the cudaGetLastError in the following configurations:
- Linux, CUDA 12.3, NVIDIA L4
- Linux, CUDA 12.6, NVIDIA A10
- Linux, CUDA 11.7, NVIDIA A100
Where I did get it to trigger was the following:
- Windows 10, CUDA 12.9, NVIDIA 5080
I am not sure if I am doing something wrong, or this is OS-specifc / driver version specific / CUDA specific. The weird thing is that in the Linux runs, the output is correct as well, so it almost seems like there’s some ghost dynamic shared memory being allocated behind my back.
Edit: adding the sharedmem.cuh in case someone asks. This was copied from some NVIDIA library/git, I can’t really remember where (maybe CCCL?)
/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
* * Neither the name of NVIDIA CORPORATION nor the names of its
* contributors may be used to endorse or promote products derived
* from this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/
#ifndef _SHAREDMEM_H_
#define _SHAREDMEM_H_
//****************************************************************************
// Because dynamically sized shared memory arrays are declared "extern",
// we can't templatize them directly. To get around this, we declare a
// simple wrapper struct that will declare the extern array with a different
// name depending on the type. This avoids compiler errors about duplicate
// definitions.
//
// To use dynamically allocated shared memory in a templatized __global__ or
// __device__ function, just replace code like this:
//
//
// template<class T>
// __global__ void
// foo( T* g_idata, T* g_odata)
// {
// // Shared mem size is determined by the host app at run time
// extern __shared__ T sdata[];
// ...
// doStuff(sdata);
// ...
// }
//
// With this
// template<class T>
// __global__ void
// foo( T* g_idata, T* g_odata)
// {
// // Shared mem size is determined by the host app at run time
// SharedMemory<T> smem;
// T* sdata = smem.getPointer();
// ...
// doStuff(sdata);
// ...
// }
//****************************************************************************
// This is the un-specialized struct. Note that we prevent instantiation of
// this
// struct by putting an undefined symbol in the function body so it won't
// compile.
template <typename T> struct SharedMemory {
// Ensure that we won't compile any un-specialized types
__device__ T *getPointer() {
extern __device__ void error(void);
error();
return NULL;
}
};
// Following are the specializations for the following types.
// int, uint, char, uchar, short, ushort, long, ulong, bool, float, and double
// One could also specialize it for user-defined types.
template <> struct SharedMemory<int> {
__device__ int *getPointer() {
extern __shared__ int s_int[];
return s_int;
}
};
template <> struct SharedMemory<unsigned int> {
__device__ unsigned int *getPointer() {
extern __shared__ unsigned int s_uint[];
return s_uint;
}
};
template <> struct SharedMemory<char> {
__device__ char *getPointer() {
extern __shared__ char s_char[];
return s_char;
}
};
template <> struct SharedMemory<unsigned char> {
__device__ unsigned char *getPointer() {
extern __shared__ unsigned char s_uchar[];
return s_uchar;
}
};
template <> struct SharedMemory<short> {
__device__ short *getPointer() {
extern __shared__ short s_short[];
return s_short;
}
};
template <> struct SharedMemory<unsigned short> {
__device__ unsigned short *getPointer() {
extern __shared__ unsigned short s_ushort[];
return s_ushort;
}
};
template <> struct SharedMemory<long> {
__device__ long *getPointer() {
extern __shared__ long s_long[];
return s_long;
}
};
template <> struct SharedMemory<unsigned long> {
__device__ unsigned long *getPointer() {
extern __shared__ unsigned long s_ulong[];
return s_ulong;
}
};
template <> struct SharedMemory<bool> {
__device__ bool *getPointer() {
extern __shared__ bool s_bool[];
return s_bool;
}
};
template <> struct SharedMemory<float> {
__device__ float *getPointer() {
extern __shared__ float s_float[];
return s_float;
}
};
template <> struct SharedMemory<double> {
__device__ double *getPointer() {
extern __shared__ double s_double[];
return s_double;
}
};
template <> struct SharedMemory<unsigned long long> {
__device__ unsigned long long *getPointer() {
extern __shared__ unsigned long long s_ulonglong[];
return s_ulonglong;
}
};
template <> struct SharedMemory<char2> {
__device__ char2 *getPointer() {
extern __shared__ char2 s_char2[];
return s_char2;
}
};
template <> struct SharedMemory<short2> {
__device__ short2 *getPointer() {
extern __shared__ short2 s_short2[];
return s_short2;
}
};
template <> struct SharedMemory<int2> {
__device__ int2 *getPointer() {
extern __shared__ int2 s_int2[];
return s_int2;
}
};
#endif //_SHAREDMEM_H_