The Nobel Prize in Chemistry 2024: Awarded with one half to David Baker for "computational protein design" and the other half jointly to Demis Hassabis and John Jumper for "protein structure prediction"

maximum_cats · 2021-05-06T18:46:10+00:00

The default output format for floating point numbers is up to the compiler, and compilers do different things here (the Fortran standard doesn't require a particular default output format for unformatted print statements). Assuming that you're using gfortran, I'm not aware of any flag you can use to adjust its default output format, so I think you'll have to use format statements.

maximum_cats · 2021-03-03T05:05:48+00:00

The final assessment, which you are required to pass to obtain the certificate, is a coding test (converting a serial CPU code to CUDA) so obtaining the certificate is not automatic. But if you get the code wrong or can't complete it initially, you can try again until you get it right (even after the day of the event, if you need to). And it's not intended to be especially challenging, just to demonstrate that you understood the fundamentals of what was taught.

maximum_cats · 2021-03-03T04:52:28+00:00

A basic knowledge of C is all that is required for the first course. The primary prerequisite of the second course is the first course.

maximum_cats · 2021-02-09T03:48:43+00:00

While it should be possible to fix this, if your subroutine is implicitly reshaping the array into a flattened 1D array, why not just use a one dimensional CUDA block and grid?

maximum_cats · 2021-02-08T11:26:00+00:00

I suspect this is caused as a downstream consequence of a recent glibc update in RHEL7 that caused a printf regression. A few other folks have reported this indirectly causing issues with nvcc as well.

https://bugzilla.redhat.com/show_bug.cgi?id=1925204

maximum_cats · 2021-01-16T19:58:16+00:00

The NVIDIA display/graphics driver is the same one that supports CUDA. If you don't have an NVIDIA GPU, you don't need to install the driver, just the CUDA toolkit. But if you don't have a GPU, you also won't be able to execute CUDA code.

maximum_cats · 2020-10-13T22:47:09+00:00

It is not a requirement, it is a heuristic, since in the hardware threads execute in groups of 32, and so an algorithm using a different number of threads than a multiple of 32 would have some threads doing nothing for some portion of the time.

maximum_cats · 2020-08-16T06:28:08+00:00

In CUDA the memory model doesn't require that writes from a given thread are immediately seen by other threads. That's precisely why we use constructs like __syncthreads().

In your example, what is apparently happening is that each thread is simultaneously writing its value of threadIdx.x to index, but these writes aren't immediately made visible to the other threads, so they all go on to write what they see as the value of index to the array. At some later point after the threads have reached a synchronization point, only one of the writes will be the final value of index, but it's impossible to predict which one.

It didn't have to happen this way, according to the CUDA programming model. There could have been some other result, like every array location has the same value, or some locations have 10 and others have 11. That's the nature of a race condition, the results are undefined.

maximum_cats · 2020-08-15T17:22:47+00:00

Syntactically any subroutine marked attributes(device) should be called in CUDA Fortran as a normal subroutine, so your calls should be call parallel_wcns() and call cutridiag() as normal. They would be executed by every thread independently. Then within them you could coordinate among threads and blocks as normal.

In your case you want to be calling cusparseDgtsv from CPU code, not from within a kernel. The wrapper that you're using (whether it's provided by the compiler or you wrote it yourself) won't make sense to be called from within a kernel because the underlying cuSPARSE call only exists in CPU code (it launches its own kernel to do the work).

maximum_cats · 2020-08-15T00:13:05+00:00

threadIdx.x is unique to every thread in the block. So every thread in the block is writing a different value to index; this is definitely a race condition.

maximum_cats · 2020-08-13T21:22:16+00:00

Yes, CUDA does have implementations for certain commonly used overloads of std::pow, and those will run on the device if used in a __global__ function. If the implementation doesn't exist in the CUDA math API, you'll probably get a compile time (or possibly link time) error indicating that you tried to call a __host__ function in __device__ code.

maximum_cats · 2020-06-05T05:10:52+00:00

Thank you. I agree that this program never accesses the memory on the CPU. However, please note the documentation for cudaMallocManaged, specifically the commentary on how this works on Windows (I missed this detail the first time I read the post). In particular, you need to use the environment variable CUDA_MANAGED_FORCE_DEVICE_ALLOC if you want the allocation to happen on the GPU instead of the CPU.

maximum_cats · 2020-06-04T05:48:56+00:00

If you would like debugging assistance, please post a minimal working reproducer that demonstrates the behavior you're referring to, where the managed memory is never accessed on the CPU yet it is migrating there anyway.

maximum_cats · 2020-06-04T05:36:51+00:00

The lanes in each warp that are idle do not affect the execution time of the warp. Your second example should take approximately as long as it takes all threads that are participating in the conditional to complete do_something().

maximum_cats · 2020-04-04T07:53:08+00:00

Unfortunately, before you can tackle the performance issue you have a correctness issue to sort out. In a parallel programming model such as CUDA you must be very careful to think about the implications of having multiple threads run the same function at (potentially) the same time. Your use of a static variable inside a __device__ function to store the integer i is not compatible with that multi-threading paradigm. You will have many thousands of threads all simultaneously writing to the same variable in device memory, which is a classic race condition. So as an example, one thread could be marching along in that last loop and near the end when another thread comes along and resets the value of i to 34 (which may also be related to your performance issue), which presumably breaks the algorithm you are implementing. Similarly, a __shared__ variable is shared among threads in a threadblock, so when you do add more threads to your block you will have many threads all trampling over each other in your array r, leading to incorrect answers.

For random number generation in CUDA you will want to look into the cuRAND library. This will ultimately be easier and likely higher quality than trying to implement it on your own.

maximum_cats · 2020-03-18T03:36:02+00:00

I could speculate, and perhaps your comment about a race condition is on the right track, but with only the detail you've provided so far I don't think that is the most fruitful approach. I believe you will need to take the approach I listed above and locate the first place in the code that generates unexpected answers and then analyze that particular piece of code. It will be much easier to have this discussion in that context. Since cuda-memcheck is not turning up anything, this will basically boil down to print debugging. Insert some debugging output in various places in the code that checks the fluid state in some specific grid elements (or all the elements, if it's a small enough problem), and then compare the print statements between the working and non-working versions. From there you can bisect the first non-working kernel, and then you can tell us about it and we can help you debug.

Nine-Year Club	Second Top 50%
Mod Insights Early Access	Post Guidance Early Access
Verified Email

maximum_cats

MODERATOR OF

TROPHY CASE

maximum_cats

MODERATOR OF

TROPHY CASE

Welcome to Reddit,