Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

CUDA compilation error: dot and vector multiplication functions use STL and cause compilation error. #1285

Open
LightShade12 opened this issue Apr 8, 2024 · 0 comments

Comments

@LightShade12
Copy link

GLM: v1.0.1

Visual Studio 2022

Windows 10, GTX 1650

CUDA: v12.0

Including glm.hpp inside .cu file with #define GLM_FORCE_CUDA defined before and compiling with nvcc causes compilation error:
glm\glm\detail\_vectorize.hpp(74): error : calling a __host__ function from a __host__ __device__ function is not allowed

inclusion snippet:

#include <cstdint>
#define __CUDACC__
#include <cuda.h>
#define GLM_FORCE_CUDA
#include <glm/glm.hpp>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <surface_indirect_functions.h>

I am trying to call glm::dot from inside a kernel:

__global__ void kernel(cudaSurfaceObject_t _surfobj, int max_x, int max_y)
{
	int i = threadIdx.x + blockIdx.x * blockDim.x;
	int j = threadIdx.y + blockIdx.y * blockDim.y;
	if ((i >= max_x) || (j >= max_y)) return;

	glm::vec2 uv = { (float(i) / max_x) ,(float(j) / max_y) };

	uv.x *= ((float)max_x / (float)max_y);
	uv.x = uv.x * 2.f - ((float)max_x / (float)max_y);
	uv.y = uv.y * 2.f - 1.f;

	glm::vec3 rayDirection(uv.x, uv.y, -1.0f);
	glm::vec3 rayOrigin(0.0f, 0.0f, 2.0f);
	uchar4 color = { 0,0,0,255 };
	float radius = 0.5f;

	float a = dot(rayDirection, rayDirection);
	float b = 2.0f * dot(rayOrigin, rayDirection);
	float c = dot(rayOrigin, rayOrigin) - radius * radius;

	float discriminant = b * b - 4.0f * a * c;
	if (discriminant < 0.0f)
	{
		surf2Dwrite(color, _surfobj, i * 4, j);
		return;
	}
	color={255,0,0,255};
	surf2Dwrite(color, _surfobj, i * 4, j);
}

If disabling the warning,
Calling glm::dot() or vec3 * vec3 operations inside kernel causes cudaError=700 IllegalAccessError with Nsight debugger pointing to ln74, _vectorize.hpp :

GLM_FUNC_QUALIFIER GLM_CONSTEXPR static vec<2, T, Q> call(Fct Func, vec<2, T, Q> const& a, vec<2, T, Q> const& b)
{
	return vec<2, T, Q>(Func(a.x, b.x), Func(a.y, b.y)); //<-HERE
}

detailed error output:

 detected during:
1>            instantiation of "vec<2, T, Q> glm::detail::functor2<vec, 2, T, Q>::call(Fct, const vec<2, T, Q> &, const vec<2, T, Q> &) [with vec=glm::vec, T=float, Q=glm::highp, Fct=glm::TMax<float>]"
1>D:\dev0\projects\DustRayTracer\DustRayTracer\thirdparty\glm\glm\detail/func_common.inl(274): here
1>            instantiation of "glm::vec<L, T, Q> glm::detail::compute_max_vector<L, T, Q, Aligned>::call(const glm::vec<L, T, Q> &, const glm::vec<L, T, Q> &) [with L=2, T=float, Q=glm::highp, Aligned=false]"
1>D:\dev0\projects\DustRayTracer\DustRayTracer\thirdparty\glm\glm\detail/func_common.inl(644): here
1>            instantiation of "glm::vec<L, T, Q> glm::max(const glm::vec<L, T, Q> &, const glm::vec<L, T, Q> &) [with L=2, T=float, Q=glm::highp]"
1>D:\dev0\projects\DustRayTracer\DustRayTracer\thirdparty\glm\glm\detail/func_common.inl(283): here
1>            instantiation of "glm::vec<L, T, Q> glm::detail::compute_clamp_vector<L, T, Q, Aligned>::call(const glm::vec<L, T, Q> &, const glm::vec<L, T, Q> &, const glm::vec<L, T, Q> &) [with L=2, T=float, Q=glm::highp, Aligned=false]"
1>D:\dev0\projects\DustRayTracer\DustRayTracer\thirdparty\glm\glm\detail/func_common.inl(659): here
1>            instantiation of "glm::vec<L, T, Q> glm::clamp(const glm::vec<L, T, Q> &, T, T) [with L=2, T=float, Q=glm::highp]"
1>D:\dev0\projects\DustRayTracer\DustRayTracer\thirdparty\glm\glm\detail/func_packing.inl(17): here

I also get host and device annotation ignore warnings:

warning #20012-D: __host__ annotation is ignored on a function("vec") that is explicitly defaulted on its first declaration
__device__ annotation is ignored on a function("vec") that is explicitly defaulted on its first declaration

and

__device__ annotation is ignored on a function("mat") that is explicitly defaulted on its first declaration
__host__ annotation is ignored on a function("mat") that is explicitly defaulted on its first declaration

for all types in detail/ folder

Please help.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

1 participant