안녕하세요, 선생님의 좋은 강의 덕분에 잘 공부하고 있습니다.강의 내용 복기 중에 한가지 의문이 생겨 질문 남깁니다.강의 17-8 matrix addition, pitched의 4분 07초 부분에서 device pitch (dev_pitch) 값이 40448 byte로 나온 것을 확인할 수 있고, 이어서 이 경우에 CUDA GPU 메모리 alignment boundary가 256 byte라고 설명 주셨습니다.여기서, 만약 alignment boundary가 256 byte 라면 device pitch 값은 40448 byte가 아닌 40192 byte 가 나와야 하지 않는지 의문이 들었습니다. 40192 byte가 40000 byte(10000개 elements) 에서 가장 가까운 256의 배수임과 동시에 40000 byte를 모두 커버할 수 있다는 근거로 이렇게 생각을 했습니다.그래서 alignment boundary 값을 512 byte로 계산을 해보니, 40448 byte가 40000 byte로부터 가장 가까운 512 배수로 계산이 되는데요. 이 경우에, 어떤 값이 맞는지 확신이 서지 않아 질문 글을 작성하게 되었습니다.제가 사용한 계산식은 pitch = ceil(row_size_in_byte / alignment_boundary_in_byte) * alignment_boundary_in_byte 입니다. 제가 아직 공부하는 단계이고, 이해가 부족했을 가능성이 큽니다. 답변에 소중한 시간 내어주시면 정말 감사드립니다. 다시 한번 멋진 강의 만들어주셔서 감사합니다.

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요.안녕하세요! C++과 CUDA 공부를 쉽고 재미있게 잘 하고 있습니다 ^^ 부족해서 하나부터 다 공부하고 있는데요.중요한 건 아닌 것 같습니다만 common.cpp의 procArg 함수의 이하 if문의 conditoin에 대한 질문이 생겨서 글을 남기게 되었습니다.if (typeid(TYPE) == typeid(float) && typeid(TYPE) == typeid(double)) 다름이 아니라 TYPE의 변수 타입을 알아내 strtof인지 strtol을 사용하는 결정의 if문의 조건에 걸린 연산자가 || (or) 연산자여야 할 것 같은데요..어차피 vecSize의 개수는 정수로 떨어지니 저렇게 만들어두셨나 싶다가도 정 궁금해서 여쭤봅니다! 그럼 답변 기다리겠습니다.감사합니다 😀

"hist-shared.cu"의 다음 부분에서 질문있습니다.if(threadIdx.x < HIST_SIZE) { atomicAdd(&(hist[threadIdx.x]), s_hist[threadIdx.x]); }threadIdx.x는 0~31로 모두 다르기 때문에 병렬 연산이긴 해도, global memory에 있는 hist 배열의 서로 다른부분으로 write가 이루어진다고 생각했습니다. 따라서 서로 겹치는 위치가 없어서 atomic한 연산이 필요없을거라 생각했는데, atomic 연산이 없으면 결과가 이상하게 나오더라구요.. 혹시 왜 그런지 알 수 있을까요? 양질의 강의 제공해주셔서 감사합니다!

강의에서 10,000 x 10,000 행렬의 덧셈을 해결하고자 했고, 이를 2D thread block(32 x 32)으로 풀어내셨습니다. 1D thread block을 사용하지 않는 것에 대한 의문이 들었습니다. [질문]1D thread block은 실무에서 잘 쓰이지 않나요? 의문을 가진 이유는 다음과 같습니다. 우선 2D로 설계할 경우, thread block의 dimension을 어떻게 설정하는지에 따라 낭비되는 thread가 달라집니다.[1 x 1,024] : GridDim(10,000 x 10 x 1), 2,400,000개 낭비[2 x 512] : GridDim(5,000 x 20 x 1), 2,400,000개 낭비[4 x 256] : GridDim(2,500 x 40 x 1), 2,400,000개 낭비[8 x 128] : GridDim(1,250 x 79 x 1), 1,120,000개 낭비[16 x 64] : GridDim(625 x 157 x 1), 480,000개 낭비[32 x 32] : GridDim(313 x 313 x 1), 320,256개 낭비 Thread 낭비가 가장 적은 block dimension은 강의에서 소개해주신 [32 x 32] 입니다. 저는 여기서 1D로 설계했을 때, 낭비되는 therad의 개수가 궁금해졌습니다.[1,024] : GridDim(97,657 x 1 x 1), 768개 낭비 1D로 thread block을 설계했을때, therad 낭비는 768개로 2D로 설계하는 경우와 비교했을때 적지않은 차이라고 생각합니다. 물론 실제로 측정해본 결과(RTX 3090) 약 40 usec 차이로 1D thread block인 경우가 더 빠르긴 했지만 생각보다 큰 차이는 아니었습니다.양질의 강의를 제공해주셔서 감사합니다!!

안녕하세요  저번 Vector Addition 수업에서는 256M(256 x 1024 x 1024)개의 연산을 수행하기 위해서<<<256*1024, 1024>>>로 커널을 런치하여 Block의 개수에는 제한이 없는 것 처럼 보였는데,15-5강의에서는 SM의 Thread Block은 32개 까지 가능이라고 하셔서 헷갈리는데두 개의 Thread Block 개념이 다른 것인지 궁금합니다.

안녕하세요. 먼저 좋은 CUDA 강의를 제공해주셔서 감사합니다. 다름이 아니라 15-4 강의에서 Warp를 동시에 돌린다는 말에 질문이 있습니다. Warp 스케줄링을 저는 아래와 같이 비교하여 이해하였습니다.CPU: Core -> 여러개의 쓰레드GPU: SM -> 여러개의 Warp 강의 내용처럼, CUDA에서 Warp별로 스케줄링을 관리를 한다면 모든 SP의 Clock들이 같은 Instruction을 수행(SM내 SP들의 Clock들은 Sync가 맞음)하여야 Warp단위로 스케줄링이 가능할 것 같은데, 제 생각이 맞는지 궁금합니다. 감사합니다. 

ppt 자료는 없나요..수업중에 ppt에 포함되어 있다고 말씀하셨는데.. ppt 자료는 어디서 구할 수 있나요

알찬 강의 계속 잘 듣고 있습니다!CPU와 GPU version의 계산시간 비교에 대해 질문 드립니다. 커널에서 계산시간만 비교하면 CUDA이 훨씬 빠릅니다. CPU version (filter-host.cu): 44,784 usCUDA version (filter-dev.cu): 830 us하지만 CUDA version에서 메모리 복사를 포함하면 337,737 us 으로 CPU version보다 훨씬 느린 것 같습니다. 혹시 메모리 복사 시간을 줄이는 방법이 있는지요?

안녕하세요!좋은 강의 잘 듣고 있습니다.common.cpp 코드 중에서 ELAPSED_TIME_BEGIN(N)과 ELAPSED_TIME_BEGIN(N)을 매크로 함수로 정의한 이유가 무었인지요?CUDA_CHECK_ERROR() 함수는 __FILE__과 __LINE__ 때문에 매크로 함수로 정의하신 걸로 이해합니다만, 시간 측정 함수들을 매크로함수로 정의하신 이유는 잘 이해되지 않습니다.답변 부탁드립니다.감사합니다.

좋은 강의 고맙습니다.병렬처리라면 일반적으로 다른 쓰레드에 영향을 주지 않고 수행이 될거라 생각이 됩니다.예제 코드에서 hello<<<1,8>>>();을 이용하여 쓰레드 8개를 병렬로 돌리는데 출력된 값은hello CUDA 0!hello CUDA 1!hello CUDA 2!,,,hello CUDA 7!이 되는 것을 볼 수 있는데 병렬로 처리된다면 출력 되는 숫자가 순차적으로 나올 수 없는 것 아닌가요?

제목 그대로 kernelMatCpy에서는 각 스레드 작업이 다른 스레드에 영향을 주지 않기 때문에 필요없을 것 같은데요. 실제로 __stncthreads()를 제거하고 돌려봐도 정상적으로 결과가 나오구요.혹시 이런 경우라도 내부적으로 꼬일 수 있어서 사용하신 것인지 아니면 대부분의 shared memory로 복사해서 쓰는 경우에 필요하기 때문에 습관(?)차원에서 사용하신 것인지 궁금합니다.

소스코드 그대로 돌리면, 아래와 같은 에러가 발생합니다. (25장 전체 소스가 모두 동일합니다) 이유와 수정 방법을 알고 싶읍니다. // 25d-gemm-alignedTile.cu/tmp/tmpg97edtlu/25d-gemm-alignedTile.cu(98): error: "beta" is ambiguous  1 error detected in the compilation of "/tmp/25d-gemm-alignedTile.cu". 

안녕하세요.먼저 좋은 강의 감사드립니다. 아래 소스는 matcpy-dev.cu에서 가져온 kernel 함수입니다.global void kernelMatCpy( float* C, const float* A, int matsize, size_t pitch_in_elem ) { register unsigned gy = blockIdx.y * blockDim.y + threadIdx.y; // CUDA-provided index if (gy < matsize) { register unsigned gx = blockIdx.x * blockDim.x + threadIdx.x; // CUDA-provided index if (gx < matsize) { register unsigned idx = gy * pitch_in_elem + gx; // in element C[idx] = A[idx]; } }} 그리고 다음 소스는 transpose-dev.cu에서 가져온 kernel 함수 입니다.global void kernelMatTranspose( float* C, const float* A, unsigned matsize, size_t pitch_in_elem ) { register unsigned gy = blockIdx.y * blockDim.y + threadIdx.y; // CUDA-provided index if (gy < matsize) { register unsigned gx = blockIdx.x * blockDim.x + threadIdx.x; // CUDA-provided index if (gx < matsize) { register unsigned idxA = gy * pitch_in_elem + gx; register unsigned idxC = gx * pitch_in_elem + gy; C[idxC] = A[idxA]; } }} 메모리 접근 관점에서 보면 각 thread별로 A에서 read해서 C에 write하는 동일한 과정으로 보이는데, 결과는 속도 차이가 많이 나는 정확한 원인이 궁금합니다. memory coalescing 때문으로 추정하고 있는데, 맞는 생각일까요?감사합니다.

예제소스코드12a-vecadd-host.cpp -> vecadd-host.cpp 로 파일명 고치고 common.cpp 랑 2개를 쿠다 프로젝트에 추가해서 컴파일 했는데 링크에러가 납니다.vecadd-host.obj : error LNK2005: "float __cdecl getRMS(float const ,float const ,int,bool)" (?getRMS@@YAMPEBM0H_N@Z) already defined in common.obj파일은 저 2개 뿐인데already defined in common.obj 이런 에러가 여러 개가 뜨는데 뭐 때문일까요?    

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요. 선생님 너무 좋은 강의 다시 한번 감사드립니다. 다름이 아니고, video에서 각각의 frame별로 filter를 처리하는 부분을 cuda도 구현하는 코드를 짠다고 가정할 때, 각 프레임은 W x H x C ( rbg) 의 차원을 가지고 있고 영상이기 때문에 영상은 W x H x C x T 의 차원을 가지고 있을 겁니다. 그렇다면, cudaMalloc3D나 cudaMemecpy3D 부분을 사용해서 각 frame을 처리하고 ( device ) 나머지는 for 문 ( host ) 으로 시간에 따라 frame을 device로 보내도록 구현하였습니다. 그런데 혹시 시간 부분 까지 device에서 처리할 수 있는 방법이 있는 지가 궁금합니다.

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요. 안녕하세요, 선생님! 궁금한 게 있어서요. device memory와 host memory는 구분되어 있잖아요?? 그래서 각각 사용되는 데이터는 각각의 메모리에서 사용하는 것으로 알고 있었어요. 그런데 vecSize 나 saxpy_a는 global로 선언 되어 있는데, 이 변수들은 device메모리에서 어떻게 사용하는 거죠?? global 변수의 위치는 host memory에 위치 하지 않나요??

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요.안녕하세요. 선생님 좋은 강의 잘 듣고 있습니다. 강의 부분에 질문이 하나 있어 이렇게 남깁니다.kernel 을 launch한 이후에 cudaDeviceSynchronize() 를 call 해주는 이유가 kernel의 작업 완료를 보장하기 위함이라는 사실은 인지하고 있습니다.그런데, 그 아래 부분에서 cudaMemcpy를 해주기 때문에 저 부분이 생략해도 되지 않나 싶어서 질문을 남깁니다. memcpy가 저번 강의 때, kernel 작업 완료 후에 이뤄진다고 말씀하셨던 부분이 있어서, cudaDeviceSynchronize() 을 이번 예제에서 빼도 상관 없지 않을까해서 질문드립니다. 좋은 강의 감사합니다.

선생님 안녕하세요. 저번에 opencl과 cuda를 같이 공부 중이던 학생이에요.제가 프로그램을 작성하는 도중에 VRAM <-> RAM의 전송 속도가 느리다는 것을 깨달았는데요.제가 보통 결과 값으로 40mb 정도를 가지는데(위의 자료에 따르면 6ms 정도) 이를 최대한 cpu에서 빨리 받기 위한 방법은 아직까지 존재하지 않는 것인지 알고 싶습니다!

알려주시면 감사하겠습니다!

SM내 warp가 24개가 이상일 경우에 overlab이 생기지 않아 효율적인 스케줄링이 가능하다는 것은 정확히 이해 했지만, Warp 갯수가 적어지면 근본적으로 작업시간 자체는 줄지 않을까? 그렇다면 빨리 끝나는게 스케줄링이 좋은거 아닐까? 라는 생각이 들었습니다.다음은 제 생각을 담은 그림입니다.다음과 같이 23개의 warp가 overhead가 걸릴 수는 있겠지만 일이 끝나는 시간은 짧지 않을까?또 반대로 warp의 수가 작다는 것은 일 양이 많지 않다는 것이니 당연히 빨리 끝나는게 맞지 않을까?라는 생각을 했습니다.