Shuffle instruction in CUDA not working

I have got problem with shuffle instruction in CUDA 5.0.

This is snippet of my kernel. It is inside the loop. Print is there only for debug purpose because I can't use ordinary debugger:

...
tex_val = tex2D(srcTexRef, threadIdx.x + w, y_pos);
if (threadIdx.x == 0)
{
    left = left_value[y_pos];
}
else
{
    printf("thread %d; shfl value: %f \n", threadIdx.x, __shfl_up(value, 1));
    left = __shfl_up(value, 1);
}

printf("thread %d; value: %f; tex_val: %f; left: %f \n", threadIdx.x, value, tex_val, left);
...

From that I get this output:

l0:  ITERATION 1
l1:  thread 0; value: 0; tex_val: 1; left: 4
l2: 
l3:  ITERATION 2
l4:  thread 1; shfl value: 0
l5:  thread 0; value: 5; tex_val: 1; left: 5
l6:  thread 1; value: 0; tex_val: 1; left: 0
l7: 
l8:  ITERATION 3
l9:  thread 1; shfl value: 0
l10: thread 2; shfl value: 1
l11: thread 0; value: 6; tex_val: 1; left: 6
l12: thread 1; value: 1; tex_val: 1; left: 0
l13: thread 2; value: 2; tex_val: 1; left: 1
...

From the output I can see that thread 1 doesn't get value from thread 0 in any iteration even though I can clearly see that it has value (line 4 - shfl value is 0; line 5 - value is 5). Thread 2 and higher can get value from lower thread. Where am I making mistake? Is it happening because of the branching?

Best How To :

Yes, it's because of the branching. Quoting from the CUDA programming guide B.14.2:

The __shfl() intrinsics permit exchanging of a variable between threads within a warp without use of shared memory. The exchange occurs simultaneously for all active threads within the warp, ...

and

Threads may only read data from another thread which is actively participating in the __shfl() command. If the target thread is inactive, the retrieved value is undefined.

In a branch, active threads are those taking the same path of execution, while those taking different ones are inactive. In your case, thread 0 is inactive, so you cannot shuffle from it.

performance issues executing list of stored procedures

c#,multithreading,performance,loops

What I would do is something like this: int openThread = 0; ConcurrentQueue<Type> queue = new ConcurrentQueue<Type>(); foreach (var sp in lstSps) { Thread worker = new Thread(() => { Interlocked.Increment(ref openThread); if(sp.TimeToRun() && sp.HasResult) { queue.add(sp); } Interlocked.Decrement(ref openThread); }) {Priority = ThreadPriority.AboveNormal, IsBackground = false}; worker.Start(); } //...

Parameters to use in a referenced function c++

c++,pointers,reference

Your code makes no sense, why are you passing someStruct twice? For the reference part, you should have something like: void names(someStruct &s) { // <<<< Pass struct once as a reference cout << "First Name: " << "\n"; cin >> s.firstname; cout << "Last Name: " << "\n"; cin...

Multiple Threads searching on same folder at same time

c#,multithreading,file-search

Instead of using ordinary foreach statement in doing your search, you should use parallel linq. Parallel linq combines the simplicity and readability of LINQ syntax with the power of parallel programming. Just like code that targets the Task Parallel Library. This will shield you from low level thread manipulation and...

Type function that returns a tuple of chosen types

c++,templates,c++11,metaprogramming

You can do this without recursion by simply expanding the parameter pack directly into a std::tuple: template<My_enum... Enums> struct Tuple { using type = std::tuple<typename Bind_type<Enums>::type...>; }; To answer your question more directly, you can declare a variadic primary template, then write two specializations: for when there are at least...

Make a triangle shape in C++

c++

This code works fine for a right angled triangle - * ** *** But I guess you want a triangle like this - * *** ***** Try this - #include <iostream> using namespace std; int main() { int i, j, k, n; cout << "Please enter number of rows you...

MFC visual c++ LNK2019 link error

c++,mfc

The header file provides enough information to let you declare variables. And for that matter to just compile (but not link) code. When you link, the linker has to resolve e.g. function references such as a reference to ServerConnection::getLicenceRefused, by bringing in the relevant machine code. You have to tell...

Test if string represents “yyyy-mm-dd”

c++,command-line-arguments

If you can use boost library you could simple do it like this: string date("2015-11-12"); string format("%Y-%m-%d"); date parsedDate = parser.parse_date(date, format, svp); You can read more about this here. If you want a pure C++ solution you can try using struct tm tm; std::string s("2015-11-123"); if (strptime(s.c_str(), "%Y-%m-%d", &tm))...

Implicit use of initializer_list

c++,c++11,initializer-list

Your program is not ill-formed because <vector> is guaranteed to include <initializer_list> (the same is true for all standard library containers) §23.3.1 [sequences.general] Header <vector> synopsis #include <initializer_list> ... Searching the standard for #include <initializer_list> reveals the header is included along with the following headers <utility> <string> <array> <deque> <forward_list>...

OpenCV - Detection of moving object C++

c++,opencv

Plenty of solutions are possible. A geometric approach would detect that the one moving blob is too big to be a single passenger car. Still, this may indicate a car with a caravan. That leads us to another question: if you have two blobs moving close together, how do you...

How can I convert an int to a string in C++11 without using to_string or stoi?

c++,string,c++11,gcc

Its not the fastest method but you can do this: #include <string> #include <sstream> #include <iostream> template<typename ValueType> std::string stringulate(ValueType v) { std::ostringstream oss; oss << v; return oss.str(); } int main() { std::cout << ("string value: " + stringulate(5.98)) << '\n'; } ...

Same function with and without template

c++,c++11

The main reason to do something like this is to specialize void integerA(int x) to do something else. That is, if the programmer provides as input argument an int to member function abc::integerA then because of the C++ rules instead of instantiating the template member function the compiler would pick...

Copy text and placeholders, variables to the clipboard

c++,qt,clipboard

You're not using the function setText correctly. The canonical prototype is text(QString & subtype, Mode mode = Clipboard) const from the documentation. What you want to do is assemble your QString ahead of time and then use that to populate the clipboard. QString message = QString("Just a test text. And...

Add more features to stack container

c++,visual-c++,stl

If this is interview question or something , and you have to do it anyways , you can do this like ,below code . derive from std::stack , and overload [] operator #include <iostream> #include <algorithm> #include <stack> #include <exception> #include <stdexcept> template <typename T> class myStack:public std::stack<T> { public:...

Get an ordered list of files in a folder

c++,boost,boost-filesystem

The fanciest way I've seen to perform what you want is straight from the boost filesystem tutorial. In this particular example, the author appends the filename/directory to the vector and then utilizes a std::sort to ensure the data is in alphabetical order. Your code can easily be updated to use...

dispatch response packet according to packet sequence id

c++,boost,boost-asio

You could use std::promise and std::future (or their boost counterparts if your are not yet on C++11). The idea is to store a std::shared_ptr<std::promise<bool>> with the current sequence id as a key in the map whenever a request is sent. In the blocking send function you wait for the corresponding...

.cpp:23: error: cannot convert ‘std::string’ to ‘const char’ for argument ‘1’ to ‘int atoi(const char)’

c++,string

Use stoi, it's the modern C++ version of C's atoi. Update: Since the original answer text above the question was amended with the following error message: ‘stoi’ was not declared in this scope Assuming this error was produced by g++ (which uses that wording), this can have two different causes:...

C++ & Qt: Random string from an array area

c++,arrays,string,qt,random

You should use the random header. #include <random> std::default_random_engine generator; std::uniform_int_distribution dist(0, 5); int StringIndex = dist(generator); std::string ChosenString = characters[StringIndex]; The above will generate a random index into your array. If you want to limit the range, change the constructor of dist, for example (dist(0,2) would only allow for...

Explicit instantiation of class template not instantiating constructor

c++,templates,constructor,explicit-instantiation

When the constructor is a template member function, they are not instantiated unless explicitly used. You would see the code for the constructor if you make it a non-template member function. template<typename T> class test { public: /*** template<typename T> test(T param) { parameter = param; }; ***/ test(T param)...

undefined reference to `vtable for implementation' error

c++,build,makefile

I think you just misspelled CFLAGS in CFLAGES=-c -Wall I'm guessing this is the case since g++ ../src/main.cpp -I ../include/ does not have the -c option...

C++11 Allocation Requirement on Strings

c++,string,c++11,memory,standards

Section 21.4.1.5 of the 2011 standard states: The char-like objects in a basic_string object shall be stored contiguously. That is, for any basic_string object s, the identity &*(s.begin() + n) == &*s.begin() + n shall hold for all values of n such that 0 <= n < s.size(). The two...

Strings vs binary for storing variables inside the file format

c++,file,hdf5,dataformat

Speaking as someone who's had to do exactly what you're talking about a number of time, rr got it basically right, but I would change the emphasis a little. For file versioning, text is basically the winner. Since you're using an hdf5 library, I assume both serializing and parsing are...

create vector of objects on the stack ? (c++)

c++,vector,heap-memory

Yes, those objects still exist and you must delete them. Alternatively you could use std::vector<std::unique_ptr<myObject>> instead, so that your objects are deleted automatically. Or you could just not use dynamic allocation as it is more expensive and error-prone. Also note that you are misusing reserve. You either want to use...

Method returning std::vector>

c++

Your error is actually coming from: array.push_back(day); This tries to put a copy of day in the vector, which is not permitted since it is unique. Instead you could write array.push_back( std::move(day) ); however the following would be better, replacing auto day...: array.emplace_back(); ...

Undefined behaviour or may be something with memset

c++,undefined-behavior

The A[32] in the method is actually just a pointer to A. Therefore, sizeof is the size of *int. Take the following test code: void szof(int A[32]) { std::cout << "From method: " << sizeof(A) << "\n"; } int main(int argc, char *argv[]) { int B[32]; std::cout << "From main:...

segfault accessing qlist element through an iterator

c++,iterator,qlist

(Edited away first "answer", this is an actual attempt at an answer) My guess: QList<Msg> messages() const { return _messages; } It's returning a copy of the QList _messages, rather than a reference to it. I'm not sure that it would give the results you're seeing, but it looks wrong...

How can I access the members of a subclass from a superclass with a different constructor?

c++,inheritance,constructor,subclass,superclass

This map: typedef map<string, Object> obj_map; only stores Object objects. When you try to put an Image in, it is sliced down and you lose everything in the Image that was not actually part of Object. The behaviour that you seem to be looking for is called polymorphism. To activate...

Marshal struct in struct from c# to c++

c#,c++,marshalling

Change this: [MarshalAs(UnmanagedType.ByValTStr, SizeConst = 36)] private string iu; to this: [MarshalAs(UnmanagedType.LPStr)] private string iu; Note that this code is good only to pass a string in the C#->C++ direction. For the opposite direction (C++->C#) it is more complex, because C# can't easily deallocate C++ allocated memory. Other important thing:...

Validate case pattern (isupper/islower) on user input string

c++,user-input

The simplest thing you can do is to use a for/while loop. A loop will basically repeat the same instruction for a number of n steps or until a certain condition is matched. The solution provided is pretty dummy, if you want to read the first name and last name...

template template class specialization

c++,templates,template-specialization

The specialization still needs to be a template template argument. You passed in a full type. You want: template <class Type, class Engine> class random_gen<std::uniform_real_distribution, Type, Engine> { ... }; Just std::uniform_real_distribution, not std::uniform_distribution<Type>. ...

C++ Isn't this a useless inline declaration?

c++,inline,private,member,protected

The Compiler can Access everything. The restrictions are only valid for the programmer. This means there are no restrictions for the Compiler to Access any variables! At the end every variable is just translated to an address which can be accessed. So for the Compiler it is no Problem to...

how to sort this vector including pairs

c++,vector

You forgot the return statement in the function bool func(const pair<int,pair<int,int> >&i , const pair<int,pair<int,int> >&j ) { i.second.first < j.second.first ; } Write instead bool func(const pair<int,pair<int,int> >&i , const pair<int,pair<int,int> >&j ) { return i.second.first < j.second.first ; } Also you should include header <utility> where class std::pair...

Checking value of deleted object

c++

It is very bad, accessing deleted objects as if they were not deleted will in the general case crash. There is no guarantee that the memory is still mapped inside the process and it could result in a virtual memory page fault. It is also likely that the memory will...

std::condition_variable – notify once but wait thread wakened twice

c++,multithreading

Converting comments into answer: condition_variable::wait(lock, pred) is equivalent to while(!pred()) wait(lock);. If pred() returns true then no wait actually takes place and the call returns immediately. Your first wake is from the notify_one() call; the second "wake" is because the second wait() call happens to execute after the Stop() call,...

Confused about returns in stack template

c++,templates,generic-programming

This depends on what you want the behaviour (protocol) of your class to be. Since you're logging into the error stream there, I assume you consider this an error condition to call pop() on an empty stack. The standard C++ way of signalling errors is to throw an exception. Something...

Translating a character array into a integer string in C++

c++,arrays,string

If you want a sequence of int, then use a vector<int>. Using the key_char string, the values of the chars in it will serve as the initial value of the ints. std::vector<int> key_num(key_char.begin(), key_char.end()); Then, iterate over each character of key_num and convert it to the equivalent int value for...

Is there standard implementation for thread block/resume in java SE?

java,multithreading,wait

have a a look at the classes in java.util.conucurrent ... CountDownLatch might be a solution for your problem if i understand your problem correctly.

pointer to pointer dynamic array in C++

c++,arrays,pointers

The valid range of indices of an array with N elements is [0, N-1]. Thus instead of for example this loop for (int i=1; i <= n; i++) ^^^^ ^^^^^^ you have to write for ( int i = 0; i < n; i++ ) As you used operator new...

ctypes error AttributeError symbol not found, OS X 10.7.5

python,c++,ctypes

Your first problem is C++ name mangling. If you run nm on your .so file you will get something like this: nm test.so 0000000000000f40 T __Z3funv U _printf U dyld_stub_binder If you mark it as C style when compiled with C++: #ifdef __cplusplus extern "C" char fun() #else char fun(void)...

Why are shaders and programs stored as integers in OpenGL?

c++,opengl,opengl-es,integer,shader

These integers are handles.This is a common idiom used by many APIs, used to hide resource access through an opaque level of indirection. OpenGL is effectively preventing you from accessing what lies behind the handle without using the API calls. From Wikipedia: In computer programming, a handle is an abstract...

3 X 3 magic square recursively

c++,algorithm,math,recursion

Basically, you are finding all permutations of the array using a recursive permutation algorithm. There are 4 things you need to change: First, start your loop from pos, not 0 Second, swap elements back after recursing (backtracking) Third, only test once you have generated each complete permutation (when pos =...

opencv window not refreshing at mouse callback

c++,opencv

your code works for me. But you used cv::waitKey(0) which means that the program waits there until you press a keyboard key. So try pressing a key after drawing, or use cv::waitKey(30) instead. If this doesnt help you, please add some std::cout in your callback function to verify it is...

Incorrect Polar - Cartesian Coordinate Conversions. What does -0 Mean?

c++,polar-coordinates,cartesian-coordinates

You are converting to cartesian the points which are in cartesian already. What you want is: std::cout << "Cartesian Coordinates:" << std::endl; std::cout << to_cartesian(to_polar(a)) << std::endl; std::cout << to_cartesian(to_polar(b)) << std::endl; //... Edit: using atan2 solves the NaN problem, (0, 0) is converted to (0, 0) which is fine....

Issue when use two type-cast operators in template class

c++

What you're trying to do makes little sense. We have subclass<int>. It is convertible to int&, but also to a lot of other reference types. char&. bool&. double&. The ambiguity arises from the fact that all the various overloads for operator<< that take any non-template argument are viable overload candidates...

How can I tell clang-format to follow this convention?

c++,clang-format

Removing BreakBeforeBraces: Allman Seems to do what you want (for me). I'm using SVN clang though. Although you probably wanted it there for a reason. According to the clang-format docs, the AllowShortBlocksOnASingleLine should do exactly what you want (regardless of brace style). This might be a bug in clang-format....

No match for 'operator*' error

c++,c++11

As @101010 hints at: pay is a string, while hours_day is a float, and while some languages allow you to multiply strings with integers, c++11 (or any other flavor of c) doesn't, much less allow strings and floats to be multiplied together.

Passing iterator's element to a function: wrong type of pointer

c++,pointers,stl,iterator

Preferred option: change isPrime to take a long (and pass *it to it). Secondary option: pass &*it instead of it. Your original code doesn't work because it is an iterator (which is a class) whereas the function expected long int * and there is no implicit conversion from iterator to...

Passing something as this argument discards qualifiers

c++,c++11

There are no operator[] of std::map which is const, you have to use at or find: template<> struct Record::getDispatcher<std::string> { static std::string impl(Record const& rec, std::string& const field) { return rec.fieldValues_.at(field); // throw if field is not in map. } }; or template<> struct Record::getDispatcher<std::string> { static std::string impl(Record const&...

C++ template template

c++,templates

Your issue is that std::deque (and other standard containers) doesn't just take a single template argument. As well as the stored type, you can specify an allocator functor type to use. If you don't care about these additional arguments, you can just take a variadic template template and be on...

Can python script know the return value of C++ main function in the Android enviroment

python,c++

For your android problem you can use fb-adb which "propagates program exit status instead of always exiting with status 0" (preferred), or use this workaround (hackish... not recommended for production use): def run_exe_return_code(run_cmd): process=subprocess.Popen(run_cmd + '; echo $?',stdout=subprocess.PIPE,shell=True) (output,err)=process.communicate() exit_code = process.wait() print output print err print exit_code return exit_code...