[CUDA Tip] Visual Studio에서 CUDA 사용시 빨간 줄 없애기

[2022/06/15] 추가 사항

본 글에서 설명하는 수정사항을 자동으로 설치할 수 있게 script로 만들었다. 

혹시 필요한 사람은 아래 repository에서 script를 받아서 사용하면 된다. 

 

https://github.com/pjueon/cuda_intellisense

GitHub - pjueon/cuda_intellisense: A simple python script to fix cuda C++ intellisense for visual studio.

 

 

Visual Studio에서 CUDA C++ 코드를 작성해보면 불편한 부분들이 있다. 정상적으로 컴파일은 되지만 문법검사오류로 인해 빨간줄이 생기고 오류메시지가 뜨는 것이다. 이런 현상의 가장 큰 문제는 심리적으로 매우 불편하다는 것이다. 또한 단순히 심리적인 문제 뿐 아니라, 이런 오류메시지가 다량으로 발생하는 환경에서는 무언가 컴파일 에러가 생겼을때 이런 메시지들에 묻혀서 실제 문제되는 오류나 warning을 놓치기 쉽다는 문제도 있다(물론 출력창을 "빌드 + Intellisense"에서 "빌드"로만 바꾸면 이 문제는 해결할 수 있다)

이 글에서는 필자가 이 문제를 해결하기 위해 사용하고 있는 방법에 대해서 공유한다.

CUDA 코드에 생기는 못생긴 빨간 줄들

 

참을 수 없는 불편함...

 

1. 문제 원인

이러한 현상은 실제 코드를 컴파일하는 컴파일러와, 문법검사/자동완성등을 담당하는 Intellisense용 컴파일러가 다르기 때문에 생기는 문제이다.(정확한 명칭은 잘 몰라서 Intellisense용 컴파일러라고 칭했지만, 기능적으로 보면 컴파일러라기보다는 파서에 가까울 것 같다. 사실 엄밀한 정의는 잘 모르겠다. 이 글에서는 이해하기 쉽게 컴파일러라고 칭하도록 하겠다.)

 

기본적으로 Visual Studio에서 CUDA C++ 코드를 포함한 프로젝트를 사용하게 되면 3가지의 컴파일러가 작동한다.

- Visual C++ : CPU에서 동작하는 일반 C++ 코드를 컴파일한다. host 컴파일러. MS에서 만듦.

- NVCC :  GPU에서 동작하는 코드를 컴파일한다. NVIDIA에서 만듦.

- Intellisense용 컴파일러: Visual Studio 상에서 문법검사/자동완성 등의 기능(Intellisense)에만 사용되는 컴파일러. 

 

여기서 Intellisense용 컴파일러는 CUDA C++의 문법을 모르기 때문에 CUDA 코드를 문법오류로 판단하고, 빨간줄이 그어지는 것이다.

 

구체적으로 살펴보면 CUDA C++ 코드에서 빨간 줄이 생기는 경우에는 크게 두 가지 정도의 패턴이 있다.

 

a) Kernel함수를 call 하는 문법. 

CUDA C++에서는 kernel함수를 call할때 다음과 같은 syntax를 사용한다. 

MyKernel<<<1, size>>>(a, b, c)

여기서 <<< >>> 이 부분. 이러한 문법은 C++에 존재하지 않는다. 그래서 이 부분에서 빨간 줄이 생긴다. 심지어, 이 부분을 자동정렬 할 경우 >>> 가 >> > 이렇게 두개로 쪼개지는 매우 불편한 현상을 볼 수 있다.(시프트 연산자와 부등호 연산자로 인식)

 

b) __CUDACC__ 및 기타 CUDA의 매크로들.

__CUDACC__ 매크로는 NVCC 컴파일러에서 CUDA C++ 코드를 빌드시 자동으로 정의해주는 매크로이다. 

CUDA Toolkit 공식 문서 중 발췌

해당 매크로는 NVCC가 빌드시 알아서 추가하는 매크로이기 때문에 Intellisense 컴파일러 입장에선 정의되어 있지 않은 매크로이다. 여기서 문제가 발생한다. 기본적으로 CUDA에서 제공하는 라이버러리들의 헤더파일을 보면, __CUDACC__ 매크로의 정의 유무를 검사해서 컴파일 분기가 변경되도록 되어있는 경우가 매우 많다. 이로 인해 몇몇 CUDA에서 제공하는 라이브러리들을 사용할때, Intellisense 컴파일러는 특정 클래스나 함수 등이 정의되지 않았다고 에러를 띄운다. 

 

texture 클래스는 __CUDACC__ 매크로가 정의되어 있을때만 컴파일 되도록 정의되어 있다.

그러면 사용자가 __CUDACC__ 매크로를 사용자 코드에서 명시적으로 정의해주면 해결되지 않을까? 불행하게도 이번에는 또 다른 문제가 생긴다.

__CUDACC__ 매크로를 명시적으로 정의하면, 각종 CUDA 매크로들에 빨간 줄이 생긴다.

 

CUDA에서 사용하는 __device__, __host__, __gloabal__, __constant__ 등의 매크로들 역시 __CUDACC__의 정의 여부에 따라 컴파일 분기가 변경된다. __CUDACC__가 정의되어 있지 않을때는 이들 매크로는 아마도 공백으로 번역되어 일반적인 C++ 코드에서 해당 매크로를 사용하는 헤더를 include 했을때도 문제가 발생하지 않게 되어있는 것 같다.(확실하진 않음) 하지만 __CUDACC__ 매크로가 정의되어 있을때는, 이 매크로들은 CUDA C++에만 존재하는(Intellisense가 알지 못하는) attribute로 번역 된다.  그렇기 때문에, 사용자가 명시적으로 __CUDACC__ 매크로를 정의하면, Intellisense 가 이러한 매크로들을 문법오류로 판단하여 못 생긴 빨간 줄이 생긴다.

 

즉, __CUDACC__ 매크로를 명시적으로 정의하든 정의하지 않든 빨간 줄이 생긴다.

 

2. 해결 방법

그렇다면 이 문제는 해결 불가능한 것인가? 다행이도 아니다. 해결책은 있다. 관련 문제로 열심히 구글링을 하다가 좋은 우회책을 찾았다. 이 방법의 핵심은,  __INTELLISENSE__ 매크로를 이용하는 것이다.

__INTELLISENSE__ 매크로는 Intellisense 컴파일러가 작동할때만 정의되는 매크로로, 실제 코드가 빌드될때는 해당 매크로는 정의되지 않는다. 

__INTELLISENSE__ 로 검색해보면 MSDN에 설명이 나온다

 

예를 들어 Visual Studio에서 아래와 같은 코드를 작성했다고 해보자. 

#include<iostream>
int main()
{
#ifdef __INTELLISENSE__
	std::cout << "Intellisense!";
#else
	std::cout << "build!";
#endif
	return 0;
}

그러면 Intellisense 컴파일러 입장에서는 __INTELLISENSE__ 매크로가 정의되어있기 때문에 위쪽 분기가 활성화하여 보여준다. 하지만 실제로 빌드해보면, 빌드시에는 __INTELLISENSE__ 매크로가 정의되어 있지 않기 때문에 아래 분기로 컴파일되면서 실행시 build! 라는 메시지가 뜨게 된다.

 

__INTELLISENSE__ 매크로 사용 예시

 

한마디로 이 매크로를 사용하면 Visual Studio의 Intellisense를 속일 수 있는 것이다!

 

a) Kernel함수 Call 관련

먼저 Kernel함수 Call 부분에서 생기는 빨간 줄을 우회해보자. 이 방법은 stackoverflow에 있는 아래 글을 참고했다.

stackoverflow.com/questions/6061565/setting-up-visual-studio-intellisense-for-cuda-kernel-calls/27992604#27992604

Setting up Visual Studio Intellisense for CUDA kernel calls

아래와 같이 cuda_intellisense.hpp 파일을 만든다(파일이름은 원하는대로 바꿔도 됨)

#pragma once
#ifdef __INTELLISENSE__
// Intellisense 컴파일러에 보여줄 부분(공백으로 정의)

//KERNEL_ARG2(grid, block) : <<< grid, block >>>
#define KERNEL_ARG2(grid, block)
//KERNEL_ARG3(grid, block, sh_mem) : <<< grid, block, sh_mem >>>
#define KERNEL_ARG3(grid, block, sh_mem)
//KERNEL_ARG4(grid, block, sh_mem, stream) : <<< grid, block, sh_mem, stream >>>
#define KERNEL_ARG4(grid, block, sh_mem, stream)

#else
//실제 코드 컴파일시에 적용되는 부분

#define KERNEL_ARG2(grid, block) <<< grid, block >>>
#define KERNEL_ARG3(grid, block, sh_mem) <<< grid, block, sh_mem >>>
#define KERNEL_ARG4(grid, block, sh_mem, stream) <<< grid, block, sh_mem, stream >>>

#endif

이제 사용자의 CUDA 코드에서, 위 헤더를 include한 후에 kernel call을 할때 다음과 같이 사용하면 된다.

// MyKernel <<<1, size>>> (a, b, c) 로 컴파일 됨
MyKernel KERNEL_ARGS2(1, size) (a, b, c)

Intellisense 컴파일러가 볼때는 위에서 우리가 정의한 매크로들은 아무런 내용이 없는 공백으로 보이게 되기 때문에 빨간 줄이 안생긴다. 물론 코드를 빌드할 때는 정상적으로 <<<  >>> 형태로 컴파일된다.

 

b) __CUDACC__ 매크로 관련

다음으로 __CUDACC__ 매크로 관련 문제도 해결해보자. 

일단 다음과 같이 cuda_intellisense_attribute.hpp 파일을 만든다(파일이름은 원하는대로 바꿔도 됨)

#pragma once
// **********실제 코드 컴파일시에는 모두 무시됨**********
#ifdef __INTELLISENSE__

// __CUDACC__ 가 define 된 것"처럼" 보여주기
#define __CUDACC__

// CUDA 관련 매크로들을 공백으로 재정의한 것"처럼" 보여주기
#define __global__ 
#define __host__ 
#define __device__ 
#define __device_builtin__
#define __device_builtin_texture_type__
#define __device_builtin_surface_type__
#define __cudart_builtin__
#define __constant__ 
#define __shared__ 
#define __restrict__
#define __noinline__
#define __forceinline__
#define __managed__

#endif 

파일 내용을 보면 먼저 모든 매크로가 #ifdef __INTELLISENSE__  분기 안에 있다. 즉, 위 코드는 intellisense에게만 보이는 코드이다. 그 다음부터는 주석의 내용처럼, __INTELLISENSE__에게도 __CUDACC__ 매크로가 정의되어 있는 것처럼 보이게 하되(각종 라이브러리의 컴파일 분기를 제대로 잡기 위함), CUDA C++에만 있는 attribute로 번역되어 버리는 각종 매크로들을 아무 내용없는 공백으로 재정의한 것처럼 보여줘서 Intellisense에서 문제가 생기지 않도록 한 것이다. 

 

이제 이 파일을 CUDA 관련 매크로가 정의되어있는 host_defines.h파일의 맨 마지막에 include 시켜야 한다. 해당 파일은 $(CUDA_PATH)\include\crt 폴더에 있다.  여기서 $(CUDA_PATH)는 자신이 CUDA를 설치한 경로를 뜻한다. 필자의 경우 수정할 파일의 전체 경로는 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.2\include\crt\host_defines.h 이다. (해당 파일을 수정하려면 관리자 권한이 필요하다.)

작성한 헤더파일을 host_defines.h 파일의 마지막에 include한다.

 

3. 적용 결과

before

after

빨간 줄은 사라지고 편안함은 오래 갈거야

4. 단점

Kernel 함수 call 과 관련한 부분에 있어서는 개인적으로 별다른 단점이 없다고 생각한다. 하지만 __CUDACC__ 매크로와 관련된 부분에서는 약간의 단점이 존재한다. 바로 NVIDIA에서 배포하는 헤더파일을 사용자가 직접 수정해야한다는 점이다. 물론 수정이라고 해봐야 실제 빌드시에는 아무런 코드도 추가되지 않기 때문에 부작용은 없지만, 그래도 공식적으로 배포된 파일을 수정해야 한다는 점 자체가 찝찝할 수 있다. 그러나 개인적으로는 그 찝찝함보다도 수많은 빨간 줄과 오류 메시지들을 없앰으로 얻는 편안함이 더 크다고 생각한다. 

 

혹시 이 방법보다 더 좋은 해결 방법을 알고 있다면 댓글로 공유 바란다.