안녕하세요  저번 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() 을 이번 예제에서 빼도 상관 없지 않을까해서 질문드립니다. 좋은 강의 감사합니다.

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

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

안녕하세요?Jetson TX2 및 Xavier를 사용하여 CUDA예제를 실행하고 있습니다.Block 예제에서 tx2와 xavier의 결과가 달라서 질문드립니다.dimgrid<<<2,2,1>>>에 의해 Block의 갯수가 4개가 되고, warpid=0인 영역은 blockid가 (0,0), (1,0), (0,1), (1,1)인 곳에서 총 4개가 있을 것입니다.하지만 Tx2에서 예제를 돌려본 결과 blockid가 (0,0), (1,0)인 곳에만 나옵니다.혹시나 해서 xavier에서 실험을 해본 결과 총 4개가 나와 정상적으로 출력됩니다.CUDA언어는 GPU구조와 상관 없이 출력되기 위해 grid-block-thread 구조를 가지므로 속도는 느리지만 동일한 결과를 출력해야하는데, 왜 tx2와 xavier가 같은 예제를 돌렸는데도 다른 결과가 나오는걸까요? (혹시나 해서 <<<2,2,1>>> => <<<4,1,1>>>로 변경하여 돌려봤지만 여전히 tx2에서 blockid 2개만 출력되었습니다.)

안녕하세요!정말 재미있게 강의를 보고 있습니다.그런데 강의 32-2를 보다 궁금증이 생겨 질문 드립니다!warp의 early terminate가 정확히 어떤 동작인지 궁금한데요. 제가 생각해 볼 때 __syncthreads() 함수를 호출하게 되면, 블록 내 모든 스레드가 해당 함수에 도달하기까지 대기해야 하고 그러면 뒤에 더 이상의 동작이 없는 Warp들이라 하더라고 해당 함수에서 대기하고 있는 상태여야 할거 같습니다.따라서 그러면 해당 warp는 다른 warp들이 동작을 마칠때까지 기다리게 되어 리소스를 반환하거나 하는 동작이 없을거 같은데 왜이렇게 성능이 향상되는지 궁금합니다.

안녕하세요. 강의 정말 너무 잘 보고 있습니다!! 다름이 아니라 28-4 강의를 보다 궁금하게 있어서 질문을 드립니다.even odd방식이 Half and Half 알고리즘보다 느린 이유Global Memory에서 Half and Half가 even odd보다 더 느린 이유위 두 가지 이유에 대해서는 이해가 되었지만, Shared Memory를 사용하는 부분에서 이해가 잘되지 않는 부분이 있습니다. Global Memory에서는 Memory Coalescing이 중요하기 때문에 Half and Half가 더 느린게 맞지만, Shared Memory를 사용하는 부분에서는 두 방식(even odd, Half and Half)에서 차이가 없지 않을까 하고 생각했습니다. 두 방식 모두 Global Memory에 접근하고 Write하는 로직은 같기 때문에 Shared Memory에서 성능을 저해하는 요소로는 Bank Conflict만 있다고 생각했는데, 제가 잘못 이해 하고 있는 걸까요??

에러처리 매크로에서 exit을 하게되면 이전에디바이스나 호스트에서 동적할당했던 자원을 해제해주는건가요?? 아니면 메모리해제를 추가적으로 구현해야하나요??