- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요. print를해서 error를 처리했다고 표현하셨는데 해결을 했다는 의미가 아니고 그냥 print한 것이 처리했다는 의미가 맞나용? print를 했으니깐 reset해주는 것이고요 PeekAtLastError()는 print도 안하고 하면 볼 수 있는 것도 없는건가요?

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. v12.8인데 samples 파일이 없습니다.- 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요.

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요. 안녕하세요 좋은 강의 감사합니다.3중 for loop 개선책 설명해주신 부분 이해가 가지 않아 질문드립니다.말씀하신건 A도 B도 C도 캐쉬 친화적으로 이동 한다고 말씀하셨는데 A배열은 아래로 움직이지 않나요? 물론 matmul-host.cpp 보다는 캐쉬 미스 확률이 좋아지겠지만 A배열의 인덱스 이동은 캐쉬 친화적인거 같지 않은데 강의에서는 A도 옆으로 움직인다고 하셔서 헷갈려서 문의드립니다.

- 출장이 잦아서, 답변에 시간이 좀 걸릴 수도 있습니다. 양해 바랍니다. - 학습 관련 질문을 남겨주세요. 상세히 작성하면 더 좋아요! - 먼저 유사한 질문이 있었는지 검색해보세요. - 서로 예의를 지키며 존중하는 문화를 만들어가요. - 잠깐! 인프런 서비스 운영 관련 문의는 1:1 문의하기를 이용해주세요. 안녕하세요, 좋은 강의 감사합니다.transpose-shared.cu가 느린 이유에 대해 질문이 있습니다.강의 23-2에서는 X축으로 읽는 것이 아니라 Y축으로 메모리를 사용했기 때문에 느려졌다고 설명해 주셨습니다.보통 이런 경우, 이중 for문을 돌릴 때 메모리 공간 지역성(spatial locality)으로 인해 X축(연속된 메모리)으로 읽는 것보다 Y축(연속되지 않은 메모리)으로 읽는 것이 더 느리다고 알고 있습니다.하지만 이 transpose-shared.cu CUDA 예제에서는 각 스레드가 배열의 단 하나의 인덱스만 사용하므로,이중 for문에서 발생하는 문제라기보다는 이후 강의에서 설명해 주신 Bank Conflict가 발생해서 느려진 것이 아닐까 생각했습니다.제 생각이 맞는지 여쭤보고 싶습니다.  수정 및 추가 질문)챕터16강의 다시 보고 제 생각이 틀렸다는것을 인지했습니다.그럼 결국 데이터를 저장할때도 캐쉬 문제인것일까요?transpose-block.cu 는 캐쉬에 저장하고 한번에 flush가 가능하지만transpose-shared.cu 는 저장할때마다 캐쉬미스가 발생해서 매번 global memory에 접근해야하니 느려지는것일까요?

안녕하세요, 선생님의 좋은 강의 덕분에 잘 공부하고 있습니다.강의 내용 복기 중에 한가지 의문이 생겨 질문 남깁니다.강의 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단위로 스케줄링이 가능할 것 같은데, 제 생각이 맞는지 궁금합니다. 감사합니다. 

알찬 강의 계속 잘 듣고 있습니다!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() 을 이번 예제에서 빼도 상관 없지 않을까해서 질문드립니다. 좋은 강의 감사합니다.