DMSP Isabelle Prim
email: isabelle.prim@dauphine.fr.

Bernard Pras

Professor

Peter S Ring

«FRIENDLY» COMPLAINING BEHAVIORS: TOWARD A RELATIONAL APPROACH

Keywords: Relationship Marketing, Complaining Behavior, Friendly Complaints, Justice Theory

The relational approach is often presented as a strategy to retain customers, but it may also be an appropriate approach to encourage customers to complain, as a review of literature shows.

Using information contained in complaints and giving the right answers (distributive, procedural and interactional) to such complaints is essential. Relational marketing may also be used to induce customers (but not all of them) to complain about the attributes of certain products/services. This article focuses on these issues and should stimulate further research in this new field.

Introduction

Over the last two decades, the nature of supplier-customer relationships has changed dramatically in many markets. Managing relationships with customers and suppliers is an issue which has also gained widespread recognition, among academic researchers as well as among industrial practitioners, mainly in industries such as the car industry or through distribution channels. Even though this evolution does not appear in all industries, it is a movement that cannot be ignored.

It becomes apparent that many industrial firms try to establish close relationships with their business partners -suppliers, distributors, clients, etc. - [START_REF] Heide | Explaining « Closeness » in Industrial Purchasing Relationships: The Effects of dependence Symmetry on Inter-Organizational[END_REF]. This relational trend is complementary and sometimes opposite to the one based on transactions. [START_REF] Macneil | The New Social Contract, An inquiry to Modern Contractual Relations[END_REF] speaks of a transactional/relational exchange continuum. Are these distinct positions, or should we instead consider a balance between the two tendencies? With a discrete contract, no relation exists between the parties apart from the simple exchange of goods. Its paradigm is the transaction of the neoclassical microeconomics (Macneil, 1980, p.10). One way or another, modern contractual relations tend to involve large numbers of people, be durable, and involve more aspects than the mere object of the contract. There is a social dimension to the exchange that is totally ignored in the transactional exchange [START_REF] Dwyer | Developing Buyer-Seller Relations[END_REF].

But the transactional as well as the relational dimensions are taken into account which lead to the idea of a continuum.

Although traditional marketing with the transactional dimension remains important, the relational dimension is fundamental in marketing today, but not for all economic sectors or customers. Some economic sectors are not concerned by this new type of marketing (e.g., industries where you buy very rarely, for instance, the real estate industry). Some customers are not ready to implement relationship marketing (e.g., they have a very low potential for long-term relationships). What one customer may consider a warm, close, «friendly» relationship, will be considered as stifling or unnecessary by others. For instance, in a focus group conducted by Barnes (1997, p.771), a bank customer indicated that he had an ideal relationship with his bank. When asked to describe the relationship, he replied: « I don't call them, and they don't call me! ».

The relational approach is often presented as a strategy to retain customers. Most firms try to increase their customer loyalty rate by developing a strategy of relationship marketing. Indeed, it has been estimated that by decreasing customer defection among dissatisfied customers by just 5 percent, a firm can achieve profit improvement of 25 to 85 percent [START_REF] Reichheld | Zero Defections: Quality Comes to Services[END_REF].

Retaining customers has become a primary concern. The success of most industrial products depends on repeat buying, and companies do their best to respond to « valid » complaints and to restore satisfaction whenever possible [START_REF] Barksdale | Complaint Voicing by Industrial Buyers[END_REF].

Customer loyalty is extremely important to companies since:

1. switching creates costs to the firm: the future revenue stream from the customer (a loss from the high margin sector very often [START_REF] Keaveney | Customer Switching Behavior in Service Industries: An Exploratory Study[END_REF]); 2. the investment to get new customers is far more costly than the one required to retain old ones [START_REF] Reichheld | Zero Defections: Quality Comes to Services[END_REF].

The study of consumer complaining behavior and its consequences is essential and critical in the explanation and prediction of consumer repurchasing intentions and brand loyalty for three reasons [START_REF] Singh | Consumer Complaint Intentions and Behavior: Definitional and Taxinomical Issues[END_REF][START_REF] Day | Modeling Choices among Alternative Responses to Dissatisfaction[END_REF][START_REF] Engel | Consumer Behavior, 4th Ed[END_REF][START_REF] Richins | Negative Word-of-Mouth by Dissatisfied Consumers: A Pilot Study[END_REF]:

1. Complaints provide valuable information about the firm's dysfunctional behaviors and its products/services failures [START_REF] Dart | Dissatisfaction Response Styles among Clients of professional Accounting Firms[END_REF][START_REF] Hansen | Encouraging « Friendly » Complaint Behavior in Industrial Markets: Preventing a Loss of Customers and Reputation[END_REF]. If a customer leaves without saying anything, he alone retains this essential information. Furthermore, voice complaints can be very interesting and constructive.

2. Moreover, if the firm answers his/her complaint properly and rapidly, the consumer can reach a second-order satisfaction. [START_REF] Bearden | The Role of Public and Private Complaining in Satisfaction with Problem Resolution[END_REF] have shown that the very act of complaining may enhance secondary satisfaction, not only through its ability to initiate redress, but also through its cathartic effects of « getting it off my chest » [START_REF] Oliver | Satisfaction: A Behavioral Perspective on the Consumer[END_REF].

3. Overall, it has been shown that when a consumer is satisfied by the response to his complaint, he becomes more loyal than other consumers [START_REF] Lapidus | Customer Complaint Situations: An Equity Theory Perspective[END_REF].

Despite this growing interest in consumer complaining behaviors, very few researchers (with a few exceptions -e.g., Fornell & Wernerfelt, 1988;[START_REF] Dacin | Understanding and Influencing Consumer Complaint Behavior: Improving Organizational Complaint Management[END_REF][START_REF] Oliver | Satisfaction: A Behavioral Perspective on the Consumer[END_REF] argue that it is important to encourage consumers to complain. This same pattern is true with the investigation of appropriate organizational responses to various complaint behaviors.

As previously stated, using information contained in complaints and giving the right answers to such complaints is important. How can a company give the right answers and stimulate a « positive » complaining behavior from the consumer? The objective is to stimulate complaining behavior leading to useful information. A «friendly» complaint provides signals about the need to reconsider attributes important to the company and on which the company believes it is able to respond quickly.

The literature on consumer complaining behaviors, on justice theory, and on relationship marketing tends to suggest that relationship marketing is a way to develop «friendly» complaints, even if some precautions must be observed.

Consumer complaining behavior conceptualization

The literature explaining complaining behaviors by consumers belongs to a larger tide of research about consumer satisfaction and dissatisfaction [START_REF] Bearden | An Investigation of Influences on Consumer Complaint Reports[END_REF][START_REF] Day | Modeling Choices among Alternative Responses to Dissatisfaction[END_REF][START_REF] Day | Toward a Theory of Consumer Complaining Behavior[END_REF][START_REF] Jacoby | The Sources, Meaning, and Validity of Consumer Complaint Behavior: A Psychological Analysis[END_REF]. In this paper, the focus is on consumer complaining behaviors and not on the sources of dissatisfaction and their elimination. Not all complaints and complaining behaviors are considered hereafter. Firstly, among complaining behavior, only « justified » or « valid » complaints are of interest in the « friendly » complaints field. It is well known that some customers consider complaining as part of a game in order to make the best possible deal with the supplier [START_REF] Bearden | An Investigation of Influences on Consumer Complaint Reports[END_REF][START_REF] Kowalski | Complaints and Complaining: Functions, Antecedents, and Consequences[END_REF]; these behaviors are well documented through the game theory -these customers complain at any time, for any reason-.

Secondly, only the complaints which are explicitly expressed to the firm are considered. They are the only ones which give an organization the opportunity to gather enough information to respond quickly and positively to a service/product deficiency or failure and to up-date the product/service information system. These « justified » and « explicit » complaints can be useful in Business-to-Business and Business-to-Consumer contexts. No systematic distinction between these two domains is required for our discussion (when the developments are valid in both types of situations).

Definition and taxonomies

Though several definitions of complaining behavior have been proposed, there is a general consensus about the conceptual meaning of consumer complaining behavior.

Consumer complaining behavior is believed to be triggered by feelings or emotions of perceived dissatisfaction [START_REF] Day | Modeling Choices among Alternative Responses to Dissatisfaction[END_REF]. Without this feeling of dissatisfaction, the complaint cannot be considered as a real complaining behavior, but as a « game theory » behavior and a « negotiation » tool.

Complaining behaviors triggered by a perceived dissatisfaction may result in some action being taken or in no action being taken. In the first case, complaining behavior is named « behavioral complaining behavior » (exit, voice, third party, e.g. any consumer actions that convey an expression of dissatisfaction). In the second case, it is named « non-behavioral complaining behavior » (there is no action -the consumer tries to forget the dissatisfaction and remains loyal-). This distinction behavioral/action -non behavioral/no action was first identified by [START_REF] Hirschman | Exit, Voice and Loyalty, Responses to Decline in Firms, Organizations and States[END_REF]. All the subsequent taxonomies (except the one by [START_REF] Richins | Negative Word-of-Mouth by Dissatisfied Consumers: A Pilot Study[END_REF]) are based on this distinction between behavioral and non behavioral complaints.

The « friendly » complaint needs to be explicit. In other words, it has to be voiced and the company has to be able to respond to it. That is the reason why, among the following taxonomies, the complaints of interest in a «friendly» complaint perspective will be identified. a) Hirschman (1970, p.81) theorizes that three options face a dissatisfied consumer: exit 1 (i.e. leave the relationship), voice (i.e. communicate a displeasure to the institution) and loyalty. Hirschman introduces the concept of loyalty as a third response. This concept can be twofold: on the one hand, it can be an outcome itself whereby individuals choose loyalty (staying) rather than exit or voice. On the other hand, loyalty can be a moderating variable where high-loyalty individuals forego exit until all efforts at voice have failed [START_REF] Oliver | Satisfaction: A Behavioral Perspective on the Consumer[END_REF].

b)

Based on this work, [START_REF] Day | Toward a Theory of Consumer Complaining Behavior[END_REF] propose a two-level hierarchical classification scheme. The first level distinguishes behavioral from non behavioral responses (i.e., action/ no action). The second level represents the distinction between public (e.g., third party, large audience, etc.) and private action (e.g., boycott, family circle, etc.). This classification has not been tested.

c)

In order to improve Day's & Landon's work, Day (1980) suggests another basis for classification at the second level of the previous taxonomy (See Figure 1). He notes that consumers complain (or do not complain) to achieve specific objectives. In fact, consumers may provide various explanations for the complaint action they undertake. So Day proposes that the 'purpose' of complaining can be used to classify consumer behavior into three categories:

1. seeking redress;

2. complaining for reasons other than seeking remedies (to affect future behavior, to persuade others, to take legal action, etc.);

3. as a means of expressing personal boycott.

Please, Insert FIGURE 1 d) [START_REF] Singh | Consumer Complaint Intentions and Behavior: Definitional and Taxinomical Issues[END_REF] introduces a slightly different taxonomy. When dissatisfaction occurs, three types of complaining behavior can be found:

1. voice responses (seeking redress from the seller or no action 2 ); 2. private responses (word-of-mouth communication);

3. third-party responses (implementing legal action).

e) Finally, [START_REF] Richins | Negative Word-of-Mouth by Dissatisfied Consumers: A Pilot Study[END_REF] notes that complaining involves at least three distinct activities: (a) switching brands/stores/service providers (defined as exit), (b) making a complaint to the seller (defined as voice), and (c) telling others about the unsatisfactory experience (which constitutes negative word-of-mouth). She indicates a fourth possibility which is complete inaction (the consumer keeps his/her dissatisfaction to himself/herself and tries to forget the experience). Most of the customers opt for this fourth possibility. Except for complete inaction, these answers can be linked together.

The dimensions operationally useful in a «friendly» complaints perspective may be illustrated as such:

1. a behavioral complaint which is voiced;

2. a complaint aimed at (1) seeking redress, (2) seeking something other than remedy 3 , or (3) personal boycott. These objectives will be reached through voice responses 4 . For this reason, an objective of the «friendly» complaint strategy is to transform non behavioral complaints into behavioral complaints and non-voiced complaints into voiced complaints [START_REF] Stephens | Why Don't Some People Complain? A Cognitive-Emotive Process Model of Consumer Complaint Behavior[END_REF]) -See figure 2-. These different types of responses are not independent of one another.

Please, Insert FIGURE 2

The attitude and propensity toward complaining

The attitude toward complaining is conceptualized as the overall affect of « goodness » or « badness » of complaining towards sellers and is not specific to a given episode of dissatisfaction.

Research has shown that the consumer's disposition to complain is correlated with some personality orientations including assertiveness, alienation, self-confidence, locus of control and self-monitoring [START_REF] Bearden | An Investigation of Influences on Consumer Complaint Reports[END_REF]. In addition, learning theories suggest that consumers develop more positive attitudes toward complaining when they become more familiar with complaining practices and environment (e.g., knowledge of unfair practices, consumer rights and complaint channels) and when they have faced complaining experiences with positive outcomes. Therefore, in a « friendly » complaints context, the firms have to work on these learning processes since they want to develop complaining behaviors. There exist several ways to do so, such as communicating through advertising the firm's complaint handling policy, giving in-store complaint policy information while the customer is making his choice, communicating through the service warranty in the contract.

The basis of a « friendly » complaint is a « voiced » and « well-handled » complaint.

One must bear in mind Hirschman's observation (1970) that customers may take no action even when dissatisfied if they are loyal to the seller or if they perceive that complaining to the seller (or the supplier) is probably pointless. But the seller may increase the propensity to complain through encouraging the complaint to be voiced and convincing the consumer that the complaint will be taken care of.

A new approach to «friendly» complaints through the justice theory

All complaints do not have the same status. As mentioned, some of them are « calculated » complaints (game theory). Others are motivated by a willingness to improve society's wellbeing. But most of them seem to be the consequence of the customers' willingness to get redress after a disconfirmation or a harmful experience. [START_REF] Hansen | Encouraging « Friendly » Complaint Behavior in Industrial Markets: Preventing a Loss of Customers and Reputation[END_REF] use the concept of « friendly » complaints. In their work, they refer to the adjective « friendly » because « complaining can be desirable for the marketer » (Hansen & al., 1996, p.271-272).

Developing and managing « friendly » complaints deserve attention since a firm which does not receive any claims probably faces important problems of upward information.

But « friendly » complaints may be desirable only in some cases. And it is valuable for the marketer to develop complaining behavior in these specific cases.

The justice theory helps in understanding the customers' expectations when complaining. The justice theory, and its potential applications in the context of « friendly » complaints, is presented below.

«Friendly» complaints: Two conditions for being efficient?

A « friendly » complaint has to attain its objective that is, be desirable for the marketer. The desirability depends on:

1. the quality and quantity of information generated by the complaints;

2. the capacity of the firm to respond positively and quickly to these complaints.

The quality of the complaints is related to the reliability of the information provided (that is, the accuracy of the information, the credibility of the source, in other words, the credibility of the customers who complain). The quantity of information rests on the capacity of the firm to stimulate the customers' « complaining » behavior in two ways: (1) sending systematically useful information about one's firm or its competitors;

(2) inciting « shy » customers or customers who do not express themselves to complain «friendly» (See Figure 2).

The capacity of the firm to respond to these complaints signifies that «friendly» complaints make sense only when the firm is capable of responding quickly and efficiently to these complaints. A «friendly» complaint may have an opposite effect if the company's capacity to respond is weak.

Therefore, in order for a complaint to be qualified as « friendly », certain conditions must be met: They are:

• the complaint must be from a reliable customer;

• the complaint must relate to attributes on which the firm has a strong capacity of response or a competitive advantage. Whenever these conditions are not fulfilled, it can be assumed that « friendly » complaints will have more negative than positive effects as the literature review on justice theory leads us to believe.

Finally, a « friendly » complaint strategy makes sense if it is applied to segments of customers who are very important to the firm, i.e. « high margin » customers or « long-term » customers.

Indeed, it is important to give high margin customers the perception that the firm cares about their problems and concerns [START_REF] Reichheld | Zero Defections: Quality Comes to Services[END_REF]) and provide long-term customers with a good knowledge of the firm, of its products/services. Their feedback is highly valued.

Justice theory and «friendly» complaints

The focus on « good » (reliable, high-margin, long-term) customers may be excellent but risky since the firm has to give the right answer at the right time; otherwise this « friendly » complaint policy might become very dangerous.

The justice literature suggests that each stage of the sequence of a buying and a complaining process is subject to fairness considerations and that each aspect of complaint resolution creates a justice episode [START_REF] Tax | Customer Evaluations of Service Complaint Experiences: Implications for Relationship Marketing[END_REF]. A literature review is important to help in understanding the responses to be given to « friendly » complaints and the sensitive dimensions of a complaining behavior. [START_REF] Tax | Customer Evaluations of Service Complaint Experiences: Implications for Relationship Marketing[END_REF] identify three dimensions to justice: distributive, procedural and interactional. Perceived justice has been found to be the main determinant of complainants' negative word-of-mouth behavior and their repatronage intentions. It has been found to mediate the effects of likelihood of success, as well as attitude toward complaining, product importance, and stability and controllability over complaining behavior [START_REF] Blodgett | The Effect of Perceived Justice on Complainants' Negative Word-of-Mouth Behavior and Repatronage Intentions[END_REF].

Theories based on distributive justice focus on the allocation of benefits and costs. They use equity principles and the customer may assess the fairness of the compensation differently on the basis of his (her):

• prior experience with the firm in question and/or other firms;

• awareness of other customers' resolutions;

• and perceptions of his (her) own loss [START_REF] Tax | Customer Evaluations of Service Complaint Experiences: Implications for Relationship Marketing[END_REF].

Distributive justice responds to economic costs but also to emotional costs and therefore the response itself will be economic (and based on equity) and affective. Distributive justice will then be found in the evaluation of compensation for financial loss (economic response) and an apology. The apology will give satisfaction on the emotional side. The importance of the apology may differ from Business-to-Business to Business-to-Consumer 5 . However, distributive justice cannot be analyzed without considering also procedural justice and interactional justice. Customers who are voicing « friendly » complaints are mainly looking for procedural and interactional justice, whereas customers who are voicing « usual » complaints are more interested in distributive justice, i.e. they are more interested in what they will get out of their complaint (See Table 1).

Please, Insert TABLE 1

Procedural justice is the perceived fairness of the means by which the ends are accomplished.

Here, the process is more important than the outcome. Procedural justice aims at resolving conflicts in ways that encourage the continuation of a productive relationship between the disputants, even when the outcome is unsatisfactory to one or both parties [START_REF] Folger | Distributive and Procedural Justice in the Work Place[END_REF][START_REF] Tax | Customer Evaluations of Service Complaint Experiences: Implications for Relationship Marketing[END_REF].

More specifically, according to Tax, Brown & Chandrashekaran (1998, p. 72), procedural justice reflects certain aspects of customer convenience and firm responsibility, follow-up and accessibility, with the adage « justice delayed is justice denied ». Swiftness and accessibility, among other attributes, are important.

Finally, interactional justice refers to the fairness of the interpersonal treatment people receive during the enactment of procedures. These interactional factors might help us understand why some people feel unfairly treated even though they would characterize the decision making procedure and the outcome as fair.

This component takes into account the human part of the relationship as well as the quality of the communication between the firm and the complaining party. Research results point out several aspects of (un)fair communication and behavior, such as honesty, empathy and politeness [START_REF] Tax | Customer Evaluations of Service Complaint Experiences: Implications for Relationship Marketing[END_REF].

The economic aspect is a necessary condition for all complaints to be satisfied. In « friendly » complaints, the relational aspects will be prominent. « Friendly » complaints are more likely to occur in a relational context. In this context, the consumer will expect to be satisfied on the three dimensions of justice, whereas in a transactional context, the consumer will not expect any interactional justice. He/she will focus his/her evaluation of the supplier on the distributive dimension of justice (which is an economic dimension, i.e. « How much do I get if I complain? »). The procedural dimension may also be of importance (See Table 2).

Limits

One of the most important limits to the strategy of developing complaints, even though they are «friendly», is that this policy may increase the customer's expectations. The more a firm responds well and quickly to complaints, the more the consumer expects this firm to do so and the more it limits this type of strategy to its attributes with the strongest competitive advantages, otherwise customers may switch to competitors able to provide a better product/service. For this reason, it may be wise for firms to limit this policy to their best customers, who tend to be more loyal 6 . By doing so, it will also limit the percentage of fraudulent claims [START_REF] Stephens | Why Don't Some People Complain? A Cognitive-Emotive Process Model of Consumer Complaint Behavior[END_REF].

Developing strong relationships with customers and giving them the feeling of being unique may represent a danger to the firm. The customer will increase his expectations according to the advantages he may get through this status.

If responses have to be quick (procedural justice), and delivered in a nice manner (interactional justice), they must also be fair (distributive justice). Tax et al. (1998, p.62) give an interesting example by illustrating what occurred to Domino's Pizza: Their decision to change their service guarantee from « Delivery within 30 minutes or receive a free pizza » to « Delivery within 30 minutes or $3.00 off the purchase price » was explained as an equity decision. Indeed, the first response (a free pizza) was too generous; some customers were embarrassed by this policy.

It seems also clear that the «friendly» complaint developing strategy should be limited to markets where the competition is not too strong. We will see in a third part how the theory of justice may be useful in managing complaints and more specifically «friendly» complaints.

Which strategy the supplier should implement to develop friendly complaints

Why sellers should encourage consumers to complain

Complaints may encourage the firm to improve its products/services so that customers as well as the firm itself benefit.

-Complaining behavior as a source of information about one's own firm and competitors

Complaints are a source of information (1) about one's company, and (2) about competitors.

It may seem obvious that encouraging complaining is probably a good way to discover consumer problems and concerns. An analysis of complaints allows the seller to understand the buyer's perception of the company's products and/or services. Firms may get useful information on product/service quality, delivery system weaknesses, etc.

A company may also gather information about competitors through the service guarantee system. Let us take the example of firms which make the following promise: « we are the cheapest on the market », and offer to refund the difference in price if the customer finds the same product cheaper anywhere else. Darty, a French household appliance distributor, built part of its success on this slogan. In this case, the seller is looking toward getting « complaints » from the customers. Indeed, coming back to say « I have found this product at a better price... » has the customer come back with information on the competitors. This constitutes part of the strategy of commercial intelligence.

-Complaints as a means of communication and loyalty building

Beyond the information side, there are other advantages for the firm in encouraging customers to complain. Managing complaints can be a means (1) of communication, and (2) of loyalty maintenance and building.

Complaints are a means of communication between the firm and its customers. According to [START_REF] Granovetter | Economic Action and Social Structure: The Problem of Embeddedness[END_REF], this communication helps build trust between the parties, which in turn helps to determine the predictability of the other party's intentions and actions. In this sense, it may be considered as a factor of stability in the relationship.

It may be assumed that the more the complaints are voiced to the firm, the less the customer will spread negative word-of-mouth complaints. Furthermore, the more the complaints are positively and quickly responded to, the more the customer is likely to spread positive wordof-mouth information. Fornell and Wernerfelt (1987, p. 344) observe that « data available indicate that customer loyalty can be increased by encouraging customers to complain ». Likewise, the emerging view in relationship marketing and service quality literature emphasizes « recovery » which can be defined as converting dissatisfied consumers into satisfied customers (hence recovered)

through quick and exemplary responsiveness by service providers [START_REF] Lapidus | Customer Complaint Situations: An Equity Theory Perspective[END_REF]. In turn, this recovery rests on managers understanding of the consumers' complaints, since recovery cannot occur without a complaint. Costs over a short-term period might be high but they should be recovered in the mid and long term with an increase in the loyalty rate and in the supplier's reputation.

How the supplier can encourage consumer complaining behaviors

It has been shown [START_REF] Barksdale | Complaint Voicing by Industrial Buyers[END_REF] that good handling of complaints is very constructive. The firm must communicate its openness to receiving complaints [START_REF] Stephens | Why Don't Some People Complain? A Cognitive-Emotive Process Model of Consumer Complaint Behavior[END_REF]. However, an optimum level of communication has to be found, since too much emphasis on this openness might make the consumers suspicious about product/service quality and the firm's delivery capabilities.

The transactional and relational aspects both influence the complaining behavior. Some are controllable by the firm, others are less controllable. There are several points of the transaction that the seller can rarely influence. For instance, [START_REF] Barksdale | Complaint Voicing by Industrial Buyers[END_REF] show that there is a link between the overall purchase price and the propensity to complain. They also demonstrate a negative link between the number of potential suppliers and the propensity to complain. These elements (the weight given to a product price in the customer portfolio or the number of players in the market) are not directly controlled by the seller. [START_REF] East | Identifying the Drivers of Complaining[END_REF] reminds us that predictive models of complaining behaviors may include demographics, biographical and situational data as factors affecting complaining behaviors.

The only tactic he may use is to communicate in such a way that the weight of these elements is decreased in the consumer's mind, but most of all, the company has to know these factors and take them into account in its strategy.

There are other points which are easier to manipulate by the supplier who wants to increase the consumer propensity to complain. The consumer must have the feeling that complaining is worth the effort [START_REF] Singh | Voice, exit and Negative Word-of-Mouth: An Investigation Across Three Service Categories[END_REF], that he will be listened to and understood. He should also be convinced that the incident was unfortunate and is not the general rule (especially when the customer attributes the origin of the problem to the supplier). The explanative models help define a strategy. Indeed, explanative models focus on motivational factors and seem more useful to those whose goals are to change customers' behaviors through communication.

Ajzen's theory (1991) of planned behaviors is very useful in identifying and exploring the antecedents of complaining behaviors. In this theory, three types of belief affect the probability of complaining behaviors by dissatisfied consumers:

• Outcome beliefs are about gains and losses, including opportunity costs, that follow complaining or not. These beliefs will be analyzed through the customer's distributive justice scheme.

• Referent beliefs are linked to the importance a customer can give to others' opinion. It is about what key people or groups think the customer should do.

• Control beliefs are linked to procedural justice. They evaluate how easy or difficult it is for the customer to complain. These are knowledge, skills and other resources that can make complaining easier. Among the other resources are key points set by the customer's environment, such as whether adequate personnel are accessible, etc. It is in some way linked to interactional justice.

In order to improve the probability of customers exhibiting complaining behaviors when dissatisfaction occurs, a firm has to play on these three levers in its communication as well as in its complaint management. The customers should be made keenly aware of the different means they have at their disposal to complain (e.g., where to go and how to complain). The French railway company S.N.C.F. used the control belief lever (procedural justice) during the winter of 1996, when most of its trains were halted by snowstorms and icestorms. Agents waited in the train stations (interactional justice) and handed out pre-stamped envelopes with the address of the customer service department to the high speed train (T.G.V.) passengers in order to get reimbursed (distributive justice 7 ). This action, even if it was quite expensive for the company, clearly helped it to recover and improve its image after this difficult time.

This suggests that attempts to promote face-to-face, mail or other medium complaints should be designed to banish unease and make such behaviors a simple, common and acceptable part of the customers' repertoire [START_REF] East | Identifying the Drivers of Complaining[END_REF].

The impact of relationship marketing on consumer complaining behaviors

According to [START_REF] Day | Toward a Theory of Consumer Complaining Behavior[END_REF], factors influencing the choice of alternative actions between exit, voice and loyalty (according to Hirschman's terminology) are:

marketing aspects (seller's reputation, circumstances of the sale, responsiveness of the seller to complaints, etc.);

consumer factors (e.g., propensity to complain);

circumstantial factors.

As explained previously, it seems clear that the firm may act mainly on the first two categories of factors, i.e. the marketing aspects and consumer factors. Relationship marketing has shown its capability to create strong links between service providers and customers 8 .

For this reason relationship marketing provides useful tools to handle «friendly» complaints.

Through the literature, one can identify four main characteristics of the relationship process [START_REF] Perrien | Dissolution of a Relationship, The Saleforce Perception[END_REF]:

• it is an asymmetrical process guided by the supplier, i.e. the overall quality and effectiveness of any relationship rests on the seller. One of the objectives of a « friendly » complaint developing strategy is to balance this asymmetry ;

• it is a personalized process: buyer's knowledge is a basic requirement to the development of a relationship;

• it is a process with shared benefits: both the seller (supplier) and the buyer must obtain a tangible outcome. Securing the desired quality of supply and ensuring customer loyalty are two of the most well known results expected from an effective relationship.

• and it is a process which requires a long-term commitment: basic objectives of a long-term relationship strategy are both to increase the customer's loyalty and to achieve some growth by cross-selling and developing relationships with high-margin customers;

It is useful, then, to analyze how relationship marketing can help develop «friendly» complaints and why it works (in other words, to assess the gains both parties can obtain in the exchange from this strategy).

Moreover the consequences of a « friendly » complaints program may be on the communication side or on the relations perspective as shown in Table 3. This table indicates which aspects of the relationship the supplier has to improve in order to develop « friendly » complaints. For instance, because the relationship is an asymmetrical process, some resources have to be spent in order to develop a bi-directional communication. The supplier can use different means to reach these objectives. All are linked to relationship marketing and can be decomposed into three categories linked to the three dimensions of justice. The gray cells, in Table 3, are not empty cells. They simply represent the fact that the link is less strong between the asymmetrical process and the type of relation (and between the three other relational characteristics and the communication side). We have only focused on the dominant links.

Please, Insert TABLE 3 -How to develop « friendly » complaints? Through relationship marketing

The more the consumer blames the seller ( attribution theory, e.g. [START_REF] Richins | Negative Word-of-Mouth by Dissatisfied Consumers: A Pilot Study[END_REF][START_REF] Boshoff | Empowerment, Attribution and Apologising as Dimensions of Service Recovery, An Experimental Study[END_REF], or believes that the dissatisfaction could have been avoided, and that the dissatisfying event is likely to happen again, the stronger his reaction and the greater his probability to complain. But in such circumstances, it will be difficult for a firm to repair the damages. Developing complaining behaviors in a less radical way would then be more interesting and efficient, which is the basis of « friendly » complaints. It would also decrease the asymmetry of the relationship.

« Friendly » complaints are direct and voiced actions to the supplier that aim to improve the product/service and the delivery process without threatening the relationship, and which ultimately help to strengthen the firm's competitive position.

To elicit such behavior the firm must:

have a better understanding of its customers;

develop bi-directional communication;

build trust;

give good and quick responses to complaints;

develop a relational way to manage its employees 9 .

All of these elements are strengthened by relationship marketing. Indeed, [START_REF] Grönroos | From Marketing Mix to Relationship Marketing: Towards a paradigm Shift in Marketing[END_REF] defines the most important elements of relationship marketing as:

the promise concept: a firm that is preoccupied with making promises may initially attract new customers and build relationships. However, if promises are not kept, the evolving relationship cannot be enhanced and maintained;

trust: there has to be a belief in the other partner's trustworthiness which results from the expertise, reliability or intentionality of that partner [START_REF] Moorman | Relationships between Providers and Users of Market Research: The Role of Personal Trust[END_REF].

It implies a personalization of the relationship. [START_REF] Macneil | The New Social Contract, An inquiry to Modern Contractual Relations[END_REF] also enhances the concepts of communication and flexibility in the relationship. [START_REF] Reichheld | The Loyalty Effect[END_REF] explains that there is no relationship marketing without any internal relationship management.

When a relationship is well established between a seller and a customer, the commitment of both parties is quite important. This long-term commitment may be a driver. In other words, the customer involved in the relationship will want to improve it and to help the supplier improve it as well.

-Why does it work? Shared benefits of « friendly » complaints Rationally, a strategy must be interesting for at least one party in order for it to survive on a long-term basis. Relationship marketing goes further since it is supposed to deliver benefits to all the parties involved. In the « friendly » complaint case, both parties perceive this relationship in its continuity. This means that the client wants to improve his/her future via complaints and the firm (the provider) wants to preserve its future revenues by retaining a satisfied customer. Both can reach their objectives by developing a system of «friendly» complaints.

In the case of services, their intangible nature makes it difficult for customers to evaluate them prior to purchase. For this reason, it is very important to give the consumer the feeling that he can report any problem to the seller and that it is not his fault if something goes wrong. The major development of numerous charge free phone numbers, postage-paid and short feedback cards reflects this view [START_REF] Oliver | Satisfaction: A Behavioral Perspective on the Consumer[END_REF]. However, they are just the beginning since there is a need to develop the relationship on a much more personal level. This is possible through techniques such as the callback: customers are contacted directly and asked pointedly if their product or service experience was satisfactory. In this manner, problems can be pinpointed, and additional diagnostic queries can be introduced [START_REF] Oliver | Satisfaction: A Behavioral Perspective on the Consumer[END_REF]. There is also a great need to develop a relational way to manage employees. They need to be given more independence and responsibility in order to be able to answer consumer requests and subsequently act on them quickly and efficiently.

Please, Insert FIGURE 3

Conclusion

The newness of the «friendly» complaint approach may be summarized as follows:

Traditionally, the complaint treatment process was:

Complaints Relational Management

Better Marketing here it becomes: Better Products/Services Relational Management

Complaints

Better Marketing

Relational Management

It should be noted that working on complaining behaviors must not become a substitute for product/service quality. Response strategies to consumer complaints must remain a last recourse. But in competing economies, product/service quality has considerably improved, so customer service programs may be a key asset in the elaboration of the strategy of differentiation. That is why this issue is important. But when a firm develops a « friendly » complaints program, it cannot become a strategy per se, that is THE strategy of the firm. It may only be an element of a broader program, including pro-active actions such as warranty programs. Moreover, the strategy has to be focused on a limited number of attributes, attributes on which the firm has a competitive advantage for a limited number of customers, those who have a relational potential. Otherwise, a «friendly» complaints program may be very costly.

To summarize, complaints may be considered as positive if the company is able to (1) obtain all of the information contained in complaints, (2) act accordingly, that is, quickly and positively, (3) make some consumers change their behavior, that is, go from an exit behavior or a no-complaint behavior or a personal complaint behavior to a voiced complaint -in a positive way-; and (4) bring the consumer to seek something more than simple redress. This strategy may also find its roots in the social dilemma theory as it is a transformation of consumer non-cooperative behaviors into consumer cooperative behaviors, and of supplier short-term orientation into long-term orientation [START_REF] Willenborg | Relationship Marketing and Consumers: A Social Dilemmas Approach[END_REF].

Actually, the development of appropriate policies and procedures is essential in reducing consumer complaints that damage the seller's reputation [START_REF] Hansen | Vendor relationships as Predictors of Organizational Buyer Complaint Response Styles[END_REF]. These damaging complaints are mainly negative word-of-mouth behavior since the firm cannot control them. Very often the firm is not even aware of them or discovers them when it is too late and when they have already been widely diffused.

Conducting both an investigation on « friendly » complaints and experiments on strategy and behavior would be useful and necessary to test the propositions made in this article and better understand complaining behaviors. Research would probably show that there exist different categories of consumers responding specifically to a « friendly » complaints strategy, some of the customers being more sensitive to the transactional side of complaining, others to the relational or to both the transactional and the relational.

Please, Insert TABLE 4

It is likely also that differences will be found between Business-to-Business and Business-to-Consumers situations. After conducting a preliminary analysis, researchers might conclude that emotional factors are less important in the first context. This could be misleading since the relational part is prominent in a Business-to-Business context and further research is needed to clearly formalize the differences between retail customers and industrial customers. • interactional justice

Transactional marketing

Relationship marketing Notes 1 -We must note that exit is either switching from one brand to another or, more radically, immediately refusing to buy the category of product or service.

2 -Voice consumer complaining behavior is directed to objects that are external to the consumer's social circle (i.e., informal relationships) and are directly involved in the dissatisfaction exchange. [START_REF] Singh | Consumer Complaint Intentions and Behavior: Definitional and Taxinomical Issues[END_REF] includes the no-action responses in this category because they appear to reflect feelings toward the seller.

In contrast, third party consumer complaining behavior includes objects that are external to the consumer, as in the voice consumer complaining behavior, but they are not directly involved in the dissatisfying transaction (e.g., legal agencies, newspapers, etc.).

Finally for the private consumer complaining behavior category the objects are not external to the consumer's social net and are also not directly involved in the dissatisfying experience (e.g., self, friends, relatives, etc.).

Thus the external/not external and involved/not involved criteria are used by [START_REF] Singh | Consumer Complaint Intentions and Behavior: Definitional and Taxinomical Issues[END_REF] to categorize consumer complaining behavior actions into the three categories of the proposed taxonomy.

3 -For instance, as stated by [START_REF] Oliver | Satisfaction: A Behavioral Perspective on the Consumer[END_REF], a second-order satisfaction can be reached just through the cathartic effects of « getting it off my chest ». Moreover, the consumer can be seeking an improvement in the product/service not for the time being but for his/her future purchases.

4 -Only in the sense of seeking redress from the seller, informing and influencing public opinion or relatives.

important for individual customers. Industrial buying decision processes are supposed to be more rational.

6 -'Loyalty' has to be taken here in a broader sense than Hirschman's definition, i.e. longterm customers.

7-Distributive justice was limited here to the economic aspect of the damage (equity). The emotional aspect (apologies) was quite neglected.

8 - Berry (1983) defines relationship marketing as « attracting, maintaining and -in multiservice organizations-enhancing customer relationships ». He stresses in this definition that the attraction of new customers should be viewed only as an intermediate step in the marketing process.

9 -Many firms have tried to develop relationship marketing. Some of them have failed because they have forgotten to develop the same 'philosophy' with their employees. They asked their employees to spend more time with their customers, to handle any incoming complaint carefully, etc. But the way these employees were managed and evaluated had remained the same (number of new customers, productivity, etc.) and they were totally inadequate for the task [START_REF] Reichheld | The Loyalty Effect[END_REF].

  FIGURE 2Possible scope of «friendly» complaints

TABLE 1 Types of justice and types of complaints

 1 Need for...

	Type of justice	distributive	procedural	interactional
	Usual complaints	important	less important	less important
	«Friendly»	less important	important	important
	complaints

TABLE 2 Types of justice, their prominent dimensions and types of relationship Type of justice

 2

	Type of relationship	distributive	procedural	interactional
	(contract)
	Transactional	economic dimension quickness of the
			response, accessibility
	Relational (*) (**)	emotional dimension fairness of the process all the aspects

(*) All the dimensions present in the transactional type of contract are present in the relational one as well, but if they are necessary they are not discriminant.

(**) « Friendly » complaints belong to this type of relationship.

TABLE 3 Relational characteristics and operational consequences for « friendly » complaints program Consequences for « friendly » complaints program Prominent relational characteristics Communication side Type of relations

 3

	• asymmetrical process	-bi-directional
		communication
		-having any problem being
		reported
		-better understanding of
		customers
		-marketing intelligence
		(competitors)
	• personalized process • shared benefits • long-term commitment	-building trust -good and quick responses (and relationship) -relational way to manage
		employees
		(employees/customers)
	Means	-emotional aspect (distributive justice) -fairness of the means (procedural aspect) -mainly interactional justice
		Relationship Marketing

taxonomy of consumer complaining behaviors (1980) Types of Complaint (or Consumer Reactions)

		Action	No action
		Behavioral	Non
			Behavioral
	Redress	Complaint for reasons	Personal boycott
	seeking	other than seeking remedy

-Further research is needed in this field, but we might assume that human factors are more


Frédéric Kuznik
email: frederic.kuznik@insa-lyon.fr

Christian Obrecht

Gilles Rusaouen

Jean-Jacques Roux

LBM Based Flow Simulation Using GPU Computing Processor

Keywords: Lattice Boltzmann Method, Graphical Processing Unit, Fluid Flow, Multi-Threaded Architecture

Graphics Processing Units (GPUs), originally developed for computer games, now provide computational power for scientific applications. In this paper, we develop a general purpose Lattice Boltzmann code that runs entirely on a single GPU. The results show that: (1) simple-precision floating point arithmetic is sufficient for LBM computation in comparison to double-precision; (2) the implementation of LBM on GPUs allows to achieve up to about one billion lattice update per second using single-precision floating point; (3) GPUs provide an inexpensive alternative to large clusters for fluid dynamics prediction.

Introduction

Nowadays, computational methods and related hardware are really inseparable. In fact, the numerical method must fit the hardware architecture to gain benefits from computational possibilities. Of course, the reciprocal is also true: the hardware architecture progress lead the numerical methods than can be used with a reasonable computational cost.

In the last two decades, the Lattice Boltzmann method (LBM) has proved its capability to simulate a large variety of fluid flows ( [START_REF] Succi | The Lattice Boltzmann -For Fluid Dynamics and Beyond[END_REF], [START_REF] Filippova | A Novel BGK Approach for Low Mach Number Combustion[END_REF], [START_REF] Mei | Lattice Boltzmann Method for 3-D Flows with Curved Boundary[END_REF], [START_REF] Guo | A Lattice Boltzmann Model for Convective Heat Transfer in Porous Media[END_REF], [START_REF] Shi | Finite-Difference-Based Lattice Boltzmann Method for Inviscid Compressible Flows[END_REF] ...).

However, it has been recognized that the LBM is both computationally expensive and memory demanding [START_REF] Mattila | Comparison of implementation of the lattice-Boltzmann method[END_REF]. But, because LBM is explicit and generally needs only nearest neighbor information, the method allows a highly efficient parallel implementation using GPU architecture ( [START_REF] Li | Implementing Lattice Boltzmann Computation on Graphics Hardware[END_REF], [START_REF] Tomov | Benchmarking and implementation of probability-based simulations on programmable graphics cards[END_REF]).

Graphics Processing Unit (GPU) is a massively multi-threaded architecture and then is widely used for graphical and now non-graphical computations( [START_REF] Anderson | General purpose molecular dynamics simulations fully implemented on graphics processing units[END_REF]).

The main advantage of GPUs is their ability to perform significantly more floating point operations (FLOPs) per unit time than a CPU (see figure 1).

Fan et al. [START_REF] Fan | GPU Cluster for High Performance Computing[END_REF] used a 32 nodes cluster made of nVIDIA GeForce 5800 ultra for LBM computations. They use the GPU vector operations and stacks of 2D textures for 3D computations. With 32 nodes, they found an efficiency of their implementation of 49.2 Million Lattice Update Per Second (MLUPS).

Tölke [START_REF] Tölke | Implementation of a Lattice Boltzmann kernel using the Compute Unified Device Architecture developed by nVIDIA[END_REF] used a nVIDIA 8800 Ultra graphics card to implement a 2D LBM code. He found very good results with a ratio between GPU time and CPU time of about 23 for the same test case. is made possible by the use of the nVIDIA CUDA C language programming environment. CUDA provides low level hardware access, avoiding the limitations imposed in fragment shaders. It works on the GT200 processor from nVIDIA, and will be supported on future devices [12]. Algorithms developed for this work will be directly applicable to newer, faster GPUs as they become available.

Lattice Boltzmann method

This part is devoted to an overview of the lattice Boltzmann model used for the purpose of this study (LBM).The model is the lattice BGK model (LBGK) from Qian, D'Humières and Lallemand [START_REF] Qian | Lattice BGK for Navier-Stokes Equation[END_REF]. The main hypothesis of the LBGK are:

• Bhatnagar, Gross and Krook approximation (BGK) ⇒ the collision operator is expressed as a single relaxation time to the local equilibrium,

• the Knudsen number is assumed to be a small parameter,

• the flow is incompressible.

The evolution of the density distribution f for a single fluid particle is then given by:

Df Dt = ∂ t f + ξ.∇ f = - f -f e τ (1)
ξ is the microscopic velocity, τ is the relaxation time and f e the Maxwell-Boltzmann equilibrium distribution function.

The macroscopic variables such as density ρ and velocity u:

ρ ( x, t) = f x, ξ, t d ξ (2) ρ ( x, t) u ( x, t) = ξf x, ξ, t d ξ (3)
To obtain the Lattice Boltzmann model, the velocity space must be discretized:

during dt, the distribution function moves along the lattice link d x i = c i dt.

In our simulations, a 9 velocities 2 dimensional (D2Q9) lattice has been used

(figure 2).

After discretization, the evolution equation becomes (figure 3):

f i ( x + c i dt, t + dt) + f i ( x, t) = - 1 τ (f i -f e i ) (4)
The macroscopic variables such as density ρ and velocity u are then given by:

ρ = i=8 i=0 f i (5)
ρ u = i=8 i=0 c i f i (6)
For the two dimensional case, applying the third-order Gauss-Hermite quadrature leads to the D2Q9 model with the following discrete velocities c i where

i = 1...8. The discrete velocities are -→ c 0 = (0, 0), -→ c 1 = --→ c 3 = c(1, 0), -→ c 2 = --→ c 4 = c(0, 1), -→ c 5 = --→ c 7 = c(1, 1), -→ c 6 = --→ c 8 = c(-1, 1).
The equilibrium density distribution function is given by:

f e i = ω i ρ 1 + 3 c i . u c 2 + 4.5 ( c i . u) 2 c 4 -1.5 (u 2 + v 2 ) c 2 (7)
with u = (u, v) and ω 0 = 4/9, ω i = 1/9 for i = 1, 2, 3, 4, ω i = 1/36 for i = 5, 6, 7, 8.

LBM can then be split into collision and propagation steps:

(1) collision:

f i ( x, t * ) = - 1 τ (f i -f e i ) (8)
(2) propagation:

f i ( x + c i dt, t + dt) = f i ( x, t * ) (9)
Of course, the third step necessary for the implementation of LBM is the determination of the boundary conditions. There are two types of boundary conditions for the case tested in our study: wall boundary condition and imposed velocity boundary condition. For the walls, the classical no-slip boundary condition is imposed by the means of bounce-back rules. A prescribed velocity is easily implemented by constantly refilling the boundary nodes with the equilibrium population corresponding to the desired value of flow speed.

The collision step, which is totally local, required about 70% of the total computational time. The propagation step required 28% of the total computation time. On the whole, 98% of the computational time can easily be parallelized.

(these values are extracted from [14]).

3 Programming overview

Hardware architecture

This paragraph is dedicated to a description of the hardware architecture (figure 4).The GPU computing processor hardware is the nVIDIA GTX280 which can easily be included in a standard workstation. The processor is composed of ten thread processing clusters (TPCs), with each broken down into three streaming multiprocessors (SMs). Threads are assigned by the thread scheduler, which addresses directly to each streaming multiprocessor through a dedicated instruction unit; the later then assigns tasks to one of the eight thread (or stream) processors (SPs). On the whole, 1 GPU is composed of 240 processors.

The memory bandwidth is 141.7 GigaByte per second and the available amount of memory is 1.0 GigaByte. The GPU can deliver about 1000 GFLOPS (Giga Floating Operations per Second) which corresponds to about 80 x86 CPU.

CUDA overview

CUDA (Nvidia) is a standard C language extension for parallel application development on the GPU, independently of the hardware target. Some definitions are necessary to understand the CUDA programming features:

• the device is the GPU,

• the host is the CPU,

• the kernel is a function that is called from the host and runs on the device,

• a CUDA kernel is executed by an array of threads (see figure 5).

In CUDA, these independent threads are organized into blocks which can contain from 32 to 512 threads each (see figure 6). A kernel is executed in a grid of thread blocks being indexed by a 2D block id in the form (row,column).

Concerning the device, one thread block is executed by one multiprocessor. In a block, each thread is indexed by a thread id in the form (row,column). Threads in a block are executed by processors within a single multiprocessor. One important consequence is that threads from different blocks cannot cooperate.

Memory access optimization

The memory access of the kernel is an important feature of the implementation performance. A schematic view of the memory access of the device is given figure 7. The closer memory is a set of 32-bit register per processor. The shared memory is on-chip memory, the size being 16KB per multiprocessor.

This memory allows data transfer between threads and is really fast as long as the number of concurrent memory accesses is a multiple of 16. This last precaution allows to avoid bank conflicts.

The global memory, which is the device memory, is large (1.5 GigaByte) but not as fast as shared memory. The host can only read and write the global memory.

In a D2Q9 lattice, each node requires at least 9 × 4 = 36 bytes of memory for single precision computations. Therefore the number of lattice nodes that can be concurrently stored into low latency shared memory is limited to approximatively 450 per multiprocessor, which on the GT200 leads to a lattice of at most 80 × 80. Hence, the use of high latency global memory is unavoidable. To efficiently hide this latency, maximizing the occupancy rate of the multiprocessors is an important issue.

Data layout in global memory can dramatically impact performances. CUDA enabled GPUs are capable of loading or storing memory segments of 32, 64 or 128 bytes in a single memory transaction. Hence, it is possible to reduce the number of global memory accesses as long as two conditions are met:

• coalescence, i.e. neighboring threads should access neighboring data,

• alignment, i.e. addresses should be a multiple of the segment's size.

The former is easily satisfied by using a separate array for each density instead of one array of structures. The later is more problematic because of the propagation step. A careful choice of the lattice's size allows to avoid misaligned memory stores when propagating along one dimension but the problem remains for the second dimension. A possible solution consists in fetching densities into shared memory. Threads can concurrently access shared memory at no cost as long as there are no bank conflicts. This is readily achieved by using a number of threads per block which is a multiple of 16.

When following the shared memory approach, special care has to be taken of densities crossing block boundaries. Incidentally, this additional step ensures global synchronization across thread blocks.

Pseudo-code

The first step of the algorithm consists in loading the data from the CPU to the GPU's global memory. This step is computational time consuming because of the CPU's RAM bandwidth.

Once data is loaded in global memory, the data grid is decomposed into threads and thread blocks. One grid point in the physical space is linked to one thread.

Each thread is identified by a thread number which depends on the row and column of the thread. The figure 8 shows the physical grid vs. the computing threads and thread blocks. The number of threads per block, which is limited by available registers and shared memory, is set in order to obtain the maximal number of concurrent threads running on each multiprocessor.

Implementation details

This section is devoted to a detailed description of the LBM GPU implementation. The case used in this section concerns the lid-driven cavity which is fully described in the part 5.1 of the paper.

CPU program

}

The listing 1 presents a part of the CPU programming stored in a source file (i.e. D2Q9_LBGK.cu). The file is then compiled with nvcc.

The explicit GPU memory allocation uses the CUDA command cudaMalloc() (similarly, the deallocation command is cudaFree()). Copy from CPU to GPU is performed using cudaMemcpy(, , , cudaMemcpyHostToDevice); this operation being slow, it must of course be minimized.

The multi-threaded architecture of nVIDIA GPU uses thread blocks and grids.

The thread block size and grid size are defined using respectively dim3 threads(, , )

and dim3 grid(, ). In order to optimize the memory access, the thread block is an array which size is a multiple of 16. The grid size is then calculated in order to execute the kernel correctly (the figure 8 shows the decomposition of the domain).

The kernels are executed using kernel<<< grid, threads >>> (...). In our code, there are two kernels considering first, the collision/propagation step and second, the exchange of information across the boundaries of the grid of threads. f e q 1=rho * ( 1 . + 3 . 0 * vx + 4 . 5 * vx * vx -s q u a r e ) ; f e q 3=f e q 1 -6.0 * vx * rho ; f e q 2=rho * ( 1 . + 3 . 0 * vy + 4 . 5 * vy * vy -s q u a r e ) ; f e q 4=f e q 2 -6.0 * vy * rho ;

Collision/propagation kernel

. . . . . . . .

{ F1 OUT [ 0 ] = F1 IN ; F3 OUT [ tx -1]=F3 IN ; F5 OUT[ 0 ] = F5 IN ; F6 OUT [ tx -1]=F6 IN ; F7 OUT [ tx -1]=F7 IN ; F8 OUT[ 0 ] = F8 IN ; } e l s e { F1 OUT [ tx +1]=F1 IN ; F3 OUT [ tx -1]=F3 IN ; F5 OUT [ tx +1]=F5 IN ; F6 OUT [ tx -1]=F6 IN ; F7 OUT [ tx -1]=F7 IN ; F8 OUT [ tx +1]=F8 IN ; } // S y n c h r o n i z e source file collision_propagation_kernel.cu.
The k parameter is the index array of the data and is calculated depending on gridID and threadID. The main idea of the kernel is to use the shared memory in order to calculate the new density distributions stored in an array

Fi_OUT[NT] which is in the shared memory. This allows to execute the propagation during the data copy to Fi_OUT[NT]. Of course, there is no propagation along the block boundaries. Then, the unknown densities at these boundaries are used to store the known ones that must pass through the block boundaries.

The next kernel is used to pass information across the blocks boundaries.

The __syncthreads() function is necessary for the multiprocessor to wait for the execution of all threads before transferring the results from the array in shared memory to the global memory.

Exchange kernel

Numerical Results

GPUs that offer support for single precision floating point arithmetic do not meet all the operations of the IEEE 754 standards [12]. It may be argued that the precision of results obtained via LBM simulation using GPUs are thus suspect. To demonstrate that this is not the case, simulations of LBM lid driven cavity are compared with the results of Ghia et al. [START_REF] Ghia | High-Re solutions for incompressible flow using the Navier-Stokes equations and a multigrid method[END_REF].

The figures 11 and 12 show the comparisons between the LBM with GPU computations and the results from [START_REF] Ghia | High-Re solutions for incompressible flow using the Navier-Stokes equations and a multigrid method[END_REF] and for R e = 1000. The LBM with GPU is precise enough to predict correctly the flow in the lid driven cavity.

Performances of LBM with GPU

Table 1 presents the performance of the implementation using single-precision floating point, measured in MLUPS (Million Lattice site Update Per Second), function of the LBM mesh grid size and the number of threads per block. From this table, two main conclusions can be done. First, the number of threads must be at least 128 to have good calculation efficiency. Second, the size of the mesh grid must be at least 512 2 : this is due to the streaming multiprocessors. They all must be used for calculation to obtain a good GPU productivity.

Table 2 presents the performance of the implementation using GPU double precision floating-point. Of course, the number of lattice sites updated per second is quite lower than for simple precision. The mean multiplication factor between the two arithmetics is about 3.8.

22

We have presented a general purpose LBM simulation fully implemented on a single GPU. We have tested our GPU implementation using GT200 processing unit. The case tested was the well known lid-driven cavity problem.

With the use of single-precision floating point numbers instead of doubleprecision floating point numbers, the accuracy of the results remains satisfactory. Moreover, the computational time is 3. 11 Vertical velocity at y = 0.5 for R e = 1000 -circle are data from [START_REF] Ghia | High-Re solutions for incompressible flow using the Navier-Stokes equations and a multigrid method[END_REF] and line is LBM. 40

12 Horizontal velocity at x = 0.5 for R e = 1000 -circle are data from [START_REF] Ghia | High-Re solutions for incompressible flow using the Navier-Stokes equations and a multigrid method[END_REF] and line is LBM. 41 Fig. 11. Vertical velocity at y = 0.5 for R e = 1000 -circle are data from [START_REF] Ghia | High-Re solutions for incompressible flow using the Navier-Stokes equations and a multigrid method[END_REF] and line is LBM.

  boundariesFollowing the pseudo-code description, in the first phase, one thread is used per grid node.

Listing 1 .

 1 CPU program int main ( int argc , char * * argv ) { // S e t s i z e s i z e m a t = nx * ny ; mem size mat = s i z e o f ( f l o a t ) * s i z e m a t ; // CPU memory a l l o c a t i o n f 0 = ( f l o a t * ) m a l l o c ( mem size mat ) ; . . . . . . . . . . f 8 = ( f l o a t * ) m a l l o c ( mem size mat ) ; unsigned int mem size mat ch ar = s i z e o f ( char ) * s i z e m a t ; geo = ( char * ) m a l l o c ( mem s ize mat char ) ; // GPU memory a l l o c a t i o n f l o a t * f 0 d e v O l d = NULL; CUDA SAFE CALL( cudaMalloc ( ( void * * ) &f 0 d e v O l d , mem size mat ) ) ; . . . . . . . . . . f l o a t * f 8 d e v O l d = NULL; CUDA SAFE CALL( cudaMalloc ( ( void * * ) &f 8 d e v O l d , mem size mat ) ) ; f l o a t * f0 dev New = NULL; CUDA SAFE CALL( cudaMalloc ( ( void * * ) &f0 dev New , mem size mat ) ) ; . . . . . . . . . . f l o a t * f8 dev New = NULL; CUDA SAFE CALL( cudaMalloc ( ( void * * ) &f8 dev New , mem size mat ) ) ; char * g e o d e v = NULL; CUDA SAFE CALL( cudaMalloc ( ( void * * ) &geo dev , mem size mat ch ar ) ) ; // I n i t i a l i z e i n i t ( ) ; i n i t g e o ( ) ; // Copy d a t a from CPU t o GPU CUDA SAFE CALL( cudaMemcpy ( f 0 d e v O l d , f0 , mem size mat , cudaMemcpyHostToDevice ) ) ; . . . . . . . . . . CUDA SAFE CALL( cudaMemcpy ( f 8 d e v O l d , f8 , mem size mat , cudaMemcpyHostToDevice ) ) ; CUDA SAFE CALL( cudaMemcpy ( f0 dev New , f0 , mem size mat , cudaMemcpyHostToDevice ) ) ; . . . . . . . . . . CUDA SAFE CALL( cudaMemcpy ( f8 dev New , f8 , mem size mat , cudaMemcpyHostToDevice ) ) ; CUDA SAFE CALL( cudaMemcpy ( geo dev , geo , mem size mat char , cudaMemcpyHostToDevice ) ) ; // D e f i n e b l o c k and g r i d s i z e s dim3 t h r e a d s ( num threads , 1 , 1 ) ; dim3 g r i d 1 ( nx/ num threads , ny ) ; dim3 g r i d 2 ( 1 , ny/ num threads ) ; while ( t<t max ) { // Execute k e r n e l c o l l i s i o n p r o p a g a t i o n c o l l i s i o n p r o p a g a t i o n <<< g r i d 1 , t h r e a d s >>> ( nx , ny , num threads , tau , geo dev , f 0 d e v O l d , f 1 d e v O l d , f 2 d e v O l d , f 3 d e v O l d , f 4 d e v O l d , f 5 d e v O l d , f 6 d e v O l d , f 7 d e v O l d , f 8 d e v O l d , f0 dev New , f1 dev New , f2 dev New , f3 dev New , f4 dev New , f5 dev New , f6 dev New , f7 dev New , f8 dev New ) ; // Execute k e r n e l exchange exchange<<< g r i d 2 , t h r e a d s >>> ( nx , ny , num threads , f1 dev New , f3 dev New , f5 dev New , f6 dev New , f7 dev New , f8 dev New ) ; } // Copy r e s u l t s back t o CPU CUDA SAFE CALL( cudaMemcpy ( f0 , f 0 d e v O l d , mem size mat , cudaMemcpyDeviceToHost ) ) ; . . . . . . . . . . CUDA SAFE CALL( cudaMemcpy ( f8 , f 8 d e v O l d , mem size mat , cudaMemcpyDeviceToHost ) ) ; . . . . . . . . . .

Listing 2 .

 2 Collision/propagation kernel g l o b a l void c o l l i s i o n p r o p a g a t i o n ( int nx , int ny , int num threads , f l o a t tau , char * geoD , . . . ) { // Setup i n d e x i n g int tx = t h r e a d I d x . x ; int bx = b l o c k I d x . x ; int by = b l o c k I d x . y ; int x S t a r t = tx + bx * num threads ; int y S t a r t = by ; int k = nx * y S t a r t+x S t a r t ; // A l l o c a t e s h a r e d memory s h a r e d f l o a t F1 OUT [NT ] ; . . . . . . . . s h a r e d f l o a t F8 OUT [NT ] ; s h a r e d f l o a t F0 IN=f 0 O l d [ k ] ; . . . . . . . . s h a r e d f l o a t F8 IN=f 8 O l d [ k ] ; // Check i f i t i s a f l u i d or boundary node i f ( geoD [ k ] == FLUID) // C o l l i s i o n { rho=F0 IN+F1 IN+F2 IN+F3 IN+F4 IN+F5 IN+F6 IN +F7 IN+F8 IN ; vx=(F1 IN-F3 IN+F5 IN+F8 IN-F6 IN-F7 IN ) / rho ; vy=(F2 IN-F4 IN+F5 IN+F6 IN-F7 IN-F8 IN ) / rho ; s q u a r e =1.5 * ( vx * vx +vy * vy ) ; f e q 0 = 4 . / 9 . * rho * ( 1 . -s q u a r e ) ; rho * =0.1111111111111111111111;

F0

  IN+=(f e q 0 -F0 IN ) * t a u i n v ; . . . . . . . . F8 IN+=(f e q 8 -F8 IN ) * t a u i n v ; } e l s e i f ( geoD [ k ] == SET U) // V e l o c i t y boundary c o n d i t i o n { . . . . . . . . } e l s e i f ( geoD [ k ] == WALL) // Wall boundary c o n d i t i o n { . . . . . . . . } // Write t o s h a r e d memory and P r o p a g a t i o n i f ( tx==0) { F1 OUT [ tx +1]=F1 IN ; F3 OUT [ num threads -1]=F3 IN ; F5 OUT [ tx +1]=F5 IN ; F6 OUT [ num threads -1]=F6 IN ; F7 OUT [ num threads -1]=F7 IN ; F8 OUT [ tx +1]=F8 IN ; } e l s e i f ( tx==num threads -1)

Listing 3 .

 3 Exchange kernel g l o b a l void exchange ( int nx , int ny , int num threads , . . . . . . . . ) { // Setup i n d e x i n g int nbx=nx / num threads ; int num threads1 = blockDim . x ; int by = b l o c k I d x . y ; int tx = t h r e a d I d x . x ; int bx ; int x S t a r t , y S t a r t ; int xStartW , xTargetW ; int xStartE , xTargetE ; int kStartW , kTargetW ; int kStartE , kTargetE ; // Exchange a c r o s s b o u n d a r i e s for ( bx=0; bx<nbx-1 ; bx++) { x S t a r t = bx * num threads ; xStartW = x S t a r t +2 * num threads -1; xTargetW = xStartW-num threads ; y S t a r t = ( by ) * num threads1 + tx ; kStartW = nx * y S t a r t+xStartW ; kTargetW = nx * y S t a r t+xTargetW ; f3 New [ kTargetW ] = f3 New [ kStartW ] ; f6 New [ kTargetW ] = f6 New [ kStartW ] ; f7 New [ kTargetW ] = f7 New [ kStartW ] ; } for ( bx=nbx -2; bx>=0 ; bx--) { x S t a r t = bx * num threads ; xStartE = x S t a r t ; xTargetE = xStartE+num threads ;

Fig. 1 .Fig. 2 .

 12 Fig.1. Performances of CPUs (circle) and GPUs (diamond) over the last few years -extracted from[12]

Fig. 3 .

 3 Fig. 3. LBM partition: collision step (left) and propagation step (right).

Fig. 4 .Fig. 5 .

 45 Fig. 4. GPU hardware architecture overview.

Fig. 6 .

 6 Fig.6. GPU programming interface.

Fig. 7 .

 7 Fig. 7. Kernel memory access.

Fig. 8 .Fig. 9 .Fig. 10 .

 8910 Fig. 8. Physical grid and thread blocks and grids.

Table 1

 1 Performance of implementation in MLUPS -simple precision floating-point. Single vs double precision floating point arithmetic: sp for single precision and dp for double precision. 39

	8 times less with simple-precision

Presentation of the test case

In order to test the implementation of the LBM model on the GPU, the lid driven cavity case is used. This case has been chosen because it has been extensively studied in the literature. The figure 9 present the lid driven cavity problem with the boundary conditions.

In order to check the convergence of the simulation, the norm used . is: Both double and single precision floating-point calculations have been carried out using the GT200 GPU. The case tested is the lid-driven cavity problem at R e = 1000.

The figure 10 shows the norm evolution of the horizontal and vertical components of the velocity. The calculation has been performed for 2×10 6 iterations.

The horizontal asymptote of the velocity norm is close to the limit that can be obtained with single precision floating-point (i.e. 1.68 × 10 -7 ). Of course, the double precision floating-point allows to obtain a lower convergence criteria than for the floating-point. The single-precision calculation exhibits oscillations of the order of 1/2 24 = 6 × 10 -8 . However, the convergence is assured because of the numerical scheme stability.

From a numerical point of view, the maximum velocity magnitude difference between the two calculations is about 10 -3 m/s i.e. a maximum relative difference of 10 -2 which validates the use of single precision arithmetic with GPUs for LBM calculations. Fig. 12. Horizontal velocity at x = 0.5 for R e = 1000 -circle are data from [START_REF] Ghia | High-Re solutions for incompressible flow using the Navier-Stokes equations and a multigrid method[END_REF] and line is LBM.